commit d0c24ee29112cd96f6385e4326e8bfc7cbdfc888 Author: Rokas Puzonas Date: Tue Nov 28 23:36:11 2023 +0200 initial commit diff --git a/.gitignore b/.gitignore new file mode 100644 index 0000000..378eac2 --- /dev/null +++ b/.gitignore @@ -0,0 +1 @@ +build diff --git a/CMakeLists.txt b/CMakeLists.txt new file mode 100644 index 0000000..5229a64 --- /dev/null +++ b/CMakeLists.txt @@ -0,0 +1,8 @@ +cmake_minimum_required(VERSION 3.12) +project(raytracing-in-a-weekend) + +set(CMAKE_CXX_STANDARD 14) +enable_language(CUDA) + +add_executable(main src/main.cu) +set_target_properties(main PROPERTIES CUDA_ARCHITECTURES all-major) diff --git a/run.sh b/run.sh new file mode 100755 index 0000000..c200fdb --- /dev/null +++ b/run.sh @@ -0,0 +1,6 @@ +#!/bin/sh +mkdir -p build +cd build +cmake .. +make +./main diff --git a/src/main.cu b/src/main.cu new file mode 100644 index 0000000..c232729 --- /dev/null +++ b/src/main.cu @@ -0,0 +1,82 @@ +#include +#include +#include +#include + +#define checkCudaErrors(val) check_cuda((val), #val, __FILE__, __LINE__) + +void check_cuda(cudaError_t result, const char *func, const char *file, int line) { + if (result) { + std::cout << "CUDA error: " << cudaGetErrorString(result) << " (error code " << static_cast(result) << ")"; + std::cout << " at " << file << ":" << line << " '" << func << "' \n"; + cudaDeviceReset(); + exit(-1); + } +} + +__global__ void render(float *fb, int max_x, int max_y) { + int x = threadIdx.x + blockIdx.x * blockDim.x; + int y = threadIdx.y + blockIdx.y * blockDim.y; + if ((x >= max_x) || (y >= max_y)) return; + int pixel_idx = 3*(y*max_x + x); + + fb[pixel_idx + 0] = float(x) / max_x; + fb[pixel_idx + 1] = float(y) / max_y; + fb[pixel_idx + 2] = 0.2; +} + +int main() { + const char *image_filename = "image.ppm"; + int nx = 1200; + int ny = 600; + int tx = 8; + int ty = 8; + + int num_pixels = nx*ny; + size_t fb_size = 3*num_pixels*sizeof(float); + + float *fb; + checkCudaErrors(cudaMallocManaged(&fb, fb_size)); + + clock_t start = clock(); + { + dim3 blocks(nx/tx+1, ny/ty+1); + dim3 threads(tx, ty); + + render<<>>(fb, nx, ny); + checkCudaErrors(cudaGetLastError()); + checkCudaErrors(cudaDeviceSynchronize()); + } + clock_t stop = clock(); + double timer_seconds = ((double)(stop - start)) / CLOCKS_PER_SEC; + std::cout << "took " << timer_seconds << " seconds.\n"; + + FILE *f = fopen(image_filename, "w"); + assert(f); + + char header[128] = { 0 }; + int header_size = snprintf(header, sizeof(header), "P3\n%d %d\n255\n", nx, ny); + fwrite(header, header_size, 1, f); + + for (int y=ny-1; y >= 0; y--) { + for (int x=0; x < nx; x++) { + size_t pixel_idx = 3*(y*nx + x); + float r = fb[pixel_idx + 0]; + float g = fb[pixel_idx + 1]; + float b = fb[pixel_idx + 2]; + int ir = int(255.99*r); + int ig = int(255.99*g); + int ib = int(255.99*b); + + char pixel_str[128] = { 0 }; + int str_size = snprintf(pixel_str, sizeof(pixel_str), "%d %d %d\n", ir, ig, ib); + fwrite(pixel_str, str_size, 1, f); + } + } + + fclose(f); + + checkCudaErrors(cudaFree(fb)); + + return 0; +}