Cloo : memory access problem

typedef struct
{
    /// <summary>
    /// The bounding box's coordinates
    /// </summary>
    float MinX, MinY, MinZ, MaxX, MaxY, MaxZ;
 
} AABB;
 
typedef struct
{
    float Maximum;
    int PrimitiveId;
    float U;
    float V;
} IntersectionResult;
 
typedef struct
{
    /// <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;
 
} BoundingVolume;
 
typedef struct MITData
{
    float4 InverseDirection;
 
    int IsXNegative;
    int IsYNegative;
    int IsZNegative;
} MITData;
 
int Intersects_BoxRay(global MITData * mitData, global float4 * origin, global float4 * direction, global AABB* aabb, float * minHit, float * maxHit)
{
	float tmin, tmax, tymin, tymax, tzmin, tzmax;
 
    if (mitData[0].IsXNegative)
    {
        tmin = (aabb[0].MaxX - origin[0].x) * mitData[0].InverseDirection.x;
        tmax = (aabb[0].MinX - origin[0].x) * mitData[0].InverseDirection.x;
    }
    else
    {
        tmin = (aabb[0].MinX - origin[0].x) * mitData[0].InverseDirection.x;
        tmax = (aabb[0].MaxX - origin[0].x) * mitData[0].InverseDirection.x;
    }
 
    if (mitData[0].IsYNegative)
    {
        tymin = (aabb[0].MaxY - origin[0].y) * mitData[0].InverseDirection.y;
        tymax = (aabb[0].MinY - origin[0].y) * mitData[0].InverseDirection.y;
    }
    else
    {
        tymin = (aabb[0].MinY - origin[0].y) * mitData[0].InverseDirection.y;
        tymax = (aabb[0].MaxY - origin[0].y) * mitData[0].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[0].IsZNegative)
    {
        tzmin = (aabb[0].MaxZ - origin[0].z) * mitData[0].InverseDirection.z;
        tzmax = (aabb[0].MinZ - origin[0].z) * mitData[0].InverseDirection.z;
    }
    else
    {
        tzmin = (aabb[0].MinZ - origin[0].z) * mitData[0].InverseDirection.z;
        tzmax = (aabb[0].MaxZ - origin[0].z) * mitData[0].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;
}
 
int Intersects_Primitive(global int * Indices, global float * Points, global float4 * origin, global float4 * direction, int primitiveId, global IntersectionResult * result)
{
	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[0].x,
        Points[i1 + 1] - origin[0].y,
        Points[i1 + 2] - origin[0].z, 0);
 
    float4 n = cross(edge0, edge1);
 
    float v = dot(*direction, n);
    float iv = 1 / v;
 
    float va = dot(n, edge2);
    float t = iv * va;
 
	if (t <= 0 || t >= result->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;
 
	result->Maximum = t;
	result->U = beta;
	result->V = gamma;
	result->PrimitiveId = primitiveId;
 
    return 1;
}
 
kernel void TreeTraversal(
	global MITData * mitData,
    global float4 * origin,
    global float4 * direction,
    float rayMaximum,
    int isShadowRay, // 1 = true, 0 = false
    global float4 * rayOffset,
    global BoundingVolume * _nodes,
    global int * primitivesIds,
    global int * Indices,
    global float * Points,
    global IntersectionResult * result)
{
    int OCLIndex = get_global_id(0);
    int bvNodeIndex = 0;
 
    //MITData mitData = new MITData(ray);
    // End of the tree
    int stopNodeIndex = _nodes[bvNodeIndex].SkipNodeIndex;
    int hasIntersection = 0;
 
    int loopCount = 0;
    while (bvNodeIndex < stopNodeIndex)
    {
        loopCount++;
 
        // Test for the best primitives
        //if (loopCount < 1 /*&& hitsCache != null*/)
        //{
            // int[] bestHits = hitsCache.Primitives;
            // if (bestHits != null)
            // for (int index = 0; index < bestHits.Length; index++)
                    // _primitiveList.IntersectPrimitive(ray, bestHits[index], istate);
        //}
 
        // 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 hasHit = Intersects_BoxRay(mitData, origin, direction, &_nodes[bvNodeIndex].BBox, &minHit, &maxHit);
        if (hasHit && minHit <= rayMaximum)
        {
            // It is a leaf -> test the primitives
            if (_nodes[bvNodeIndex].PrimitiveId > -1)
            {
                origin[0].x = origin[0].x + rayOffset[0].x;
                origin[0].y = origin[0].y + rayOffset[0].y;
                origin[0].z = origin[0].z + rayOffset[0].z;
                //origin = origin + rayOffset;
 
                // Contains a simple primitive
                if(_nodes[bvNodeIndex].Count < 1)
                {
                    if(Intersects_Primitive(Indices, Points, origin, direction, _nodes[bvNodeIndex].PrimitiveId, result))
                    {
                        // If shadow ray
                        if (isShadowRay > 0)
                            return;
                        //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 = 0;//bv.PrimitiveId;
                    int endIndex = startIndex + 0;//bv.Count - 1;
                    for (int index = startIndex; index <= endIndex; index++)
                    if(Intersects_Primitive(Indices, Points, origin, direction, primitivesIds[index], result))
                    {
 
                        // If shadow ray
                        if (isShadowRay > 0)
                            return;
                        //if (hitsCache != null)
                        // hitsCache.AddHit(_primitivesIds[index], ray.Maximum);
                        hasIntersection = 1;
                    }
                }
 
                //Note : Arithmetic on user defined variables
                //origin[0] -= rayOffset
                origin[0].x -= rayOffset[0].x;
                origin[0].y -= rayOffset[0].y;
                origin[0].z -= rayOffset[0].z;
            }
            // Next node at the same level OR
            // the next sibbling of the parent.
            bvNodeIndex++;
            if ((bvNodeIndex >= stopNodeIndex || bvNodeIndex == _nodes[bvNodeIndex].SkipNodeIndex) &&
                hasIntersection)
                return;
 
        }
 
        // Continue at the same level
        else
            bvNodeIndex = _nodes[bvNodeIndex].SkipNodeIndex;
    }
 
    return;
}