Unable to build with Clootils

Help
VO 01
2010-04-15
2012-12-21
  • VO 01
    VO 01
    2010-04-15

    Hi,

    I have 2 problems :

    1) In my application, when I call Program.Build I got the following error :

    Intel(R) Core(TM)2 Duo CPU     E7400  @ 2.80GHz

    Internal error: Compilation failed.

    and a window open, then disappear and display :

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

    I have add the following code at the begining of my application, but I still have the problem :

            public void ApplyFixForAMDSDK()
            {
                if (_fixApplied)
                    return;

                bool runningWin32NT = (Environment.OSVersion.Platform == PlatformID.Win32NT) ? true : false;
                if (runningWin32NT)
                    _fixApplied = AllocConsole();
            }

    2) I have a kernel that I can't build (at least under AMD) with Clootils. You can reach me at viewon01@viewon.tv, I will send you my kernel (a little bit too long for the forum). Also, with NVidia I got some compilations problems, but not the sames !

     
  • nythrix
    nythrix
    2010-04-19

    Internal error: Compilation failed.
    AFAIK this points to an error in the OpenCL source. Try building an empty program. If the error goes away it is definitely caused by faulty code, which makes the compiler go bananas.

    You can reach me at viewon01@viewon.tv, I will send you my kernel (a little bit too long for the forum.
    I've tried mailing you before but the message never passed through.

    Also, with NVidia I got some compilations problems, but not the sames !
    Different compiler outputs are quite normal. If the code carries more bugs you will get very different logs.

    Your code is not ok. Unless something got broken, when copy pasting, there are a couple of bugs in there.
    I'm afraid Clootils Syntax Checker is not up to the job for larger sources. In fact, starting from the next release it will not be available any more. The target of this project is not to provide an advanced IDE. Besides better tools are emerging. Why don't you try out http://www.cmsoft.com.br/index.php?option=com_content&view=article&id=65:opencltemplate-106-released&catid=41:ocltemplate-contents&Itemid=75

     
  • VO 01
    VO 01
    2010-04-20

    Hi,

    I have download this tools and it is just a warning, there is no error… I can't see why you tell that my code is not ok ?

    Have I miss something ?

    NB: For the OpenCL code, I have just avoid a method call, here is the new version with this method call.

    #pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable
    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;
    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,
    &GB.Matrix4,
    (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 = t1;
    results = t0;
    }
    else
    {
    results = t0;
    results = 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;
    if (!Solvers_SolveQuadric(qa, qb, qc, &t))
    return 0;
    int GID = get_global_id(0);
    if (t >= GB.Intersections.Maximum || t <= 0)
    return 0;
    if (t > 0)
    GB.Intersections.Maximum = (float)t;
    else
    GB.Intersections.Maximum = (float)t;
    GB.Intersections.PrimitiveId = 0;
    GB.Intersections.InstanceId = GB.Intersections.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 = MAXFLOAT;
    maxHit = 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 = MAXFLOAT;
    maxHit = MAXFLOAT;
    return 0;
    }
    if (tzmin > tmin)
    tmin = tzmin;
    if (tzmax < tmax)
    tmax = tzmax;
    minHit = tmin;
    maxHit = 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;
    int i2 = 3 * Indices;
    int i3 = 3 * Indices;
    float4 edge0 = (float4)(
    Points - Points,
    Points - Points,
    Points - Points, 0);
    float4 edge1 = (float4)(
    Points - Points,
    Points - Points,
    Points - Points, 0);
    float4 edge2 = (float4)(
    Points - origin.x,
    Points - origin.y,
    Points - 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.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.Maximum = t;
    GB.Intersections.U = beta;
    GB.Intersections.V = gamma;
    GB.Intersections.PrimitiveId = primitiveId;
    GB.Intersections.InstanceId = GB.Intersections.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.SkipNodeIndex;
    int hasIntersection = 0;
    if (primitivesList->Type != 1)
    return 0;
    TriangleMesh triangleMesh = GB.TriangleMeshes;
    while (bvNodeIndex < stopNodeIndex)
    {
    float minHit = 0;
    float maxHit = 0;
    int hasHit = Intersects_BoxRay(GB, &mitData, origin, direction, &_nodes.BBox, &minHit, &maxHit);
    if (hasHit && minHit <= GB.Intersections.Maximum)
    {
    if (_nodes.PrimitiveId > -1)
    {
    if(_nodes.Count < 1)
    {
    if (Intersects_Primitive_TriangleMesh(GB, &GB.Indices, &GB.Points, origin, direction, _nodes.PrimitiveId))
    {
    if (GB.Rays.IsShadowRay)
    return 1;
    hasIntersection = 1;
    }
    }
    else if (_nodes.Count > 0)
    {
    int startIndex = _nodes.PrimitiveId;
    int endIndex = startIndex + _nodes.Count - 1;
    for (int index = startIndex; index <= endIndex; index++)
    {
    if (Intersects_Primitive_TriangleMesh(GB, &GB.Indices, &GB.Points, origin, direction, primitivesIds))
    {
    if (GB.Rays.IsShadowRay)
    return 1;
    hasIntersection = 1;
    }
    }
    }
    }
    bvNodeIndex++;
    if ((bvNodeIndex >= stopNodeIndex || bvNodeIndex == _nodes.SkipNodeIndex) &&
    hasIntersection)
    return 1;
    }
    else
    bvNodeIndex = _nodes.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,
    &GB.AcceleratorBVH_PrimitivesIds,
    &GB.PrimitivesList,
    origin,
    direction);
    }
    static int Accelerator_NoAcceleration(GlobalBuffer GB, __global Accelerator_NA * accelerator, float4 origin, float4 direction)
    {
    PrimitivesList primitiveList = GB.PrimitivesList;
    if (primitiveList.Type == 1)
    {
    TriangleMesh triangleMesh = GB.TriangleMeshes;
    int hasIntersection = 0;
    for(int i = 0; i < triangleMesh.Indices_Length/3; i++)
    hasIntersection &= Intersects_Primitive_TriangleMesh(GB, &GB.Indices, &GB.Points, 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, origin, direction);
    }
    if (geometry.AcceleratorBVH_Index > -1)
    return Accelerator_BVH_TreeTraversal_TriangleMesh(GB, &GB.AcceleratorBVHs, origin, direction);
    return 0;
    }
    static int Intersects_Instance(GlobalBuffer GB, __global Instance * instance, int primitiveId)
    {
    int GID = get_global_id(0);
    GB.Intersections.CurrentInstanceId = primitiveId;
    Matrix4 transform = MovingMatrix_Sample(GB, &instance->w2o, GB.Intersections.Time);
    Ray localRay = Ray_Transform(&transform, &GB.Rays);
    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;
    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.Direction.x;
    mitData.InverseDirection.y = 1.0f / GB.Rays.Direction.y;
    mitData.InverseDirection.z = 1.0f / GB.Rays.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.SkipNodeIndex;
    int hasIntersection = 0;
    while (bvNodeIndex < stopNodeIndex)
    {
    float minHit = 0;
    float maxHit = 0;
    int hasHit = Intersects_BoxRay(GB, &mitData, GB.Rays.Origin, GB.Rays.Direction, &_nodes.BBox, &minHit, &maxHit);
    if (hasHit && minHit <= GB.Intersections.Maximum)
    {
    if (_nodes.PrimitiveId > -1)
    {
    if(_nodes.Count < 1)
    {
    if(Intersects_InstanceList(GB, instancesList, _nodes.PrimitiveId))
    {
    if (GB.Rays.IsShadowRay)
    return 1;
    hasIntersection = 1;
    }
    }
    else if (_nodes.Count > 0)
    {
    int startIndex = _nodes.PrimitiveId;
    int endIndex = startIndex + _nodes.Count - 1;
    for (int index = startIndex; index <= endIndex; index++)
    {
    if(Intersects_InstanceList(GB, instancesList, primitivesIds))
    {
    if (GB.Rays.IsShadowRay)
    return 1;
    hasIntersection = 1;
    }
    }
    }
    }
    bvNodeIndex++;
    if ((bvNodeIndex >= stopNodeIndex || bvNodeIndex == _nodes.SkipNodeIndex) &&
    hasIntersection)
    return 1;
    }
    else
    bvNodeIndex = _nodes.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,
    &GB.AcceleratorBVH_PrimitivesIds,
    &GB.InstancesList);
    }
    __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);
    }

     
  • mux85
    mux85
    2010-04-21

    hi, i have downloaded the tool that you suggested but running OpenCLTemplate.exe no platform is found, i only get "OpenCL ERROR". I am usign the latest developer driver from nvidia. what should i do to make it work? thanks

     
  • mux85
    mux85
    2010-04-21

    resolved. i reinstalled the driver and uninstalled the stream sdk. now everything works

     
  • nythrix
    nythrix
    2010-04-22

    It all got a bit complicated. Builds using nVidia+Clootils and ATI+manually(*) but crashes on nVidia+manually (access violation) and ATI+Clootils(this probably caused by either Clootils or ATI Stream). I can fix the nVidia crash if I comment out the last line of the kernel but I'm unable to tell what exactly is going wrong.
    I'm out of ideas. You can try commenting different pieces of the source to track down the faulty line (if any). However this might caused by a lot of factors including drivers.

    *) Manually means I built it with a short test program in Visual Studio.

     
  • nythrix
    nythrix
    2010-04-22

    Clootils Syntax Checker was pestered by a nasty bug so I've updated the package. The new version (0.7.2) correctly builds your source with both ATI Stream and CUDA!
    Try updating your CUDA installation. I suggest you go for the developer drivers, version 197.13. They look pretty solid since I haven't managed to crash them yet. Not many OpenCL drivers can say that :)