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.
thanks in advance,
Alex
Its alot of code And I dont know where is the bug. Thanks for your help I will send you the code.
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, &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
+= color * shadowMask; }