From c8be8496c8a15d0ede8338939a7512109b8e5e46 Mon Sep 17 00:00:00 2001 From: 3gg <3gg@shellblade.net> Date: Wed, 27 Nov 2024 13:41:09 -0800 Subject: Initial commit. --- CMakeLists.txt | 6 +++ hello/CMakeLists.txt | 11 +++++ hello/hello.cu | 59 +++++++++++++++++++++++++ julia/CMakeLists.txt | 11 +++++ julia/julia.cu | 108 ++++++++++++++++++++++++++++++++++++++++++++++ vector_sum/CMakeLists.txt | 11 +++++ vector_sum/main.cu | 62 ++++++++++++++++++++++++++ 7 files changed, 268 insertions(+) create mode 100644 CMakeLists.txt create mode 100644 hello/CMakeLists.txt create mode 100644 hello/hello.cu create mode 100644 julia/CMakeLists.txt create mode 100644 julia/julia.cu create mode 100644 vector_sum/CMakeLists.txt create mode 100644 vector_sum/main.cu diff --git a/CMakeLists.txt b/CMakeLists.txt new file mode 100644 index 0000000..c3ae680 --- /dev/null +++ b/CMakeLists.txt @@ -0,0 +1,6 @@ +cmake_minimum_required(VERSION 3.28) + +add_subdirectory(hello) +add_subdirectory(julia) +add_subdirectory(ptracer) +add_subdirectory(vector_sum) diff --git a/hello/CMakeLists.txt b/hello/CMakeLists.txt new file mode 100644 index 0000000..e4b4acc --- /dev/null +++ b/hello/CMakeLists.txt @@ -0,0 +1,11 @@ +cmake_minimum_required(VERSION 3.28) + +project(cuda_hello LANGUAGES CUDA CXX) + +add_executable(cuda_hello + hello.cu) + +# -Wpedantic causes warnings due to nvcc emitting non-standard (gcc-specific) +# host code. +# https://stackoverflow.com/questions/31000996/warning-when-compiling-cu-with-wpedantic-style-of-line-directive-is-a-gcc-ex +target_compile_options(cuda_hello PRIVATE -Wall -Wextra -Wno-pedantic) diff --git a/hello/hello.cu b/hello/hello.cu new file mode 100644 index 0000000..691b18c --- /dev/null +++ b/hello/hello.cu @@ -0,0 +1,59 @@ +#include + +void logDevices() { + int count; + if (cudaGetDeviceCount(&count) != cudaSuccess) { + printf("No CUDA devices found\n"); + return; + } + + printf("CUDA devices found: %d\n", count); + for (int i = 0; i < count; ++i) { + cudaDeviceProp properties; + if (cudaGetDeviceProperties(&properties, i) == cudaSuccess) { + printf("Device [%d]: %s\n", i, properties.name); + } + } +} + +__global__ void kernel(int* array, int N) { + for (int i = 0; i < N; ++i) { + array[i] = i; + } +} + +int main() { + logDevices(); + + constexpr int N = 100; + + int* host_array = new int[N]; + int* device_array = nullptr; + bool success = false; + + if (cudaMalloc(&device_array, N * sizeof(int)) != cudaSuccess) { + goto cleanup; + } + + kernel<<<1, 1>>>(device_array, N); + + if (cudaMemcpy( + host_array, device_array, N * sizeof(int), cudaMemcpyDeviceToHost) != + cudaSuccess) { + goto cleanup; + } + + for (int i = 0; i < N; ++i) { + printf("%d ", host_array[i]); + } + printf("\n"); + + success = true; + +cleanup: + delete[] host_array; + if (device_array != nullptr) { + cudaFree(device_array); + } + return success ? 0 : 1; +} diff --git a/julia/CMakeLists.txt b/julia/CMakeLists.txt new file mode 100644 index 0000000..e5428fb --- /dev/null +++ b/julia/CMakeLists.txt @@ -0,0 +1,11 @@ +cmake_minimum_required(VERSION 3.28) + +project(cuda_julia LANGUAGES CUDA CXX) + +add_executable(cuda_julia + julia.cu) + +# -Wpedantic causes warnings due to nvcc emitting non-standard (gcc-specific) +# host code. +# https://stackoverflow.com/questions/31000996/warning-when-compiling-cu-with-wpedantic-style-of-line-directive-is-a-gcc-ex +target_compile_options(cuda_julia PRIVATE -Wall -Wextra -Wno-pedantic) diff --git a/julia/julia.cu b/julia/julia.cu new file mode 100644 index 0000000..f3ecb80 --- /dev/null +++ b/julia/julia.cu @@ -0,0 +1,108 @@ +#include +#include +#include + +struct Pixel { + uint8_t r, g, b; +}; + +struct Complex { + float r, i; + + __device__ float norm2() const { return r * r + i * i; } +}; + +__device__ Complex operator*(Complex a, Complex b) { + return Complex{(a.r * b.r) - (a.i * b.i), (a.i * b.r) + (a.r * b.i)}; +} + +__device__ Complex operator+(Complex a, Complex b) { + return Complex{a.r + b.r, a.i + b.i}; +} + +__device__ int julia(int width, int height, int x, int y) { + constexpr float scale = 1.5; + constexpr int N = 200; + + const float jx = scale * (width / 2 - x) / (width / 2); + const float jy = scale * (height / 2 - y) / (height / 2); + + const Complex c{-0.8, 0.156}; + Complex a{jx, jy}; + + for (int i = 0; i < N; ++i) { + a = a * a + c; + if (a.norm2() > 1000) { + return 0; + } + } + return 1; +} + +__global__ void juliaMain(int width, int height, Pixel* image) { + const int x = blockIdx.x; + const int y = blockIdx.y; + + constexpr Pixel background{41, 95, 152}; + constexpr Pixel juliaColour{228, 192, 135}; + + const Pixel pixel = + julia(width, height, x, y) == 1 ? juliaColour : background; + + image[y * width + x] = pixel; +} + +bool write_pbm(const Pixel* image, int width, int height, const char* path) { + const size_t num_pixels = width * height; + + FILE* file = fopen(path, "wb"); + if (!file) { + return false; + } + + fprintf(file, "P6\n%d %d\n255\n", width, height); + if (fwrite(image, sizeof(Pixel), num_pixels, file) != num_pixels) { + fclose(file); + return false; + } + + fclose(file); + return true; +} + +int main(int argc, const char** argv) { + const int width = argc > 1 ? atoi(argv[1]) : 1920; + const int height = argc > 2 ? atoi(argv[2]) : 1080; + + bool success = false; + + const dim3 dim(width, height); + const int image_size_bytes = width * height * sizeof(Pixel); + auto image_host = new Pixel[width * height]; + Pixel* image_dev = nullptr; + + if (cudaMalloc(&image_dev, image_size_bytes) != cudaSuccess) { + goto cleanup; + } + + juliaMain<<>>(width, height, image_dev); + + if (cudaMemcpy( + image_host, image_dev, image_size_bytes, cudaMemcpyDeviceToHost) != + cudaSuccess) { + goto cleanup; + } + + if (!write_pbm(image_host, width, height, "julia.pbm")) { + goto cleanup; + } + + success = true; + +cleanup: + delete[] image_host; + if (image_dev) { + cudaFree(image_dev); + } + return success ? 0 : 1; +} diff --git a/vector_sum/CMakeLists.txt b/vector_sum/CMakeLists.txt new file mode 100644 index 0000000..1eea51b --- /dev/null +++ b/vector_sum/CMakeLists.txt @@ -0,0 +1,11 @@ +cmake_minimum_required(VERSION 3.28) + +project(vector_sum LANGUAGES CUDA CXX) + +add_executable(vector_sum + main.cu) + +# -Wpedantic causes warnings due to nvcc emitting non-standard (gcc-specific) +# host code. +# https://stackoverflow.com/questions/31000996/warning-when-compiling-cu-with-wpedantic-style-of-line-directive-is-a-gcc-ex +target_compile_options(vector_sum PRIVATE -Wall -Wextra -Wno-pedantic) diff --git a/vector_sum/main.cu b/vector_sum/main.cu new file mode 100644 index 0000000..ba2e964 --- /dev/null +++ b/vector_sum/main.cu @@ -0,0 +1,62 @@ +#include + +__global__ void add(int N, int* a, int* b, int* out) { + const int id = blockIdx.x; + out[id] = a[id] + b[id]; +} + +int main() { + constexpr int N = 100; + + bool success = false; + int host_array[N] = {0}; + int* dev_arrays[3] = {nullptr}; + + // Allocate device arrays. + for (int i = 0; i < 3; ++i) { + if (cudaMalloc(&dev_arrays[i], N * sizeof(int)) != cudaSuccess) { + goto cleanup; + } + } + + // Fill the host array with values 0..N-1. + for (int i = 0; i < N; ++i) { + host_array[i] = i; + } + + // Copy the host array to each of the first two device arrays. + for (int i = 0; i < 2; ++i) { + if (cudaMemcpy( + dev_arrays[i], host_array, N * sizeof(int), + cudaMemcpyHostToDevice) != cudaSuccess) { + goto cleanup; + } + } + + // Add the first two arrays. + // N blocks, 1 thread per block. + add<<>>(N, dev_arrays[0], dev_arrays[1], dev_arrays[2]); + + // Copy the result from the third array to the host. + if (cudaMemcpy( + host_array, dev_arrays[2], N * sizeof(int), cudaMemcpyDeviceToHost) != + cudaSuccess) { + goto cleanup; + } + + // Print the result. + for (int i = 0; i < N; ++i) { + printf("%d ", host_array[i]); + } + printf("\n"); + + success = true; + +cleanup: + for (int i = 0; i < 3; ++i) { + if (dev_arrays[i] != nullptr) { + cudaFree(dev_arrays[i]); + } + } + return success ? 0 : 1; +} -- cgit v1.2.3