Archives Discussions

[spectral]

Hi,

In my kernel when I do this in a specific method... I got a crash :

    int GID = get_global_id(0);

    printf("%d", GID);

    return 0;

But, when I do this I have no crash :

    int GID = get_global_id(0);

    /* printf("%d", GID); */

    return 0;

I do this test because it sounds that I got some problems to access to the GID variable !!!

int GID = get_global_id(0); /* printf("%d", GID); */ return 0;

[genaganna]

Originally posted by: viewon01  

Hi,

In my kernel when I do this in a specific method... I got a crash :

     int GID = get_global_id(0);

    printf("%d", GID);

    return 0;

 But, when I do this I have no crash :

     int GID = get_global_id(0);

    /* printf("%d", GID); */

    return 0;

 I do this test because it sounds that I got some problems to access to the GID variable !!!

Printf is valid only for CL_DEVICE_TYPE_CPU. printf is not supported in windows.  It is known issue. developers are working on this.

[spectral]

Okay,

In fact my problem come that I got this error :

"Unhandled exception at 0x042b1056 (OCL6C3.tmp.dll) in RSRenderer.exe: 0xC0000005: Access violation reading location 0x3f512c80."

There is a memory access violation somewhere !

When I try to do this in some methods I got a crash:

GB.Intersections[0].Maximum = 0;

but when I do the same in th kernel all is fine.

Also, even when I do this line in the kernel... all is fine... but I got a crash from other methods

printf("hello");

I cannot find a way to debug this problem ! Maybe there is a bug when you call some methods ! Note sure that the problem is on my side !

Thanks

[genaganna]

Originally posted by: viewon01 Okay,

In fact my problem come that I got this error :

"Unhandled exception at 0x042b1056 (OCL6C3.tmp.dll) in RSRenderer.exe: 0xC0000005: Access violation reading location 0x3f512c80."

There is a memory access violation somewhere !

When I try to do this in some methods I got a crash:

GB.Intersections[0].Maximum = 0;

but when I do the same in th kernel all is fine.

Also, even when I do this line in the kernel... all is fine... but I got a crash from other methods

printf("hello");

I cannot find a way to debug this problem ! Maybe there is a bug when you call some methods ! Note sure that the problem is on my side !

Thanks

Please past your kernel code which shows crash.

[spectral]

Here it is (Maybe it is still a problem with the GB variable). Notes that I only use the CPU (no GPU).

Thanks

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 { // Collection position and length int Instances_Index; int Instances_Length; int pad1; int pad2; } InstanceList; typedef struct __attribute__ ((packed)) _MovingMatrix4 { // Collection position and length 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 { // Accelerator 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 { /// <summary> /// The bounding box's coordinates /// </summary> float MinX, MinY, MinZ, MaxX, MaxY, MaxZ; } AABB; typedef struct __attribute__ ((packed)) _BoundingVolume { /// <summary> /// The bounding box's coordinates /// </summary> AABB BBox; /// <summary> /// The index of the sibbling node (next node at the same level). /// </summary> int SkipNodeIndex; /// <summary> /// The primitive Id. /// </summary> /// <remarks>If it is a set (Count > -1) it is the first index in the primitives' ids list (_primitivesIds)</remarks> int PrimitiveId; /// <summary> /// The number of primitives, used when it is a set. /// </summary> ushort Count; ushort pad1; int pad2; int pad3; int pad4; } BoundingVolume; typedef struct MITData { 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; __constant GlobalBuffer GB; //-------------------------------------------------------------------------------- // Matrix4_Blend //-------------------------------------------------------------------------------- static Matrix4 Matrix4_Blend(__global Matrix4 * m0, __global Matrix4 * m1, float t) { Matrix4 m; m.M30 = m.M31 = m.M32 = m.M33 = 0.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; } //-------------------------------------------------------------------------------- // MovingMatrix_Sample //-------------------------------------------------------------------------------- static Matrix4 MovingMatrix_Sample(__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)); } //-------------------------------------------------------------------------------- // Ray_Transform //-------------------------------------------------------------------------------- #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) { if (isnan(m->M00)) return *ray; Ray 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); //r._minimum = _minimum; //r._maximum = _maximum; return r; } //-------------------------------------------------------------------------------- // Intersects_Sphere //-------------------------------------------------------------------------------- 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.5 * (b - disc) : -0.5 * (b + disc)); float t0 = q / a; float t1 = c / q; // return sorted array if (t0 > t1) { results[0] = t1; results[1] = t0; } else { results[0] = t0; results[1] = t1; } return 1; } static int Intersects_Sphere(float4 origin, float4 direction, int primitiveId) { // intersect in local space 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); // early rejection if (t[0] >= GB.Intersections[GID].Maximum || t[1] <= 0) /* t[1] <= GB.Intersections[GID].Minimum */ return 0; if (t[0] > 0) /* t[0] > GB.Intersections[GID].Minimum */ GB.Intersections[GID].Maximum = (float)t[0]; else GB.Intersections[GID].Maximum = (float)t[1]; //GB.Intersections[GID].U = 0; //GB.Intersections[GID].V = 0; GB.Intersections[GID].PrimitiveId = 0; GB.Intersections[GID].InstanceId = GB.Intersections[GID].CurrentInstanceId; return 1; } //-------------------------------------------------------------------------------- // Intersects_BoxRay //-------------------------------------------------------------------------------- static int Intersects_BoxRay(struct 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; } //-------------------------------------------------------------------------------- // Intersects_Primitive_TriangleMesh //-------------------------------------------------------------------------------- static int Intersects_Primitive_TriangleMesh(__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; } //-------------------------------------------------------------------------------- // Accelerator_BVH_TreeTraversal_TriangleMesh_Internal //-------------------------------------------------------------------------------- static int Accelerator_BVH_TreeTraversal_TriangleMesh_Internal( __global BoundingVolume * _nodes, __global int * primitivesIds, __global PrimitivesList * primitivesList, float4 origin, float4 direction) { int GID = get_global_id(0); int bvNodeIndex = 0; // 1 = true, 0 = false 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; // End of the tree int stopNodeIndex = _nodes[bvNodeIndex].SkipNodeIndex; int hasIntersection = 0; // We process a triangle mesh here if (primitivesList->Type != 1) return 0; TriangleMesh triangleMesh = GB.TriangleMeshes[primitivesList->Index]; while (bvNodeIndex < stopNodeIndex) { // Do a hit test with the bounding volume. // We use the "closest intersection" to check against the BV intersection. // If the 'closest intersection' < 'bv intersection' then there is no // primitive intersection possible ! float minHit = 0; float maxHit = 0; int hasHit = Intersects_BoxRay(&mitData, origin, direction, &_nodes[bvNodeIndex].BBox, &minHit, &maxHit); if (hasHit && minHit <= GB.Intersections[GID].Maximum) { // It is a leaf -> test the primitives if (_nodes[bvNodeIndex].PrimitiveId > -1) { // Contains a simple primitive if(_nodes[bvNodeIndex].Count < 1) { if (Intersects_Primitive_TriangleMesh(&GB.Indices[triangleMesh.Indices_Index], &GB.Points[triangleMesh.Points_Index], origin, direction, _nodes[bvNodeIndex].PrimitiveId)) { // If shadow ray if (GB.Rays[GID].IsShadowRay) return 1; hasIntersection = 1; } } // Contains a set of primitive else if (_nodes[bvNodeIndex].Count > 0) { //Note : bv variable not defined int startIndex = _nodes[bvNodeIndex].PrimitiveId; int endIndex = startIndex + _nodes[bvNodeIndex].Count - 1; for (int index = startIndex; index <= endIndex; index++) { if (Intersects_Primitive_TriangleMesh(&GB.Indices[triangleMesh.Indices_Index], &GB.Points[triangleMesh.Points_Index], origin, direction, primitivesIds[index])) { // If shadow ray if (GB.Rays[GID].IsShadowRay) return 1; hasIntersection = 1; } } } } // Next node at the same level OR // the next sibbling of the parent. bvNodeIndex++; if ((bvNodeIndex >= stopNodeIndex || bvNodeIndex == _nodes[bvNodeIndex].SkipNodeIndex) && hasIntersection) return 1; } // Continue at the same level else bvNodeIndex = _nodes[bvNodeIndex].SkipNodeIndex; } return hasIntersection; } //-------------------------------------------------------------------------------- // Accelerator_BVH_TreeTraversal_TriangleMesh //-------------------------------------------------------------------------------- static int Accelerator_BVH_TreeTraversal_TriangleMesh(__global Accelerator_BVH * _acceleratorBVHs, float4 origin, float4 direction) { return Accelerator_BVH_TreeTraversal_TriangleMesh_Internal( &GB.AcceleratorBVH_Nodes[_acceleratorBVHs->Nodes_Index], &GB.AcceleratorBVH_PrimitivesIds[_acceleratorBVHs->PrimitivesIds_Index], &GB.PrimitivesList[_acceleratorBVHs->List_Index], origin, direction); } //-------------------------------------------------------------------------------- // Accelerator_NoAcceleration //-------------------------------------------------------------------------------- static int Accelerator_NoAcceleration(__global Accelerator_NA * accelerator, float4 origin, float4 direction) { PrimitivesList primitiveList = GB.PrimitivesList[accelerator->PrimitiveIndex]; // Sphere if (primitiveList.Type == 2) { return Intersects_Sphere(origin, direction, 0); } /* Plane if (primitiveList.Type == 3) { PrimitivesList primitivesList = GB.PrimitivesList[accelerator->PrimitiveIndex]; return Intersects_Plane(&GB.Planes[primitivesList.Index], origin, direction, 0); }*/ return 0; } //-------------------------------------------------------------------------------- // Intersects_Geometry //-------------------------------------------------------------------------------- static int Intersects_Geometry(Geometry geometry, float4 origin, float4 direction) { if (geometry.AcceleratorNA_Index > -1) { // No acceleration return Accelerator_NoAcceleration(&GB.AcceleratorNAs[geometry.AcceleratorNA_Index], origin, direction); } // BVH //if (geometry.AcceleratorBVH_Index > -1) // return Accelerator_BVH_TreeTraversal_TriangleMesh(&GB.AcceleratorBVHs[geometry.AcceleratorBVH_Index], origin, direction); return 0; } //-------------------------------------------------------------------------------- // Intersects_Instance //-------------------------------------------------------------------------------- static int Intersects_Instance(__global Instance * instance, int primitiveId) { int GID = get_global_id(0); GB.Intersections[GID].CurrentInstanceId = primitiveId; // Project the ray to the object space Matrix4 transform = MovingMatrix_Sample(&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(instance->Geometry, origin, direction); } //-------------------------------------------------------------------------------- // Intersects_InstanceList //-------------------------------------------------------------------------------- static int Intersects_InstanceList(__global InstanceList * instanceList, int primitiveId) { if (primitiveId < instanceList->Instances_Length) { __global Instance * instance = &GB.Instances[(instanceList->Instances_Index + primitiveId)]; return Intersects_Instance(instance, primitiveId); } return 0; //if (primitiveId < Instances.Length) // return Instances[primitiveId].Intersect(r, state); //return _lights[primitiveId - Instances.Length].Intersect(r, state); } //-------------------------------------------------------------------------------- // Accelerator_BVH_TreeTraversal_InstanceList_Internal //-------------------------------------------------------------------------------- static int Accelerator_BVH_TreeTraversal_InstanceList_Internal( __global BoundingVolume * _nodes, __global int * primitivesIds, __global InstanceList * instancesList) { int GID = get_global_id(0); int bvNodeIndex = 0; // 1 = true, 0 = false 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; // End of the tree int stopNodeIndex = _nodes[bvNodeIndex].SkipNodeIndex; int hasIntersection = 0; while (bvNodeIndex < stopNodeIndex) { // Do a hit test with the bounding volume. // We use the "closest intersection" to check against the BV intersection. // If the 'closest intersection' < 'bv intersection' then there is no // primitive intersection possible ! float minHit = 0; float maxHit = 0; int hasHit = Intersects_BoxRay(&mitData, GB.Rays[GID].Origin, GB.Rays[GID].Direction, &_nodes[bvNodeIndex].BBox, &minHit, &maxHit); if (hasHit && minHit <= GB.Intersections[GID].Maximum) { // It is a leaf -> test the primitives if (_nodes[bvNodeIndex].PrimitiveId > -1) { // Contains a simple primitive if(_nodes[bvNodeIndex].Count < 1) { // Instance List if(Intersects_InstanceList(instancesList, _nodes[bvNodeIndex].PrimitiveId)) { // If shadow ray if (GB.Rays[GID].IsShadowRay) return 1; //if (hitsCache != null) // hitsCache.AddHit(_nodes[bvNodeIndex].PrimitiveId, ray.Maximum); hasIntersection = 1; } } // Contains a set of primitive else if (_nodes[bvNodeIndex].Count > 0) { //Note : bv variable not defined int startIndex = _nodes[bvNodeIndex].PrimitiveId; int endIndex = startIndex + _nodes[bvNodeIndex].Count - 1; for (int index = startIndex; index <= endIndex; index++) { if(Intersects_InstanceList(instancesList, primitivesIds[index])) { // If shadow ray if (GB.Rays[GID].IsShadowRay) return 1; //if (hitsCache != null) // hitsCache.AddHit(_primitivesIds[index], ray.Maximum); hasIntersection = 1; } } } } // Next node at the same level OR // the next sibbling of the parent. bvNodeIndex++; if ((bvNodeIndex >= stopNodeIndex || bvNodeIndex == _nodes[bvNodeIndex].SkipNodeIndex) && hasIntersection) return 1; } // Continue at the same level else bvNodeIndex = _nodes[bvNodeIndex].SkipNodeIndex; } return hasIntersection; } //-------------------------------------------------------------------------------- // Accelerator_BVH_TreeTraversal_InstanceList //-------------------------------------------------------------------------------- static int Accelerator_BVH_TreeTraversal_InstanceList(__global Accelerator_BVH * _acceleratorBVHs) { return Accelerator_BVH_TreeTraversal_InstanceList_Internal( &GB.AcceleratorBVH_Nodes[_acceleratorBVHs->Nodes_Index], &GB.AcceleratorBVH_PrimitivesIds[_acceleratorBVHs->PrimitivesIds_Index], &GB.InstancesList[_acceleratorBVHs->List_Index]); } //-------------------------------------------------------------------------------- // UploadScene //-------------------------------------------------------------------------------- __kernel void UploadScene( 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) { 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; } //-------------------------------------------------------------------------------- // Trace //-------------------------------------------------------------------------------- __kernel void Trace(__global Ray * _rays, __global IntersectionResult * _intersections) { GB.Rays = _rays; GB.Intersections = _intersections; // Global accelerator ... always a bvh Accelerator_BVH_TreeTraversal_InstanceList(&GB.AcceleratorBVHs[GB.RootAcceleratorIndex]); } // OpenCL : // http://www.khronos.org/opencl/sdk/1.0/docs/man/xhtml/ // http://www.nvidia.com/content/GTC/documents/1077_GTC09.pdf // http://www.prace-project.eu/hpc-training/prace-stream-computing-workshop/opencl-20training.pdf // http://www.touchdreams.net/blog/2009/08/28/common-mistake-in-using-opencl-2-mis-aligned-vector-field/