Poor CUDA performance, raytracing

Question

I use CUDA 7.5 to learn the basics of raytracing. So far I've done nothing but constant color no-lighting spheres-only raytracing. But my delta time is already 14 - 16 ms for 800x600 resolution - 1 ray for each pixel.

Some pc specs, if needed: GPU - nvidia titan black, CPU - intel core i7..., enough RAM.

Here is the kernel:

__global__ void runKernel(Camera* cam, SphereData* objs, uint* size, uchar* out_img)
{

vec3 pos = cam->getPos();
rs::Rectangle rec = cam->genNearPlaneRectangle();
vec3 start = rec.center - rec.right - rec.up;
float pixelsizex = 1.0F / cam->getWidth();
float pixelsizey = 1.0F / cam->getHeight();
uint width = cam->getWidth();



uint id = threadIdx.x + blockIdx.x * blockDim.x;
uint i = id % width;
uint j = id / width;

vec3 pixelW = start + rec.right * 2 * (pixelsizex / 2 + pixelsizex*i) + rec.up * 2 * (pixelsizey / 2 + pixelsizey*j);
vec3 dir = (pixelW - pos).normalize();

Ray ray = Ray(pos, dir);


vec3 resColor = Tracer::rayTraceSpheres(ray, objs, *size);

out_img[4 * width * j + 4 * i + 0] = 255 * resColor.x();
out_img[4 * width * j + 4 * i + 1] = 255 * resColor.y();
out_img[4 * width * j + 4 * i + 2] = 255 * resColor.z();
out_img[4 * width * j + 4 * i + 3] = 255;

}

I have a really basic standart setup for OpenGL, I use it to render a fullscreen texture of raytracing results.

Main loop's update and render functions:

void Graphics::update()
{
cameraUpdate();

//println(camera.genLookingRay().dir.toString());

uint objSize = objs.size();
cudaMemcpy(cudaData.d_cam, &camera, sizeof(Camera), cudaMemcpyHostToDevice);
cudaMemcpy(cudaData.d_objs, objs.getArray(), sizeof(SphereData)*objs.size(), cudaMemcpyHostToDevice);
cudaMemcpy(cudaData.d_size, &objSize, sizeof(uint), cudaMemcpyHostToDevice);

uint threadsPerBlock = 128;
uint blocks = ((uint)(width*height) + threadsPerBlock - 1)/threadsPerBlock;

runKernel <<< blocks, threadsPerBlock >>>(cudaData.d_cam, cudaData.d_objs, cudaData.d_size, cudaData.d_out_img);


// Check for any errors launching the kernel
cudaError cudaStatus = cudaGetLastError();
if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
}

cudaStatus = cudaDeviceSynchronize();

cudaMemcpy(image, cudaData.d_out_img, sizeof(uchar)*4*width*height, cudaMemcpyDeviceToHost);

RenderFullScreen::setupTexture(width, height, image, 0);
RenderFullScreen::build();

}


void Graphics::render()
{
glClearColor(0,0,0,1);
glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);


Shader::setCurrent(Shader::SHADER_HUD);
Shader::getCurrent().enable();
Shader::getCurrent().setUniformMat4f("P", camera.getMatOrthographic());
RenderFullScreen::render();
Shader::getCurrent().disable();

glfwSwapBuffers(window);
}

Function "rayTraceSpheres(...)":

__CUDA__ vec3 Tracer::rayTraceSpheres(Ray& ray, SphereData* objs, uint size)
{
    vec3 color(0,0,0);
    float tmin = INFINITY;

    for(size_t i = 0; i < size; i++)
    {
        Sphere obj = Sphere(objs[i]);

        float ct = tmin;
        vec3 cc(0,0,0);
        bool res = obj.hit(ray, ct, cc);

        if(ct < tmin) {
            color = cc;
            tmin = ct;
        }
    }

    return color;
}

__CUDA__ bool Sphere::hit(Ray& ray, float& tmin, vec3& color)
{
    float t;
    vec3 temp = ray.start - center;
    float a = ray.dir * ray.dir;
    float b = temp * 2 * ray.dir;
    float c = temp * temp - radius * radius;
    float disc = b * b - 4 * a * c;

    if(disc < 0){
        return false;
    }
    else
    {
        float e = sqrt(disc);
        float denom = 2 * a;
        t = (-b - e)/denom;
        if(t > 0){
            tmin = t;
            color = this->color;
            return true;
        }
        t = (-b + e)/denom;
        if(t > 0)
        {
            tmin = t;
            color = this->color;
            return true;
        }
    }
    return false;
}

What could be the issue ?

Answer 1

I would guess that the amount of branching you're doing is causing the performance loss. When there's warp divergence, divergent threads need to be thrown out and processed again so be careful about using conditionals and loops in CUDA. Try indexing into a list of function pointers and unraveling loops to avoid branching.

Answer 2

If You are working on Windows, CUDA tend to lag a bit due to WDDM feature. With my own CUDA raytracing implementation on Ubuntu 14.04 I get about 55 FPS on GTX 560 Ti, while on Windows 7 with GTX 970 on board I get only about 16 FPS, with the exact same code in both cases (i.e. a single ray per pixel solution). See this for more info about the issue.