|
| 1 | +// Capstone scaffold: GPU level-set ray-march renderer. |
| 2 | +// |
| 3 | +// The host plumbing is DONE (build SDF, upload to device, launch, read back, |
| 4 | +// write PPM) and the light direction + ambient level are handed to the kernel |
| 5 | +// as parameters. Your job is the kernel body — the per-pixel ray-march + shade. |
| 6 | +// As shipped it compiles and runs but only writes a placeholder gradient. |
| 7 | +// |
| 8 | +// Build: cmake -S . -B build && cmake --build build |
| 9 | +// Run: ./build/raymarch [ambient] [lx ly lz] (writes sphere.ppm) |
| 10 | +// View: python3 ppm2png.py sphere.ppm sphere.png (then open sphere.png) |
| 11 | +#include <nanovdb/NanoVDB.h> |
| 12 | +#include <nanovdb/tools/CreatePrimitives.h> |
| 13 | +#include <nanovdb/cuda/DeviceBuffer.h> |
| 14 | +#include <nanovdb/cuda/GridHandle.cuh> |
| 15 | +#include <nanovdb/math/Ray.h> |
| 16 | +#include <nanovdb/math/HDDA.h> // nanovdb::math::ZeroCrossing |
| 17 | +#include <nanovdb/math/Stencils.h> // nanovdb::math::GradStencil |
| 18 | + |
| 19 | +#include <cstdio> |
| 20 | +#include <cstdlib> |
| 21 | +#include <cmath> |
| 22 | +#include <vector> |
| 23 | + |
| 24 | +using GridT = nanovdb::NanoGrid<float>; |
| 25 | + |
| 26 | +__global__ void render(const GridT* grid, unsigned char* img, int W, int H, |
| 27 | + nanovdb::Vec3f lightDir, float ambient) |
| 28 | +{ |
| 29 | + const int x = blockIdx.x * blockDim.x + threadIdx.x; |
| 30 | + const int y = blockIdx.y * blockDim.y + threadIdx.y; |
| 31 | + if (x >= W || y >= H) return; |
| 32 | + const int pid = (y * W + x) * 3; |
| 33 | + |
| 34 | + // ===================== YOUR WORK STARTS HERE ===================== |
| 35 | + // Placeholder so it compiles & runs: a horizontal grey gradient. |
| 36 | + unsigned char shade = (unsigned char)(255.f * x / W); |
| 37 | + |
| 38 | + // The aesthetic inputs are given to you as parameters: `lightDir` and |
| 39 | + // `ambient` (set in main, overridable on the command line). You write the |
| 40 | + // ray-march and the shading math. Signatures are in the cheat sheet. |
| 41 | + // |
| 42 | + // 1. Camera ray (WORLD space). Pinhole: an eye point and, per pixel, a |
| 43 | + // direction through the image plane. You decide the math — remember |
| 44 | + // aspect ratio and that image y grows downward. Wrap it in a Ray. |
| 45 | + // |
| 46 | + // 2. The SDF lives in INDEX space, so convert your world ray to index |
| 47 | + // space before you trace it. (Ray has a method for this.) |
| 48 | + // |
| 49 | + // 3. Find where the ray first crosses the surface (the SDF sign change). |
| 50 | + // Module 5 gave you a helper in math/HDDA.h that drives an HDDA and |
| 51 | + // reports the hit voxel, the value, and the ray parameter. |
| 52 | + // |
| 53 | + // 4. On a hit, the surface normal is the (normalized) SDF gradient at the |
| 54 | + // hit voxel. Module 5's gradient stencil computes it. (Watch its |
| 55 | + // template parameter.) |
| 56 | + // |
| 57 | + // 5. Shade: a Lambertian intensity from the normal and the provided |
| 58 | + // `lightDir`, lifted off the floor by `ambient`; write it as greyscale. |
| 59 | + // On a MISS, leave the background. |
| 60 | + // ====================== YOUR WORK ENDS HERE ====================== |
| 61 | + |
| 62 | + img[pid] = shade; img[pid + 1] = shade; img[pid + 2] = shade; |
| 63 | +} |
| 64 | + |
| 65 | +int main(int argc, char** argv) |
| 66 | +{ |
| 67 | + // Shading parameters passed into the kernel. Defaults, overridable on the |
| 68 | + // command line: ./raymarch [ambient] [lx ly lz] |
| 69 | + float ambient = 0.15f; |
| 70 | + nanovdb::Vec3f lightDir(-0.577f, 0.577f, -0.577f); // (-1,1,-1) normalized |
| 71 | + if (argc >= 2) ambient = (float)std::atof(argv[1]); |
| 72 | + if (argc >= 5) lightDir = nanovdb::Vec3f((float)std::atof(argv[2]), |
| 73 | + (float)std::atof(argv[3]), |
| 74 | + (float)std::atof(argv[4])); |
| 75 | + lightDir.normalize(); |
| 76 | + |
| 77 | + // Build a level-set sphere SDF on the host (radius 100, voxel size 1). |
| 78 | + auto handle = nanovdb::tools::createLevelSetSphere<float>( |
| 79 | + /*radius=*/100.0, /*center=*/nanovdb::Vec3d(0.0), |
| 80 | + /*voxelSize=*/1.0, /*halfWidth=*/3.0); |
| 81 | + |
| 82 | + // Move to the device. copy<DeviceBuffer>() fills the HOST side of the dual |
| 83 | + // buffer; deviceUpload() pushes it to the GPU — skip it and deviceGrid() |
| 84 | + // returns nullptr. |
| 85 | + auto devHandle = handle.copy<nanovdb::cuda::DeviceBuffer>(); |
| 86 | + devHandle.deviceUpload(); |
| 87 | + const GridT* dGrid = devHandle.deviceGrid<float>(); |
| 88 | + if (!dGrid) { std::printf("no device grid\n"); return 1; } |
| 89 | + |
| 90 | + const int W = 512, H = 512; |
| 91 | + unsigned char* dImg = nullptr; |
| 92 | + cudaMalloc(&dImg, size_t(W) * H * 3); |
| 93 | + const dim3 block(16, 16), gridDim((W + 15) / 16, (H + 15) / 16); |
| 94 | + render<<<gridDim, block>>>(dGrid, dImg, W, H, lightDir, ambient); |
| 95 | + cudaDeviceSynchronize(); |
| 96 | + if (auto e = cudaGetLastError(); e != cudaSuccess) { |
| 97 | + std::printf("CUDA error: %s\n", cudaGetErrorString(e)); |
| 98 | + return 1; |
| 99 | + } |
| 100 | + |
| 101 | + std::vector<unsigned char> img(size_t(W) * H * 3); |
| 102 | + cudaMemcpy(img.data(), dImg, img.size(), cudaMemcpyDeviceToHost); |
| 103 | + cudaFree(dImg); |
| 104 | + |
| 105 | + FILE* f = std::fopen("sphere.ppm", "wb"); |
| 106 | + std::fprintf(f, "P6\n%d %d\n255\n", W, H); |
| 107 | + std::fwrite(img.data(), 1, img.size(), f); |
| 108 | + std::fclose(f); |
| 109 | + std::printf("wrote sphere.ppm (%dx%d)\n", W, H); |
| 110 | + return 0; |
| 111 | +} |
0 commit comments