OpenCL crsh

Polar_01 · February 1, 2010, 9:33pm

Hi,

I have write some OpenCL code… but when I call the “Build” method… I got a “memory access violation”…

I’m not sure that it is related to my kernel code , but I put it here :

NB: Is therea way to debug some code on XP today ?

typedef struct
{
float4 Origin;
float4 Direction;
float Maximum;
int IsShadowRay;

int pad1;
int pad2;

} Ray;

typedef struct
{
// Collection position and length
int Instances_Index;
int Instances_Length;

int pad1;
int pad2;

} InstanceList;

typedef struct
{
// Collection position and length
int Transforms_Index;
int Transforms_Length;

int pad1;
int pad2;

} MovingMatrix4;

typedef struct
{
int Type;

int InstanceList_Index;
int TriangleMesh_Index;

int pad1;

} PrimitivesList;

typedef struct
{
// Accelerator
int AcceleratorBVH_Index;

int pad1;
int pad2;
int pad3;

} Geometry;

typedef struct
{
MovingMatrix4 o2w;
MovingMatrix4 w2o;

Geometry Geometry;

} Instance;

typedef struct
{
int Points_Index;
int Points_Length;

int Indices_Index;
int Indices_Length;

} TriangleMesh;

typedef struct
{
float V01;
float V02;
float V03;
float V04;
float V05;
float V06;
float V07;
float V08;
float V09;
float V10;
float V11;
float V12;
float V13;
float V14;
float V15;
float V16;
} Matrix4;

typedef struct
{
int PrimitivesList_Index;

int Nodes_Index;
int Nodes_Length;

int PrimitivesIds_Index;
int PrimitivesIds_Length;

int pad1;
int pad2;
int pad3;

} Accelerator_BVH;

typedef struct
{
///

/// The bounding box’s coordinates
///

float MinX, MinY, MinZ, MaxX, MaxY, MaxZ;
} AABB;

typedef struct
{
float Maximum;
int PrimitiveId;
float U;
float V;

int CurrentInstanceId;
int InstanceId;

int pad1;
int pad2;

} IntersectionResult;

typedef struct
{
///

/// The bounding box’s coordinates
///

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;

} GlobalBuffer;

__constant GlobalBuffer GB;

//--------------------------------------------------------------------------------
// Intersects_BoxRay
//--------------------------------------------------------------------------------

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
//--------------------------------------------------------------------------------

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)
    return 0;

float v2 = dot(i, edge0);
if ((v1 + v2) * v > v * v)
    return 0;

float gamma = iv * v2;
if (gamma < 0)
    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
//--------------------------------------------------------------------------------

int Accelerator_BVH_TreeTraversal_TriangleMesh_Internal(
__global BoundingVolume * _nodes,
__global int * primitivesIds,
__global PrimitivesList * primitivesList)
{
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;

// 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;
    float maxHit;
    int hacrap = Intersects_BoxRay(&mitData, GB.Rays[GID].Origin, GB.Rays[GID].Direction, &_nodes[bvNodeIndex].BBox, &minHit, &maxHit);
    if (hacrap && minHit <= GB.Rays[GID].Maximum)
    {
        // It is a leaf -> test the primitives
        if (_nodes[bvNodeIndex].PrimitiveId > -1)
        {
            // Contains a simple primitive
            if(_nodes[bvNodeIndex].Count < 1)
            {
				TriangleMesh triangleMesh = GB.TriangleMeshes[primitivesList[0].TriangleMesh_Index];
				if (Intersects_Primitive_TriangleMesh(&GB.Indices[triangleMesh.Indices_Index], &GB.Points[triangleMesh.Points_Index], GB.Rays[GID].Origin, GB.Rays[GID].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++)
                {
					TriangleMesh triangleMesh = GB.TriangleMeshes[primitivesList->TriangleMesh_Index];
					if (Intersects_Primitive_TriangleMesh(&GB.Indices[triangleMesh.Indices_Index], &GB.Points[triangleMesh.Points_Index], GB.Rays[GID].Origin, GB.Rays[GID].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
//--------------------------------------------------------------------------------

int Accelerator_BVH_TreeTraversal_TriangleMesh(__global Accelerator_BVH * _acceleratorBVHs)
{
//if (_acceleratorBVHs->Nodes_Index >= 84951)
// return 0;

//if (GB.PrimitivesList[_acceleratorBVHs->PrimitivesList_Index].TriangleMesh_Index < 0)
//	return 0;

return Accelerator_BVH_TreeTraversal_TriangleMesh_Internal(
					&GB.AcceleratorBVH_Nodes[_acceleratorBVHs->Nodes_Index],
					&GB.AcceleratorBVH_PrimitivesIds[_acceleratorBVHs->PrimitivesIds_Index],
					&GB.PrimitivesList[_acceleratorBVHs->PrimitivesList_Index]);

}

//--------------------------------------------------------------------------------
// Intersects_Primitive_Geometry
//--------------------------------------------------------------------------------

int Intersects_Primitive_Geometry(Geometry geometry, float4 origin, float4 direction)
{
if (geometry.AcceleratorBVH_Index < 0)
return 0;
return Accelerator_BVH_TreeTraversal_TriangleMesh( &GB.AcceleratorBVHs[geometry.AcceleratorBVH_Index] );
}

//--------------------------------------------------------------------------------
// Intersects_Instance
//--------------------------------------------------------------------------------

int Intersects_Instance(__global Instance * instance, int primitiveId, float4 origin, float4 direction)
{
int GID = get_global_id(0);
GB.Intersections[GID].CurrentInstanceId = primitiveId;

return Intersects_Primitive_Geometry(instance->Geometry, origin, direction);

		/*
        // Project the ray to the object space
        Matrix4 transform = w2o.Sample(intersection.Time);
        Ray localRay = ray.Transform(ref transform);
        intersection.Current = this;

        // Test the intersection with the geometry
        bool hasIntersection = Geometry.Intersect(localRay, intersection);

        // FIXME: transfer max distance to current ray
        ray.Maximum = localRay.Maximum;

        return hasIntersection;
        */

}

//--------------------------------------------------------------------------------
// Intersects_InstanceList
//--------------------------------------------------------------------------------

int Intersects_InstanceList(__global InstanceList * instanceList, float4 origin, float4 direction, int primitiveId)
{
// 6 2 4 7 3
/if (primitiveId == 7)
{
GB.Intersections[get_global_id(0)].InstanceId = primitiveId;
GB.Intersections[get_global_id(0)].Maximum = 324.940979f;
GB.Intersections[get_global_id(0)].PrimitiveId = primitiveId;
GB.Intersections[get_global_id(0)].U = 0.15913143f;
GB.Intersections[get_global_id(0)].V = 0.7451545616f;
return 1;
}
return 0;/

if (primitiveId < instanceList->Instances_Length)
{	
	__global Instance * instance = &GB.Instances[(instanceList->Instances_Index + primitiveId)];
	return Intersects_Instance(instance, primitiveId, origin, direction);
}

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
//--------------------------------------------------------------------------------

int Accelerator_BVH_TreeTraversal_InstanceList_Internal(
__global BoundingVolume * _nodes,
__global int * primitivesIds,
__global PrimitivesList * primitivesList)
{
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;

// 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;
    float maxHit;
    int hacrap = Intersects_BoxRay(&mitData, GB.Rays[GID].Origin, GB.Rays[GID].Direction, &_nodes[bvNodeIndex].BBox, &minHit, &maxHit);
    if (hacrap && minHit <= GB.Rays[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
				__global InstanceList * instanceList = &GB.InstancesList[primitivesList->InstanceList_Index];               
                if(Intersects_InstanceList(instanceList, GB.Rays[GID].Origin, GB.Rays[GID].Direction, _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++)
                {
					__global InstanceList * instanceList = &GB.InstancesList[primitivesList->InstanceList_Index];
					if(Intersects_InstanceList(instanceList, GB.Rays[GID].Origin, GB.Rays[GID].Direction, 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_NoAcceleration
//--------------------------------------------------------------------------------

void Accelerator_NoAcceleration(__global int * primitivesIds)
{
}

//--------------------------------------------------------------------------------
// Accelerator_BVH_TreeTraversal_InstanceList
//--------------------------------------------------------------------------------

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.PrimitivesList[_acceleratorBVHs->PrimitivesList_Index]);
}

//--------------------------------------------------------------------------------
// Trace
//--------------------------------------------------------------------------------

__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

)
{
GB.Rays = _rays;
GB.Intersections = _intersections;
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;

// Global accelerator ... always a bvh
Accelerator_BVH_TreeTraversal_InstanceList(&GB.AcceleratorBVHs[_rootAcceleratorIndex]);

/*GB.Intersections[get_global_id(0)].InstanceId = 0;
GB.Intersections[get_global_id(0)].Maximum = 100;
GB.Intersections[get_global_id(0)].PrimitiveId = 0;*/

}

zzviola · February 14, 2010, 10:04pm

I have the same problem, did you solve it?

Topic		Replies	Views
Memory errors when writing to local variable in kernel CUDA Programming and Performance	21	4384	January 20, 2017
CL_OUT_OF_RESOURCES CUDA Programming and Performance	6	4041	March 11, 2010
CUDA Project output problem CUDA Programming and Performance	15	1986	July 20, 2011
Unknown 8GB memory getting allocated on GPU Legacy PGI Compilers	12	9663	December 7, 2020
HYDRA_FULL_ERROR=1 in training TAO Toolkit	10	2929	August 29, 2023
Memory Access Error with C++ Parallel Algorithm on NVIDIA HPC SDK, but not with CUDA nvc, nvc++ and nvfortran cuda	11	1103	October 9, 2023
Unspecified Device Error with OpenGL CUDA Programming and Performance	9	6540	October 25, 2011
Optix device code returning ptxas error OptiX	5	1223	June 14, 2022
Failed call to cuInit: CUDA_ERROR_UNKNOWN: unknown error TAO Toolkit	5	9913	October 12, 2021
cudaDeviceSynchronize() takes time CUDA Programming and Performance	11	1990	July 28, 2023

OpenCL crsh

Related topics