A program itt is több részből áll. A hatékonyabb kezelés érdekében van definiálva egy Sphere nevű struktúra a gömbök tárolását, és használatát teszi Egyszerűbbé.

A Sphere struktúrában található egy RGB színkód, ami a gömb színét adja majd, illetve a gömb (X,Y,Z) koordinátái és sugara.

struct Sphere {
    float   r,b,g;
    float   radius;
    float   x,y,z;
    __device__ float hit( float ox, float oy, float *n ) {
        float dx = ox - x;
        float dy = oy - y;
        if (dx*dx + dy*dy < radius*radius) {
            float dz = sqrtf( radius*radius - dx*dx - dy*dy );
            *n = dz / sqrtf( radius * radius );
            return dz + z;
        }
        return -INF;
    }
};

A struktúrához tartozim még egy hit() függvény ami egy sugár és egy gömb találkozását képes detektálni. A függvény bemenetül két normalizálr koordinátát kap ox, és oy néven, és visszadja az ütközés a sugárral való metszés távolságát, és a normálvektor és a sugár által bezárt szög koszinuszát a metszés pontjában. (A programot az Nvidia programozói készítették, jól optimalizált, viszont ettől nagyon bonyolult. Megértése csak azoknak célszerű, akik sugárkövetéssel akarnak foglalkozni.)

A porgramban definiálva van még egy konstans tömb, ami 20 darab gömböt tárol.

#define SPHERES 20

__constant__ Sphere s[SPHERES];

A CPU-n futó fő program nem túl édrekes. Annyit csinál mont a korábbi képeket előállító kódok. Lefoglalja a képtömbböt, és elindítja a kernelt. Illetve ezeken kívül legenerál még 20 db véletlenszerű gömböt, amiket átmásol a konstans memóriába.

    Sphere *temp_s = (Sphere*)malloc( sizeof(Sphere) * SPHERES );
    for (int i=0; i<SPHERES; i++) {
        temp_s[i].r = rnd( 1.0f );
        temp_s[i].g = rnd( 1.0f );
        temp_s[i].b = rnd( 1.0f );
        temp_s[i].x = rnd( 1000.0f ) - 500;
        temp_s[i].y = rnd( 1000.0f ) - 500;
        temp_s[i].z = rnd( 1000.0f ) - 500;
        temp_s[i].radius = rnd( 100.0f ) + 20;
    }
    cudaMemcpyToSymbol( s, temp_s, sizeof(Sphere) * SPHERES);
    free( temp_s );

A kernel - habár bonyolultnak is tánet - nem csinál sok dolgot.

Kiszámítja a szál kooridnátáit, és a hozzájuk  kapcsolódó pixel indexet. Illetve kiszámítja az x-y koordinátákat a kép középpontjához képest is.

    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;

    int offset = x + y * blockDim.x * gridDim.x;
    
    float   ox = (x - N/2);
    float   oy = (y - N/2);

Ezek után egy ciklus végigmegy a gömbökön, és a legközelebbi mettszéspont alapján beállítja a szníeket.

    float   r=0, g=0, b=0;
    float   maxz = -INF;
    for(int i=0; i<SPHERES; i++) {
        float   n;
        float   t = s[i].hit( ox, oy, &n );
        if (t > maxz) {
            float fscale = n;
            r = s[i].r * fscale;
            g = s[i].g * fscale;
            b = s[i].b * fscale;
            maxz = t;
        }
    }

A konstans memória kezelése ezen a ponton egyszerűen annyiből áll, hogy használni kell a tömböt észben tartva, hogy nem lehet bele írni.

/* This software contains source code provided by NVIDIA Corporation.
 * The program is based on the example code ray_const.cu of Chapter 6
 * of the "CUDA by Example" book.
 * 
 * A program az NVIDIA Corporation által készített programkódot tartalmaz.
 * A program a "Cuda by Example" tankönyv 6. fejezetének ray_const.cu 
 * példaprogramjára épül.
 */

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <stdlib.h>

#include <FreeImage.h>

#define N           1024
#define BLOCKDIM    16

#define rnd( x ) (x * rand() / RAND_MAX)
#define INF     2e10f

void writeRGBImageToFile(char* fileName, unsigned char* bytes, int w, int h);

struct Sphere {
    float   r,b,g;
    float   radius;
    float   x,y,z;
    __device__ float hit( float ox, float oy, float *n ) {
        float dx = ox - x;
        float dy = oy - y;
        if (dx*dx + dy*dy < radius*radius) {
            float dz = sqrtf( radius*radius - dx*dx - dy*dy );
            *n = dz / sqrtf( radius * radius );
            return dz + z;
        }
        return -INF;
    }
};
#define SPHERES 20

__constant__ Sphere s[SPHERES];

__global__ void kernel( unsigned char *ptr ) {
    // map from threadIdx/BlockIdx to pixel position
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;

    int offset = x + y * blockDim.x * gridDim.x;
    
    float   ox = (x - N/2);
    float   oy = (y - N/2);

    float   r=0, g=0, b=0;
    float   maxz = -INF;
    for(int i=0; i<SPHERES; i++) {
        float   n;
        float   t = s[i].hit( ox, oy, &n );
        if (t > maxz) {
            float fscale = n;
            r = s[i].r * fscale;
            g = s[i].g * fscale;
            b = s[i].b * fscale;
            maxz = t;
        }
    } 

    ptr[offset*3 + 0] = (int)(r * 255);
    ptr[offset*3 + 1] = (int)(g * 255);
    ptr[offset*3 + 2] = (int)(b * 255);
}

// globals needed by the update routine
struct DataBlock {
    unsigned char   *dev_bitmap;
};

int main( void ) {
    unsigned char* bitmap = new unsigned char[N * N * 3];
    unsigned char* dev_bitmap;

    // allocate memory on the GPU for the output bitmap
    cudaMalloc( (void**)&dev_bitmap, N * N * 3 * sizeof(unsigned char));

    // allocate temp memory, initialize it, copy to constant
    // memory on the GPU, then free our temp memory
    Sphere *temp_s = (Sphere*)malloc( sizeof(Sphere) * SPHERES );
    for (int i=0; i<SPHERES; i++) {
        temp_s[i].r = rnd( 1.0f );
        temp_s[i].g = rnd( 1.0f );
        temp_s[i].b = rnd( 1.0f );
        temp_s[i].x = rnd( 1000.0f ) - 500;
        temp_s[i].y = rnd( 1000.0f ) - 500;
        temp_s[i].z = rnd( 1000.0f ) - 500;
        temp_s[i].radius = rnd( 100.0f ) + 20;
    }
    cudaMemcpyToSymbol( s, temp_s, sizeof(Sphere) * SPHERES);
    free( temp_s );

    // generate a bitmap from our sphere data
    dim3 blockDim = dim3(BLOCKDIM, BLOCKDIM, 1);
    dim3 gridDim = dim3((N + BLOCKDIM - 1) / BLOCKDIM, (N + BLOCKDIM - 1) / BLOCKDIM, 1);

    kernel<<<gridDim, blockDim >>>( dev_bitmap );

    // copy our bitmap back from the GPU for display
    cudaMemcpy( bitmap, dev_bitmap, N * N * 3 * sizeof(unsigned char), cudaMemcpyDeviceToHost );

    writeRGBImageToFile("image.png", bitmap, N, N);

    cudaFree( dev_bitmap );

    return 0;
}


void writeRGBImageToFile(char* fileName, unsigned char* bytes, int w, int h)
{
    FreeImage_Initialise();

    FIBITMAP* image = FreeImage_Allocate(w, h, 24);

    unsigned int pitch = FreeImage_GetPitch(image);
    unsigned int lineWidth = w * 3;

    for (int i = 0; i < h; i++)
    {
        unsigned char* src = bytes + i*lineWidth;
        unsigned char* dst = FreeImage_GetScanLine(image, i);

        memcpy(dst, src, lineWidth);
    }

    FreeImage_Save(FIF_PNG, image, fileName);

    FreeImage_Unload(image);

    FreeImage_DeInitialise();
}