toys/sww/source/platform/cuda_render.cu

#include <stdio.h>
#include <time.h>

#include "cuda_render.h"

// limited version of CheckCudaErrors from helper_cuda.h in CUDA examples
#define CheckCudaErrors(val) CheckCuda((val), #val, __FILE__, __LINE__)

void CheckCuda(cudaError_t result, char const *const func, const char *const file, int const line) {
    if (result) {
        printf("CUDA error = %d at %s:%d '%s'\n", (unsigned int)(result), file, line, func);
        // Make sure we call CUDA Device Reset before exiting
        cudaDeviceReset();
        exit(111);
    }
}

__global__ void Render(float *fb, int max_x, int max_y) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    int j = threadIdx.y + blockIdx.y * blockDim.y;

    if ((i >= max_x) || (j >= max_y)) return;

    int pixel_index = j * max_x * 3 + i * 3;
    fb[pixel_index + 0] = float(i) / max_x;
    fb[pixel_index + 1] = float(j) / max_y;
    fb[pixel_index + 2] = 0.2;
}

void CudaRender_Create(int w, int h, int tx, int ty) {
    int nx = 1200;
    int ny = 600;
    int tx = 8;
    int ty = 8;

    int num_pixels = w * h;
    size_t fb_size = num_pixels * sizeof(float) * 4;

    // allocate FB
    vec3 *fb;
    CheckCudaErrors(cudaMallocManaged((void **)&fb, fb_size));

    clock_t start, stop;
    start = clock();
    // Render our buffer
    dim3 blocks(w / tx + 1,h / ty + 1);
    dim3 threads(tx, ty);

    Render<<<blocks, threads>>>(fb, w, h);
    CheckCudaErrors(cudaGetLastError());
    CheckCudaErrors(cudaDeviceSynchronize());

    stop = clock();
    double timer_seconds = ((double)(stop - start)) / CLOCKS_PER_SEC;
    printf("took %lf seconds.\n", timer_seconds);

    for (int j = h - 1; j >= 0; j--) {
        for (int i = 0; i < w; i++) {
            size_t pixel_index = j*nx + i;
            int ir = int(255.99 * fb[pixel_index].r());
            int ig = int(255.99 * fb[pixel_index].g());
            int ib = int(255.99 * fb[pixel_index].b());
        }
    }

    CheckCudaErrors(cudaFree(fb));
}