CUDA memo (2)

cuda memo


まずCで左側から右側に向かって黒から白にグラデーションしてるPFM画像*1を吐くプログラムを記述。

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

#define WIDTH  (256)
#define HEIGHT (256)

#define SCALE (1.0)
#define ENDIAN (-1)

void writeppm(FILE *output, const size_t size) {
    int x, y;
    float *ppmbuf = malloc(sizeof(float)*size);
    for (x=0;x<WIDTH;++x) {
        for (y=0;y<HEIGHT;++y) {
            ppmbuf[y*WIDTH+x] = (float)x / (float)WIDTH;
        }
    }
    fprintf(output, "Pf\n%d %d\n%f\n", WIDTH, HEIGHT, ENDIAN*SCALE);
    fwrite(ppmbuf, sizeof(float), size, output);
    free(ppmbuf);
}

int main() {
    const size_t ppmsize = WIDTH * HEIGHT;
    FILE *ppm = fopen("white2black.ppm", "w");
    if (ppm) {
        writeppm(ppm, ppmsize);
    }
    return 0;
}

それを適当にCUDAで書き直したもの。

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

#define WIDTH  (256)
#define HEIGHT (256)

#define SCALE (1.0)
#define ENDIAN (-1)

__global__ void white2black(float *ppmbuf) {
    int index = blockIdx.x*blockDim.x + threadIdx.x;
    float gray = (threadIdx.x / (float)WIDTH);
    ppmbuf[index] = gray;
}
void writeppm(FILE *output, const size_t size) {
    float *ppmbuf = (float*)malloc(sizeof(float)*size);
    float *d_ppmbuf = NULL;

    dim3 blockNum(HEIGHT);
    dim3 threadNum(WIDTH);

    cudaMalloc((void**)&d_ppmbuf, sizeof(float)*size);

    if (cudaConfigureCall(blockNum, threadNum)) {
        fprintf(stderr, "%s\n", cudaGetErrorString(cudaGetLastError()));
        exit(EXIT_FAILURE);
    }
    white2black<<<blockNum, threadNum>>>(d_ppmbuf);

    cudaMemcpy(ppmbuf, d_ppmbuf, sizeof(float)*size, cudaMemcpyDeviceToHost);
    cudaFree(d_ppmbuf);

    fprintf(output, "Pf\n%d %d\n%f\n", WIDTH, HEIGHT, ENDIAN*SCALE);
    fwrite(ppmbuf, sizeof(float), size, output);
    free(ppmbuf);
}

int main(int argc, char **argv) {
    const size_t size = WIDTH * HEIGHT;
    FILE *ppm = fopen("white2black.ppm", "w");
    if (ppm) {
        writeppm(ppm, size);
    }
    return 0;
}

特にどうというところはないのだけれども、スレッドが512個までしか作れないということを知らずに書いていて途中しばらくハマった。

*1:このフォーマットの画像を開けるプログラムがImageMagicぐらいしか見当たらない

test