2 Replies Latest reply on Mar 9, 2010 2:47 PM by nou

    compiler crashing

    nou
      on GPU. CPU is fine

      attached kernel crash when i try build it.

      0   llc             0x00000000009e18cf
      1   llc             0x00000000009e2aa9
      2   libpthread.so.0 0x00007fcc93134080
      3   llc             0x00000000005ffdca
      4   llc             0x00000000006033e1
      5   llc             0x00000000005f676a
      6   llc             0x00000000007d251a
      7   llc             0x00000000007d2d60
      8   llc             0x00000000007d2f9c
      9   llc             0x000000000041ba06
      10  libc.so.6       0x00007fcc92dd15a6 __libc_start_main + 230
      11  llc             0x0000000000419779
      Stack dump:
      0.    Program arguments: /usr/local/bin/x86_64/llc -mcpu=teddy -mattr=mwgs-3-256-1-1 -regalloc=linearscan -mtriple=amdil-pc-amdopencl /tmp/OCLWjJuMD.bc -f -o /tmp/OCLWjJuMD.il
      1.    Running pass 'AMD IL Control Flow Graph structurizer Pass' on function '@__OpenCL_raytrace_kernel'
      Segmentation fault


      BUILD LOG:
      ====================
      llc has failed

      windows crash too.

      #define EPSILON 1e-6f #define MAX_TRIANGLES_PER_BVH_LEAF 16 // TODO #define MAX_BVH_DEPTH 20 // TODO typedef struct _Ray { float4 o; float4 d; int4 sign; float4 invDir; } Ray; typedef struct _Camera { float4 orig; float4 dir; float4 up; float4 right; uint2 dimXY; float fovY; // rad } Camera; typedef struct _Vertex { float4 position; // 4B float4 normal; // 4B float4 texcoord12; // 4B float4 texcoord34; // 4B } Vertex; typedef struct _AABB { float4 corners[2]; } AABB; typedef struct _Triangle { AABB aabb; uint vertexAidx; uint vertexBidx; uint vertexCidx; uint padding; } Triangle; typedef struct _TriangleHit { float u; float v; float t; int triIndex; } TriangleHit; typedef struct _BVHNODE { AABB aabb; // 8B uint firstTriangleIndex; // 4B uint packedNumTrisDepthLeaf; // 4B } BVHNode; float4 AABBextents(const AABB* aabb) { return aabb->corners[1] - aabb->corners[0]; } float4 AABBcentre(const AABB* aabb) { return (aabb->corners[0] + aabb->corners[1]) * 0.5f; } void createRayFromPixel(const uint x, const uint y, const Camera* cam, Ray* ray) { float4 dir, dx, dy, forward; float mx = (float)((x - cam->dimXY.x * 0.5f) * (1.0f / cam->dimXY.y) * tan(cam->fovY)); float my = (float)((y - cam->dimXY.y * 0.5f) * (1.0f / cam->dimXY.y) * tan(cam->fovY)); dx = cam->right * mx; dy = cam->up * my; forward = normalize(cross(cam->up, cam->right)); dir = forward + dx*2.0f + dy*2.0f; ray->o = cam->orig; ray->d = normalize(dir); ray->invDir = native_divide((float4)(1.0f), ray->d); ray->sign = ray->invDir < (float4)(0); } float rayTriangle(const Ray* r, const float4 A, const float4 B, const float4 C, float* u, float* v) { float4 edge1 = B - A; // TODO - optimise float4 edge2 = C - A; float4 tvec = r->o - A; float4 pvec = cross(r->d, edge2); float det = dot(edge1, pvec); //det = __fdividef(1.0f, det); det = 1.0f / det; *u = dot(tvec, pvec) * det; if (*u < 0.0f || *u > 1.0f) return -1.0f; float4 qvec = cross(tvec, edge1); *v = dot(r->d, qvec) * det; if (*v < 0.0f || (*u + *v) > 1.0f) return -1.0f; return dot(edge2, qvec) * det; } bool getAABBRayIntersection(const Ray* r, const AABB* aabb, float* tIn) { float tmin, tmax, tymin, tymax, tzmin, tzmax; tmin = (aabb->corners[r->sign[0]][0] - r->o[0]) * r->invDir[0]; tmax = (aabb->corners[1-r->sign[0]][0] - r->o[0]) * r->invDir[0]; tymin = (aabb->corners[r->sign[1]][1] - r->o[1]) * r->invDir[1]; tymax = (aabb->corners[1-r->sign[1]][1] - r->o[1]) * r->invDir[1]; if ( (tmin > tymax) || (tymin > tmax) ) return false; if (tymin > tmin) tmin = tymin; if (tymax < tmax) tmax = tymax; tzmin = (aabb->corners[r->sign[2]][2] - r->o[2]) * r->invDir[2]; tzmax = (aabb->corners[1-r->sign[2]][2] - r->o[2]) * r->invDir[2]; if ( (tmin > tzmax) || (tzmin > tmax) ) return false; if (tzmin > tmin) tmin = tzmin; if (tzmax < tmax) tmax = tzmax; *tIn=tmin; return (tmin < MAXFLOAT) && (tmax > EPSILON); } void stackPush(uint* stack, uint* head, uint val) { stack[++(*head)] = val; } uint stackPop(uint* stack, uint* head) { return stack[(*head)--]; } bool stackEmpty(uint* stack, uint* head) { return (*head) == 0; } bool intersectAllTriangles( const Ray* ray, const uint numTris, const uint firstTriIndex, __global const Triangle* triangles, __global const unsigned int* triangleIndices, __global const Vertex* vertices, TriangleHit* thOut) { TriangleHit minHit; minHit.t = MAXFLOAT; float u, v; for(uint f = 0; f < numTris; f++) { uint triIdx = triangleIndices[firstTriIndex+f]; Vertex A = vertices[triangles[triIdx].vertexAidx]; // optimise!! Vertex B = vertices[triangles[triIdx].vertexBidx]; // TODO - global!? Vertex C = vertices[triangles[triIdx].vertexCidx]; float t = rayTriangle(ray, A.position, B.position, C.position, &u, &v); if(t > 0.0f) { if(t < minHit.t) { minHit.t = t; minHit.u = u; minHit.v = v; minHit.triIndex = triIdx; } } } if(minHit.t < MAXFLOAT) { *thOut = minHit; return true; } return false; } float4 tracePrimaryRay( Ray* ray, const uint numTriangles, const uint numVertices, const uint maxBVHdepth, __global const Triangle* triangles, __global const uint* triangleIndices, __global const Vertex* vertices, __global const BVHNode* bvhNodes) { float4 outColour; __private uint stack[24]; uint stackHead; TriangleHit th, minTh; minTh.t = MAXFLOAT; float hitDist; bool hit; //float4 lightPos = (float4)(10, 100, 10, 0); // TODO //float4 N, P; //float2 TC; //float4 toLight; stackPush(stack, &stackHead, 0); while(!stackEmpty(stack, &stackHead)) { unsigned int nodeIndex = stackPop(stack, &stackHead); BVHNode node = bvhNodes[nodeIndex]; hit = getAABBRayIntersection(ray, &node.aabb, &hitDist); if(!hit || hitDist > minTh.t) continue; uint numtris = (node.packedNumTrisDepthLeaf >> 8) & 0x00FFFFFF; uint depth = (node.packedNumTrisDepthLeaf >> 2) & 0x0000003F; uint leftRightLeaf = (node.packedNumTrisDepthLeaf) & 0x00000003; if(leftRightLeaf == 0) { // LEAF //ray.minDist = hitDist; bool hit = intersectAllTriangles(ray, numtris, node.firstTriangleIndex, triangles, triangleIndices, vertices, &th); if(hit) { if(th.t < minTh.t) { minTh = th; } } continue; } unsigned int leftIndex = nodeIndex + 1; unsigned int rightIndex = nodeIndex + (1 << (maxBVHdepth - depth)); stackPush(stack, &stackHead, leftIndex); stackPush(stack, &stackHead, rightIndex); } // while if(minTh.t < MAXFLOAT) { outColour = (float4)(0,0,0,0); } else outColour = (float4)(1,1,1,0); // TODO - background colour return outColour; } uint camMap2Drow(const Camera* cam, uint x, uint y) { return y*cam->dimXY.x + x; } uint camMap2Dcol(const Camera* cam, uint x, uint y) { return x*cam->dimXY.y + y; } __kernel void raytrace( const Camera cam, const uint numTriangles, const uint numVertices, const uint maxBVHdepth, __global const Triangle* triangles, __global const uint* triangleIndices, __global const Vertex* vertices, __global const BVHNode* bvhNodes, __global float4 * img) { unsigned int tidx = get_global_id(0); unsigned int tidy = get_global_id(1); Ray ray; createRayFromPixel(tidx, tidy, &cam, &ray); //TriangleHit th; float4 color = tracePrimaryRay(&ray, numTriangles, numVertices, maxBVHdepth, triangles, triangleIndices, vertices, bvhNodes); //img[camMap2Drow(&cam, tidx, tidy)] = color; //img[tidx * cam.dimXY.y + tidy] = (float4)1.0f/mint; //img[tidx * cam.dimXY.y + tidy] = (float4)((float)tidx/128.0f, (float)tidy/128, 0, 1); }

        • compiler crashing
          MicahVillmow
          nou,
          This does not compile with an internal build do to these errors.
          bug_forum.cl(201): error: vector subscript not support
          tmin = (aabb->corners[r->sign[0]][0] - r->o[0]) * r->invDir[0];
          ^

          bug_forum.cl(201): error: vector subscript not support
          tmin = (aabb->corners[r->sign[0]][0] - r->o[0]) * r->invDir[0];
          ^

          bug_forum.cl(201): error: vector subscript not support
          tmin = (aabb->corners[r->sign[0]][0] - r->o[0]) * r->invDir[0];
          ^

          bug_forum.cl(203): error: vector subscript not support
          tmax = (aabb->corners[1-r->sign[0]][0] - r->o[0]) * r->invDir[0];
          ^

          bug_forum.cl(203): error: vector subscript not support
          tmax = (aabb->corners[1-r->sign[0]][0] - r->o[0]) * r->invDir[0];
          ^

          bug_forum.cl(203): error: vector subscript not support
          tmax = (aabb->corners[1-r->sign[0]][0] - r->o[0]) * r->invDir[0];
          ^

          bug_forum.cl(205): error: vector subscript not support
          tymin = (aabb->corners[r->sign[1]][1] - r->o[1]) * r->invDir[1];
          ^

          bug_forum.cl(205): error: vector subscript not support
          tymin = (aabb->corners[r->sign[1]][1] - r->o[1]) * r->invDir[1];
          ^

          bug_forum.cl(205): error: vector subscript not support
          tymin = (aabb->corners[r->sign[1]][1] - r->o[1]) * r->invDir[1];
          ^

          bug_forum.cl(207): error: vector subscript not support
          tymax = (aabb->corners[1-r->sign[1]][1] - r->o[1]) * r->invDir[1];
          ^

          bug_forum.cl(207): error: vector subscript not support
          tymax = (aabb->corners[1-r->sign[1]][1] - r->o[1]) * r->invDir[1];
          ^

          bug_forum.cl(207): error: vector subscript not support
          tymax = (aabb->corners[1-r->sign[1]][1] - r->o[1]) * r->invDir[1];
          ^

          bug_forum.cl(221): error: vector subscript not support
          tzmin = (aabb->corners[r->sign[2]][2] - r->o[2]) * r->invDir[2];
          ^

          bug_forum.cl(221): error: vector subscript not support
          tzmin = (aabb->corners[r->sign[2]][2] - r->o[2]) * r->invDir[2];
          ^

          bug_forum.cl(221): error: vector subscript not support
          tzmin = (aabb->corners[r->sign[2]][2] - r->o[2]) * r->invDir[2];
          ^

          bug_forum.cl(223): error: vector subscript not support
          tzmax = (aabb->corners[1-r->sign[2]][2] - r->o[2]) * r->invDir[2];
          ^

          bug_forum.cl(223): error: vector subscript not support
          tzmax = (aabb->corners[1-r->sign[2]][2] - r->o[2]) * r->invDir[2];
          ^

          bug_forum.cl(223): error: vector subscript not support
          tzmax = (aabb->corners[1-r->sign[2]][2] - r->o[2]) * r->invDir[2];
          ^


          It is not legal in OpenCL to index into a vector. A vector is not an array and array syntax cannot be used on them.