5 Replies Latest reply on Apr 26, 2010 4:44 PM by omkaranathan

    Strange error message during compilation

    spectral

      During the compilation of my opencl code, I got the following error :

       

      "The handle could not be opened"
      "during redirection of handle 1."

      Is it a problem with the output console ?I use the kernel32.AllocConsole for this !



        • Strange error message during compilation
          omkaranathan

          Could you explain a bit more about your application?

          This seems to be a VC++/.NET error, and not an error produced by OpenCL code. 

            • Strange error message during compilation
              spectral

              The error is displayed when I call "buildprogram" !

              When I do a "get-info" with the "build-log" I got the following log from the OpenCL driver :

               

              -----------------------------------------------
              Intel(R) Core(TM)2 Duo CPU     E7400  @ 2.80GHz
              -----------------------------------------------
              Internal error: Compilation failed.



                • Strange error message during compilation
                  omkaranathan

                  Could you post your kernel code?

                   

                    • Strange error message during compilation
                      spectral

                      Yes,

                       

                      I have attach the code

                       

                       

                      typedef struct __attribute__ ((packed)) _Ray { float4 Origin; float4 Direction; int IsShadowRay; int pad1; int pad2; int pad3; } Ray; typedef struct __attribute__ ((packed)) _IntersectionResult { float Maximum; int PrimitiveId; float U; float V; int CurrentInstanceId; int InstanceId; float Time; int pad1; } IntersectionResult; typedef struct __attribute__ ((packed)) _InstanceList { int Instances_Index; int Instances_Length; int pad1; int pad2; } InstanceList; typedef struct __attribute__ ((packed)) _MovingMatrix4 { int Transforms_Index; int Transforms_Length; float _t0; float _t1; float _inv; int pad1; int pad2; int pad3; } MovingMatrix4; typedef struct __attribute__ ((packed)) _PrimitivesList { int Type; int Index; int pad1; int pad2; } PrimitivesList; typedef struct __attribute__ ((packed)) _Geometry { int AcceleratorBVH_Index; int AcceleratorNA_Index; int pad1; int pad2; } Geometry; typedef struct __attribute__ ((packed)) _Instance { MovingMatrix4 o2w; MovingMatrix4 w2o; Geometry Geometry; } Instance; typedef struct __attribute__ ((packed)) _TriangleMesh { int Points_Index; int Points_Length; int Indices_Index; int Indices_Length; } TriangleMesh; typedef struct __attribute__ ((packed)) _Matrix4 { float M00; float M01; float M02; float M03; float M10; float M11; float M12; float M13; float M20; float M21; float M22; float M23; float M30; float M31; float M32; float M33; } Matrix4; typedef struct __attribute__ ((packed)) _Accelerator_BVH { int List_Index; int Nodes_Index; int Nodes_Length; int PrimitivesIds_Index; int PrimitivesIds_Length; int pad1; int pad2; int pad3; } Accelerator_BVH; typedef struct __attribute__ ((packed)) _Accelerator_NA { int PrimitiveIndex; int pad1; int pad2; int pad3; } Accelerator_NA; typedef struct __attribute__ ((packed)) _AABB { float MinX, MinY, MinZ, MaxX, MaxY, MaxZ; } AABB; typedef struct __attribute__ ((packed)) _BoundingVolume { AABB BBox; int SkipNodeIndex; int PrimitiveId; ushort Count; ushort pad1; int pad2; int pad3; int pad4; } BoundingVolume; typedef struct { float4 InverseDirection; int IsXNegative; int IsYNegative; int IsZNegative; } MITData; typedef struct { __global Ray * Rays; __global IntersectionResult * Intersections; __global InstanceList * InstancesList; __global PrimitivesList * PrimitivesList; __global Instance * Instances; __global Matrix4 * Matrix4; __global TriangleMesh * TriangleMeshes; __global float * Points; __global int * Indices; __global Accelerator_BVH * AcceleratorBVHs; __global BoundingVolume * AcceleratorBVH_Nodes; __global int * AcceleratorBVH_PrimitivesIds; __global Accelerator_NA * AcceleratorNAs; int RootAcceleratorIndex; } GlobalBuffer; static Matrix4 Matrix4_Blend(__global Matrix4 * m0, __global Matrix4 * m1, float t) { Matrix4 m; m.M30 = m.M31 = m.M32 = m.M33 = 1.f; float tinv = 1.f - t; m.M00 = tinv * m0->M00 + t * m1->M00; m.M01 = tinv * m0->M01 + t * m1->M01; m.M02 = tinv * m0->M02 + t * m1->M02; m.M03 = tinv * m0->M03 + t * m1->M03; m.M10 = tinv * m0->M10 + t * m1->M10; m.M11 = tinv * m0->M11 + t * m1->M11; m.M12 = tinv * m0->M12 + t * m1->M12; m.M13 = tinv * m0->M13 + t * m1->M13; m.M20 = tinv * m0->M20 + t * m1->M20; m.M21 = tinv * m0->M21 + t * m1->M21; m.M22 = tinv * m0->M22 + t * m1->M22; m.M23 = tinv * m0->M23 + t * m1->M23; return m; } static Matrix4 MovingMatrix_Sample(GlobalBuffer GB, __global MovingMatrix4 * movingMatrix, float time) { if (movingMatrix->Transforms_Length == 1 || movingMatrix->_t0 >= movingMatrix->_t1) return GB.Matrix4[movingMatrix->Transforms_Index]; float nt = (clamp(time, movingMatrix->_t0, movingMatrix->_t1) - movingMatrix->_t0) * movingMatrix->_inv; int idx0 = (int)nt; int idx1 = min(idx0 + 1, movingMatrix->Transforms_Length - 1); return Matrix4_Blend( &GB.Matrix4[movingMatrix->Transforms_Index + idx0], &GB.Matrix4[movingMatrix->Transforms_Index + idx1], (float)(nt - idx0)); } #define Matrix4_TransformPX(m,x,y,z) (m->M00 * x + m->M01 * y + m->M02 * z + m->M03) #define Matrix4_TransformPY(m,x,y,z) (m->M10 * x + m->M11 * y + m->M12 * z + m->M13) #define Matrix4_TransformPZ(m,x,y,z) (m->M20 * x + m->M21 * y + m->M22 * z + m->M23) #define Matrix4_TransformVX(m,x,y,z) (m->M00 * x + m->M01 * y + m->M02 * z) #define Matrix4_TransformVY(m,x,y,z) (m->M10 * x + m->M11 * y + m->M12 * z) #define Matrix4_TransformVZ(m,x,y,z) (m->M20 * x + m->M21 * y + m->M22 * z) static Ray Ray_Transform(Matrix4 * m, __global Ray * ray) { Ray r; if (isnan(m->M00)) { r.Origin.x = ray->Origin.x; r.Origin.y = ray->Origin.y; r.Origin.z = ray->Origin.z; r.Direction.x = ray->Direction.x; r.Direction.y = ray->Direction.y; r.Direction.z = ray->Direction.z; return r; } r.Origin.x = Matrix4_TransformPX(m, ray->Origin.x, ray->Origin.y, ray->Origin.z); r.Origin.y = Matrix4_TransformPY(m, ray->Origin.x, ray->Origin.y, ray->Origin.z); r.Origin.z = Matrix4_TransformPZ(m, ray->Origin.x, ray->Origin.y, ray->Origin.z); r.Direction.x = Matrix4_TransformVX(m, ray->Direction.x, ray->Direction.y, ray->Direction.z); r.Direction.y = Matrix4_TransformVY(m, ray->Direction.x, ray->Direction.y, ray->Direction.z); r.Direction.z = Matrix4_TransformVZ(m, ray->Direction.x, ray->Direction.y, ray->Direction.z); return r; } static int Solvers_SolveQuadric(float a, float b, float c, float * results) { float disc = b * b - 4 * a * c; if (disc < 0) return 0; disc = sqrt(disc); float q = ((b < 0) ? -0.5f * (b - disc) : -0.5f * (b + disc)); float t0 = q / a; float t1 = c / q; if (t0 > t1) { results[0] = t1; results[1] = t0; } else { results[0] = t0; results[1] = t1; } return 1; } static int Intersects_Sphere(GlobalBuffer GB, float4 origin, float4 direction, int primitiveId) { float qa = direction.x * direction.x + direction.y * direction.y + direction.z * direction.z; float qb = 2.f * ((direction.x * origin.x) + (direction.y * origin.y) + (direction.z * origin.z)); float qc = ((origin.x * origin.x) + (origin.y * origin.y) + (origin.z * origin.z)) - 1.f; float t[2]; if (!Solvers_SolveQuadric(qa, qb, qc, &t[0])) return 0; int GID = get_global_id(0); if (t[0] >= GB.Intersections[GID].Maximum || t[1] <= 0) return 0; if (t[0] > 0) GB.Intersections[GID].Maximum = (float)t[0]; else GB.Intersections[GID].Maximum = (float)t[1]; GB.Intersections[GID].PrimitiveId = 0; GB.Intersections[GID].InstanceId = GB.Intersections[GID].CurrentInstanceId; return 1; } static int Intersects_BoxRay(GlobalBuffer GB, MITData * mitData, float4 origin, float4 direction, __global AABB* aabb, float * minHit, float * maxHit) { float tmin, tmax, tymin, tymax, tzmin, tzmax; if (mitData->IsXNegative) { tmin = (aabb->MaxX - origin.x) * mitData->InverseDirection.x; tmax = (aabb->MinX - origin.x) * mitData->InverseDirection.x; } else { tmin = (aabb->MinX - origin.x) * mitData->InverseDirection.x; tmax = (aabb->MaxX - origin.x) * mitData->InverseDirection.x; } if (mitData->IsYNegative) { tymin = (aabb->MaxY - origin.y) * mitData->InverseDirection.y; tymax = (aabb->MinY - origin.y) * mitData->InverseDirection.y; } else { tymin = (aabb->MinY - origin.y) * mitData->InverseDirection.y; tymax = (aabb->MaxY - origin.y) * mitData->InverseDirection.y; } if (tmin > tymax || tymin > tmax) { minHit[0] = MAXFLOAT; maxHit[0] = MAXFLOAT; return 0; } if (tymin > tmin) tmin = tymin; if (tymax < tmax) tmax = tymax; if (mitData->IsZNegative) { tzmin = (aabb->MaxZ - origin.z) * mitData->InverseDirection.z; tzmax = (aabb->MinZ - origin.z) * mitData->InverseDirection.z; } else { tzmin = (aabb->MinZ - origin.z) * mitData->InverseDirection.z; tzmax = (aabb->MaxZ - origin.z) * mitData->InverseDirection.z; } if (tmin > tzmax || tzmin > tmax) { minHit[0] = MAXFLOAT; maxHit[0] = MAXFLOAT; return 0; } if (tzmin > tmin) tmin = tzmin; if (tzmax < tmax) tmax = tzmax; minHit[0] = tmin; maxHit[0] = tmax; return 1; } static int Intersects_Primitive_TriangleMesh(GlobalBuffer GB, __global int * Indices, __global float * Points, float4 origin, float4 direction, int primitiveId) { int tri = 3 * primitiveId; int i1 = 3 * Indices[tri + 0]; int i2 = 3 * Indices[tri + 1]; int i3 = 3 * Indices[tri + 2]; float4 edge0 = (float4)( Points[i2 + 0] - Points[i1 + 0], Points[i2 + 1] - Points[i1 + 1], Points[i2 + 2] - Points[i1 + 2], 0); float4 edge1 = (float4)( Points[i1 + 0] - Points[i3 + 0], Points[i1 + 1] - Points[i3 + 1], Points[i1 + 2] - Points[i3 + 2], 0); float4 edge2 = (float4)( Points[i1 + 0] - origin.x, Points[i1 + 1] - origin.y, Points[i1 + 2] - origin.z, 0); float4 n = cross(edge0, edge1); float v = dot(direction, n); float iv = 1.0f / v; float va = dot(n, edge2); float t = iv * va; int GID = get_global_id(0); if (t <= 0 || t >= GB.Intersections[GID].Maximum) return 0; float4 i = cross(edge2, direction); float v1 = dot(i, edge1); float beta = iv * v1; if (beta < 0.f) return 0; float v2 = dot(i, edge0); if ((v1 + v2) * v > v * v) return 0; float gamma = iv * v2; if (gamma < 0.f) return 0; GB.Intersections[GID].Maximum = t; GB.Intersections[GID].U = beta; GB.Intersections[GID].V = gamma; GB.Intersections[GID].PrimitiveId = primitiveId; GB.Intersections[GID].InstanceId = GB.Intersections[GID].CurrentInstanceId; return 1; } static int Accelerator_BVH_TreeTraversal_TriangleMesh_Internal( GlobalBuffer GB, __global BoundingVolume * _nodes, __global int * primitivesIds, __global PrimitivesList * primitivesList, float4 origin, float4 direction) { int GID = get_global_id(0); int bvNodeIndex = 0; MITData mitData; mitData.InverseDirection.x = 1.0f / direction.x; mitData.InverseDirection.y = 1.0f / direction.y; mitData.InverseDirection.z = 1.0f / direction.z; if (mitData.InverseDirection.x < 0.f) mitData.IsXNegative = 1; else mitData.IsXNegative = 0; if (mitData.InverseDirection.y < 0.f) mitData.IsYNegative = 1; else mitData.IsYNegative = 0; if (mitData.InverseDirection.z < 0.f) mitData.IsZNegative = 1; else mitData.IsZNegative = 0; mitData.InverseDirection.w = 0; int stopNodeIndex = _nodes[bvNodeIndex].SkipNodeIndex; int hasIntersection = 0; if (primitivesList->Type != 1) return 0; TriangleMesh triangleMesh = GB.TriangleMeshes[primitivesList->Index]; while (bvNodeIndex < stopNodeIndex) { float minHit = 0; float maxHit = 0; int hasHit = Intersects_BoxRay(GB, &mitData, origin, direction, &_nodes[bvNodeIndex].BBox, &minHit, &maxHit); if (hasHit && minHit <= GB.Intersections[GID].Maximum) { if (_nodes[bvNodeIndex].PrimitiveId > -1) { if(_nodes[bvNodeIndex].Count < 1) { if (Intersects_Primitive_TriangleMesh(GB, &GB.Indices[triangleMesh.Indices_Index], &GB.Points[triangleMesh.Points_Index], origin, direction, _nodes[bvNodeIndex].PrimitiveId)) { if (GB.Rays[GID].IsShadowRay) return 1; hasIntersection = 1; } } else if (_nodes[bvNodeIndex].Count > 0) { int startIndex = _nodes[bvNodeIndex].PrimitiveId; int endIndex = startIndex + _nodes[bvNodeIndex].Count - 1; for (int index = startIndex; index <= endIndex; index++) { if (Intersects_Primitive_TriangleMesh(GB, &GB.Indices[triangleMesh.Indices_Index], &GB.Points[triangleMesh.Points_Index], origin, direction, primitivesIds[index])) { if (GB.Rays[GID].IsShadowRay) return 1; hasIntersection = 1; } } } } bvNodeIndex++; if ((bvNodeIndex >= stopNodeIndex || bvNodeIndex == _nodes[bvNodeIndex].SkipNodeIndex) && hasIntersection) return 1; } else bvNodeIndex = _nodes[bvNodeIndex].SkipNodeIndex; } return hasIntersection; } static int Accelerator_BVH_TreeTraversal_TriangleMesh(GlobalBuffer GB, __global Accelerator_BVH * _acceleratorBVHs, float4 origin, float4 direction) { return Accelerator_BVH_TreeTraversal_TriangleMesh_Internal( GB, &GB.AcceleratorBVH_Nodes[_acceleratorBVHs->Nodes_Index], &GB.AcceleratorBVH_PrimitivesIds[_acceleratorBVHs->PrimitivesIds_Index], &GB.PrimitivesList[_acceleratorBVHs->List_Index], origin, direction); } static int Accelerator_NoAcceleration(GlobalBuffer GB, __global Accelerator_NA * accelerator, float4 origin, float4 direction) { PrimitivesList primitiveList = GB.PrimitivesList[accelerator->PrimitiveIndex]; if (primitiveList.Type == 1) { TriangleMesh triangleMesh = GB.TriangleMeshes[primitiveList.Index]; int hasIntersection = 0; for(int i = 0; i < triangleMesh.Indices_Length/3; i++) hasIntersection &= Intersects_Primitive_TriangleMesh(GB, &GB.Indices[triangleMesh.Indices_Index], &GB.Points[triangleMesh.Points_Index], origin, direction, i); return hasIntersection; } if (primitiveList.Type == 2) { return Intersects_Sphere(GB, origin, direction, 0); } return 0; } static int Intersects_Geometry(GlobalBuffer GB, Geometry geometry, float4 origin, float4 direction) { if (geometry.AcceleratorNA_Index > -1) { return Accelerator_NoAcceleration(GB, &GB.AcceleratorNAs[geometry.AcceleratorNA_Index], origin, direction); } return 0; } static int Intersects_Instance(GlobalBuffer GB, __global Instance * instance, int primitiveId) { int GID = get_global_id(0); GB.Intersections[GID].CurrentInstanceId = primitiveId; Matrix4 transform = MovingMatrix_Sample(GB, &instance->w2o, GB.Intersections[GID].Time); Ray localRay = Ray_Transform(&transform, &GB.Rays[GID]); float4 origin = localRay.Origin; float4 direction = localRay.Direction; return Intersects_Geometry(GB, instance->Geometry, origin, direction); } static int Intersects_InstanceList(GlobalBuffer GB, __global InstanceList * instanceList, int primitiveId) { if (primitiveId < instanceList->Instances_Length) { __global Instance * instance = &GB.Instances[(instanceList->Instances_Index + primitiveId)]; return Intersects_Instance(GB, instance, primitiveId); } return 0; } static int Accelerator_BVH_TreeTraversal_InstanceList_Internal( GlobalBuffer GB, __global BoundingVolume * _nodes, __global int * primitivesIds, __global InstanceList * instancesList) { int GID = get_global_id(0); int bvNodeIndex = 0; MITData mitData; mitData.InverseDirection.x = 1.0f / GB.Rays[GID].Direction.x; mitData.InverseDirection.y = 1.0f / GB.Rays[GID].Direction.y; mitData.InverseDirection.z = 1.0f / GB.Rays[GID].Direction.z; if (mitData.InverseDirection.x < 0) mitData.IsXNegative = 1; else mitData.IsXNegative = 0; if (mitData.InverseDirection.y < 0) mitData.IsYNegative = 1; else mitData.IsYNegative = 0; if (mitData.InverseDirection.z < 0) mitData.IsZNegative = 1; else mitData.IsZNegative = 0; mitData.InverseDirection.w = 0; int stopNodeIndex = _nodes[bvNodeIndex].SkipNodeIndex; int hasIntersection = 0; while (bvNodeIndex < stopNodeIndex) { float minHit = 0; float maxHit = 0; int hasHit = Intersects_BoxRay(GB, &mitData, GB.Rays[GID].Origin, GB.Rays[GID].Direction, &_nodes[bvNodeIndex].BBox, &minHit, &maxHit); if (hasHit && minHit <= GB.Intersections[GID].Maximum) { if (_nodes[bvNodeIndex].PrimitiveId > -1) { if(_nodes[bvNodeIndex].Count < 1) { if(Intersects_InstanceList(GB, instancesList, _nodes[bvNodeIndex].PrimitiveId)) { if (GB.Rays[GID].IsShadowRay) return 1; hasIntersection = 1; } } else if (_nodes[bvNodeIndex].Count > 0) { int startIndex = _nodes[bvNodeIndex].PrimitiveId; int endIndex = startIndex + _nodes[bvNodeIndex].Count - 1; for (int index = startIndex; index <= endIndex; index++) { if(Intersects_InstanceList(GB, instancesList, primitivesIds[index])) { if (GB.Rays[GID].IsShadowRay) return 1; hasIntersection = 1; } } } } bvNodeIndex++; if ((bvNodeIndex >= stopNodeIndex || bvNodeIndex == _nodes[bvNodeIndex].SkipNodeIndex) && hasIntersection) return 1; } else bvNodeIndex = _nodes[bvNodeIndex].SkipNodeIndex; } return hasIntersection; } static int Accelerator_BVH_TreeTraversal_InstanceList(GlobalBuffer GB, __global Accelerator_BVH * _acceleratorBVHs) { return Accelerator_BVH_TreeTraversal_InstanceList_Internal(GB, &GB.AcceleratorBVH_Nodes[_acceleratorBVHs->Nodes_Index], &GB.AcceleratorBVH_PrimitivesIds[_acceleratorBVHs->PrimitivesIds_Index], &GB.InstancesList[_acceleratorBVHs->List_Index]); } __kernel void Trace( __global Ray * _rays, __global IntersectionResult * _intersections, int _rootAcceleratorIndex, __global InstanceList * _instancesList, __global PrimitivesList * _primitivesList, __global Instance * _instances, __global Matrix4 * _matrix4, __global TriangleMesh * _triangleMeshes, __global float * _points, __global int * _indices, __global Accelerator_BVH * _acceleratorBVHs, __global BoundingVolume * _acceleratorBVH_Nodes, __global int * _acceleratorBVH_PrimitivesIds, __global Accelerator_NA * _acceleratorNAs) { GlobalBuffer GB; GB.Rays = _rays; GB.Intersections = _intersections; GB.RootAcceleratorIndex = _rootAcceleratorIndex; GB.InstancesList = _instancesList; GB.PrimitivesList = _primitivesList; GB.Instances = _instances; GB.Matrix4 = _matrix4; GB.TriangleMeshes = _triangleMeshes; GB.Points = _points; GB.Indices = _indices; GB.AcceleratorBVHs = _acceleratorBVHs; GB.AcceleratorBVH_Nodes = _acceleratorBVH_Nodes; GB.AcceleratorBVH_PrimitivesIds = _acceleratorBVH_PrimitivesIds; GB.AcceleratorNAs = _acceleratorNAs; Accelerator_BVH_TreeTraversal_InstanceList(GB, &GB.AcceleratorBVHs[GB.RootAcceleratorIndex]); }