4 Replies Latest reply on May 4, 2010 3:15 PM by MicahVillmow

    raytracer artifacs with 2.01 and crashing with 2.1

    abahena

      Hi

      Any idea why we are getting this artifacts and why its crashing on 2.1 version. Please look at the picture below. if we use a i7 as opencl we get more artifacst than using a ati card, also the code doesnt work in Nvidia opencl implementation, with an nvidia card.

      [img]http://www.alchemist3d.com/raytracer.jpg[/img]

      thanks in advance,

      Alex

        • raytracer artifacs with 2.01 and crashing with 2.1
          MicahVillmow
          abahena,
          Without the source code it is hard to determine what is wrong or why it is crashing. If you can simplify it to a test case and post it here I can look into why it is failing. If you cannot do that and don't want to post your whole source code, please send an email to streamdeveloper@amd.com CC: Micah and I'll look into it.
            • raytracer artifacs with 2.01 and crashing with 2.1
              abahena

              Its alot of code And I dont know where is the bug. Thanks for your help I will send you the code.

                • raytracer artifacs with 2.01 and crashing with 2.1
                  abahena

                  Hi

                  I hope you got the code but anyways I think the problem its in raytracekernel:

                   

                  void createRay(const uint x, const uint y, __global RTCamera* camera, __global RTRay* ray) { //TODO: now we use a fixed view direction -z, we should change to camera model float4 dir = (float4)((float)(x)/(camera->screenSize.x-1) - 0.5f, (float)(y)/(camera->screenSize.y-1) - 0.5f, -camera->near, 0.0f); ray->origin = camera->position; ray->direction = normalize(dir); ray->pixel.x = x; ray->pixel.y = y; ray->depth = 0; ray->blendFactor = (float4)(1.0f, 1.0003f, 0.0f, 0.0f); } __kernel void spawnPrimaryRays( __global RTCamera* camera, __global RTRay* newRays) { uint x = get_global_id(0); uint y = get_global_id(1); uint width = get_global_size(0); //Create a primary ray __global RTRay* pRay = &newRays[width * y + x]; createRay(x, y, camera, pRay); } //Intersection between ray and triangle void intersectRayTri(const RTRay* ray, __constant RTPrimitive* primitive, RTIntersection* result) { float4 diff = ray->origin - primitive->vertexA.position; float4 edge1 = primitive->vertexB.position - primitive->vertexA.position; float4 edge2 = primitive->vertexC.position - primitive->vertexA.position; float4 normal = cross(edge1, edge2); float fDdN = dot(ray->direction, normal); float fSign; if(fDdN > FLT_EPSILON) { fSign = 1.0f; } else if(fDdN < -FLT_EPSILON) { fSign = -1.0f; fDdN = -fDdN; } else { //Parallel ray (No intersection) return; } float fDdQxE2 = fSign * dot(ray->direction, cross(diff, edge2)); if(fDdQxE2 >= 0.0f) { float fDdE1xQ = fSign * dot(ray->direction, cross(edge1, diff)); if(fDdE1xQ >= 0.0f) { if(fDdQxE2 + fDdE1xQ <= fDdN) { //line intersects -> check ray float fQdN = -fSign * dot(diff, normal); if(fQdN >= 0.0f) { //ray intersects float fInv = 1.0f/fDdN; result->t = fQdN * fInv; result->b1 = fDdQxE2 * fInv; result->b2 = fDdE1xQ * fInv; result->b0 = 1.0f - result->b1 - result->b2; result->p = primitive; } } } } } //Find the closest intersection (min(t)) void findIntersection(const RTRay* ray, const uint numPrimitives, __constant RTPrimitive* primitives, RTIntersection* result) { result->t = FLT_MAX; result->p = 0; RTIntersection intersection; intersection.p = 0; intersection.t = FLT_MAX; for(int i=0; i<numPrimitives; i++) { intersectRayTri(ray, &primitives[i], &intersection); if(intersection.t < result->t) { *result = intersection; } } } //Get surface normal from intersection (Now support only triangle) void getSurfaceNormal(const RTRay* ray, RTIntersection* intersection, float4* normal) { //Intersection Point float4 p = ray->origin + ray->direction * intersection->t; //Vertex Normals __constant RTPrimitive* tri = intersection->p; *normal = normalize(tri->vertexA.normal * intersection->b0 + tri->vertexB.normal * intersection->b1 + tri->vertexC.normal * intersection->b2); } //Do shading + shadow float4 shade( const float4 position, const float4 normal, const uint idx, __global RTCamera* camera, const uint numPrimitives, __constant RTObject* objects, //List of objects in the scene __constant RTPrimitive* primitives, //TODO: now it's a list, change to kd-tree or BHV-tree __constant RTLight* lights, //List of lights float* shadowMask ) { float4 color = (float4)(0.0f); __constant RTObject* object = objects + idx; float4 lightDir = lights[0].position - position; float lightDistance = length(lightDir); lightDir = lightDir/lightDistance; float NdL = max(0.0f, dot(normal, lightDir)); //Hard shadow RTRay shadowRay; RTIntersection shadowResult; shadowRay.origin = position + lightDir * 0.00001f; shadowRay.direction = lightDir; findIntersection(&shadowRay, numPrimitives, primitives, &shadowResult); __constant RTObject* occluder = objects + shadowResult.p->objectIndex; *shadowMask = shadowResult.t < lightDistance ? 1.0f - occluder->material.opacity: 1.0f; //Specular //float4 E = normalize(camera->position - position); //float4 R = -lightDir - 2.0f * dot(normal, -lightDir) * normal; //float specularI = pow( max(dot(R, E), 0.0f), object->material.specularPow); //color += object->material.specular * specularI * shadowMask; //Diffuse color += object->material.diffuse * NdL; return color; } //Trace rays #pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable __kernel void raytrace( __global RTRay* rays, //Current Rays const uint numObjects, const uint numPrimitives, const uint numLights, __global RTCamera* camera, __constant RTObject* objects, //List of objects in the scene __constant RTPrimitive* primitives, //Primitives TODO: now it's a list, change to kd-tree or BHV-tree __constant RTLight* lights, //List of lights __global RTRayInfo* rayInfo, //number of spawned rays __global RTRay* newRays, //spawned/to be traced rays __global float4* fragments) { //Get current ray int id = (int)get_global_id(0); int n = rayInfo->numRays; if(n <= 0) { return; } RTRay ray; __global const RTRay* pRay = rays + id; ray.origin = pRay->origin; ray.direction = pRay->direction; ray.depth = pRay->depth; ray.pixel = pRay->pixel; ray.blendFactor = pRay->blendFactor; atom_dec(&rayInfo->numRays); //rayInfo->numRays--; //if(id == 0) //{ // TRACE3("Depth: %d Ray: %d id: %d\n", ray.depth, rayInfo->numRays, id); // TRACE3("Ray: %f %f %f\n", ray.origin.x, ray.origin.y, ray.origin.z); // TRACE3("Dir: %f %f %f\n", ray.direction.x, ray.direction.y, ray.direction.z); //} //Find the closest intersection RTIntersection result; findIntersection(&ray, numPrimitives, primitives, &result); uint c = ray.pixel.x + ray.pixel.y * 800; if(result.p == 0) { return; } float4 position; float4 normal; position = ray.origin + ray.direction * result.t; getSurfaceNormal(&ray, &result, &normal); uint idx = result.p->objectIndex; __constant RTObject* object = objects + idx; //shade float shadowMask; float4 color = object->material.opacity * ray.blendFactor.x * shade(position, normal, idx, camera, numPrimitives, objects, primitives, lights, &shadowMask); if(ray.depth < 3) { float4 position = ray.origin + ray.direction * result.t; //spawn reflection ray if(object->material.reflectCoeff > 0.0f && shadowMask > 0.0f) { float4 R = normalize(ray.direction - 2.0f * dot(normal, ray.direction) * normal); //rayInfo->numRays++; //REQUIRE ATOMIC!!!! but link fails - LQ ATI > < atom_inc(&rayInfo->numRays); atom_inc(&rayInfo->newRayIndex); int rayIndex = rayInfo->newRayIndex; newRays[rayIndex].origin = position + (R * 0.000001f); newRays[rayIndex].direction = R; newRays[rayIndex].pixel.x = ray.pixel.x; newRays[rayIndex].pixel.y = ray.pixel.y; newRays[rayIndex].depth = ray.depth + 1; newRays[rayIndex].blendFactor.x = object->material.reflectCoeff; newRays[rayIndex].blendFactor.y = ray.blendFactor.y; } //spawn refraction ray float cosI = -dot(normal, ray.direction); float n = ray.blendFactor.x/object->material.ior; //if(object->material.ior != 1.0003f) //TRACE1("%f\n", n); float ref = 1.0f - n*n*(1.0f - cosI*cosI); if(ref > 0.0f && object->material.opacity < 1.0f) { float4 T = normalize(n*ray.direction + (n * cosI - sqrt(ref)) * normal); //float4 T = ray.direction; //rayInfo->numRays++; //REQUIRE ATOMIC!!!! but link fails - LQ ATI > < atom_inc(&rayInfo->numRays); atom_inc(&rayInfo->newRayIndex); int rayIndex = rayInfo->newRayIndex; newRays[rayIndex].origin = position + (T * 0.000001f); newRays[rayIndex].direction = T; newRays[rayIndex].pixel.x = ray.pixel.x; newRays[rayIndex].pixel.y = ray.pixel.y; newRays[rayIndex].depth = ray.depth + 1; newRays[rayIndex].blendFactor.x = object->material.opacity; newRays[rayIndex].blendFactor.y = 1.0003f;//object->material.ior; } } fragments[c] += color * shadowMask; }

              • raytracer artifacs with 2.01 and crashing with 2.1
                MicahVillmow
                abahena,
                I have it and will look into it when I can. One thing you can try is this in raytracekernel:
                __constant RTObject* objects, //List of objects in the scene
                change all pointers using constant address space to declare the max size they will use.
                __constant RTObject* objects __attribute__((max_constant_size(N))). This is in section 4.5 of the OpenCL Programming guide.