CUDAのプログラムをOpenCLに移植(Ray-trace)

CUDA by Exampleには、レイトレースをCUDAで実装する方法が記載されている。

CUDA by Example 汎用GPUプログラミング入門

これを、OpenCLに移植してみよう。

改訂新版 OpenCL入門 1.2対応 マルチコアCPU・GPUのための並列プログラミング

基本的なレイトレースのアルゴリズムをGPUに移植

各ビットについて、円の情報を元に、その円が当該ビットに入っているかどうかを判定する。 これを、ドット毎にCUDAコアへ転送して移植する。

#define INF 2e10f
#include "./raytrace.h"

float hit(__global struct Sphere *s, float ox, float oy, float *n)
{
    float dx = ox - s->x;
    float dy = oy - s->y;

    float radius = s->radius;

    if (dx*dx + dy*dy < radius * radius) {
        float dz = sqrt (radius * radius - dx*dx - dy*dy);
        *n = dz / sqrt (radius * radius);
        return dz + s->z;
    }
    return -INF;
}

__kernel void calc_sphere(__global struct Sphere *s, __global unsigned char *ptr)
{
    int x = get_group_id(0);
    int y = get_local_id(0);
    
    int offset = x + y * IMAGE_WIDTH;

    float ox = (x - IMAGE_WIDTH / 2);
    float oy = (y - IMAGE_WIDTH / 2);

    float r = 0, g = 0, b = 0;
    float maxz = -INF;

    for (int i = 0; i < SPHERES; i++) {
        float n;
        float t = hit(&(s[i]), 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 * 4 + 0] = (int)(r * 255);
    ptr[offset * 4 + 1] = (int)(g * 255);
    ptr[offset * 4 + 2] = (int)(b * 255);
    ptr[offset * 4 + 3] = 255;
}

calc_sphereがドット毎に呼ばれ、Sphere s[i]で定義された情報を探索して色を決定する。

• hit ()

     float dz = sqrt (radius * radius - dx*dx - dy*dy);

sqrtはCUDA OpenCLで定義されているようだ。平方根を計算する。 最初は分からなくてmath.hをインクルードして、sqrtf()を呼ぼうとしたが、コンパイルできなかった。 ググってみると、この方法が正しいようだ。

github.com