djairoh
/
cuda-based-raytrace
mirror of https://github.com/MartinOpat/cuda-based-raytrace.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

Refactored functionality into separate files

This commit is contained in:
Martin Opat 2024-11-14 17:54:43 +01:00
parent efb58fba30
commit 1974d7bcc0
8 changed files with 241 additions and 40 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

69
.vscode/settings.json vendored Normal file
View File

@ -0,0 +1,69 @@
{
"files.associations": {
"iostream": "cpp",
"iosfwd": "cpp",
"array": "cpp",
"atomic": "cpp",
"bit": "cpp",
"*.tcc": "cpp",
"bitset": "cpp",
"cctype": "cpp",
"charconv": "cpp",
"chrono": "cpp",
"clocale": "cpp",
"cmath": "cpp",
"codecvt": "cpp",
"compare": "cpp",
"concepts": "cpp",
"condition_variable": "cpp",
"cstdarg": "cpp",
"cstddef": "cpp",
"cstdint": "cpp",
"cstdio": "cpp",
"cstdlib": "cpp",
"cstring": "cpp",
"ctime": "cpp",
"cwchar": "cpp",
"cwctype": "cpp",
"deque": "cpp",
"string": "cpp",
"unordered_map": "cpp",
"vector": "cpp",
"exception": "cpp",
"algorithm": "cpp",
"functional": "cpp",
"iterator": "cpp",
"memory": "cpp",
"memory_resource": "cpp",
"numeric": "cpp",
"optional": "cpp",
"random": "cpp",
"ratio": "cpp",
"string_view": "cpp",
"system_error": "cpp",
"tuple": "cpp",
"type_traits": "cpp",
"utility": "cpp",
"fstream": "cpp",
"initializer_list": "cpp",
"iomanip": "cpp",
"istream": "cpp",
"limits": "cpp",
"mutex": "cpp",
"new": "cpp",
"numbers": "cpp",
"ostream": "cpp",
"semaphore": "cpp",
"span": "cpp",
"sstream": "cpp",
"stdexcept": "cpp",
"stop_token": "cpp",
"streambuf": "cpp",
"thread": "cpp",
"cinttypes": "cpp",
"typeinfo": "cpp",
"variant": "cpp",
"format": "cpp",
"__nullptr": "cpp"
}
}

13
img/handler.h Normal file
View File

@ -0,0 +1,13 @@
#pragma once
#include <fstream>
void saveImage(const char* filename, unsigned char* framebuffer, int width, int height) {
std::ofstream imageFile(filename, std::ios::out | std::ios::binary);
imageFile << "P6\n" << width << " " << height << "\n255\n";
for (int i = 0; i < width * height * 3; i++) {
imageFile << framebuffer[i];
}
imageFile.close();
}

9
linalg/linalg.h Normal file
View File

@ -0,0 +1,9 @@
#pragma once
#ifndef LINALG_H
#define LINALG_H
#include "vec.h"
#include "mat.h"
#endif // LINALG_H

1
linalg/mat.h Normal file
View File

@ -0,0 +1 @@
#pragma once

18
linalg/vec.h Normal file
View File

@ -0,0 +1,18 @@
#pragma once
#include <cuda_runtime.h>
#include <cmath>
struct Vec3 { // TODO: Maybe make this into a class
double x, y, z;
__host__ __device__ Vec3() : x(0), y(0), z(0) {}
__host__ __device__ Vec3(double x, double y, double z) : x(x), y(y), z(z) {}
__host__ __device__ Vec3 operator+(const Vec3& b) const { return Vec3(x + b.x, y + b.y, z + b.z); }
__host__ __device__ Vec3 operator-(const Vec3& b) const { return Vec3(x - b.x, y - b.y, z - b.z); }
__host__ __device__ Vec3 operator*(double b) const { return Vec3(x * b, y * b, z * b); }
__host__ __device__ Vec3 operator-() const { return Vec3(-x, -y, -z); }
__host__ __device__ double dot(const Vec3& b) const { return x * b.x + y * b.y + z * b.z; }
__host__ __device__ Vec3 normalize() const { double len = sqrt(x * x + y * y + z * z); return Vec3(x / len, y / len, z / len); }
};

102
main.cu Normal file
View File

@ -0,0 +1,102 @@
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <iostream>
#include <cmath>
#include "linalg/linalg.h"
#include "objs/sphere.h"
#include "img/handler.h"
#define WIDTH 3840
#define HEIGHT 2160
#define SAMPLES_PER_PIXEL 8
__device__ Vec3 phongShading(const Vec3& point, const Vec3& normal, const Vec3& lightDir, const Vec3& viewDir, const Vec3& color) {
double ambientStrength = 0.1;
double diffuseStrength = 0.8;
double specularStrength = 0.5;
int shininess = 64;
Vec3 ambient = color * ambientStrength;
double diff = max(normal.dot(lightDir), 0.0);
Vec3 diffuse = color * (diffuseStrength * diff);
Vec3 reflectDir = (normal * (2.0 * normal.dot(lightDir)) - lightDir).normalize();
double spec = pow(max(viewDir.dot(reflectDir), 0.0), shininess);
Vec3 specular = Vec3(1.0, 1.0, 1.0) * (specularStrength * spec);
return ambient + diffuse + specular;
}
__global__ void renderKernel(unsigned char* framebuffer, Sphere* spheres, int numSpheres, Vec3 lightPos) {
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x >= WIDTH || y >= HEIGHT) return;
int pixelIndex = (y * WIDTH + x) * 3;
Vec3 rayOrigin(0, 0, 0);
Vec3 colCum(0, 0, 0);
double spp = static_cast<double>(SAMPLES_PER_PIXEL);
for (int sample = 0; sample < SAMPLES_PER_PIXEL; sample++) {
double u = (x + (sample / spp) - WIDTH / 2.0) / WIDTH;
double v = (y + (sample / spp) - HEIGHT / 2.0) / HEIGHT;
Vec3 rayDir(u, v, 1.0);
rayDir = rayDir.normalize();
for (int i = 0; i < numSpheres; ++i) {
double t;
if (spheres[i].intersect(rayOrigin, rayDir, t)) {
Vec3 hitPoint = rayOrigin + rayDir * t;
Vec3 normal = (hitPoint - spheres[i].center).normalize();
Vec3 lightDir = (lightPos - hitPoint).normalize();
Vec3 viewDir = -rayDir;
colCum = colCum + phongShading(hitPoint, normal, lightDir, viewDir, spheres[i].color);
}
}
}
// Average color across all samples
Vec3 color = colCum * (1.0 / SAMPLES_PER_PIXEL);
framebuffer[pixelIndex] = static_cast<unsigned char>(fmin(color.x, 1.0) * 255);
framebuffer[pixelIndex + 1] = static_cast<unsigned char>(fmin(color.y, 1.0) * 255);
framebuffer[pixelIndex + 2] = static_cast<unsigned char>(fmin(color.z, 1.0) * 255);
}
int main() {
Sphere spheres[] = {
{ Vec3(0, 0, 5), 1.0, Vec3(1.0, 0.0, 0.0) }, // Red sphere
{ Vec3(-2, 1, 7), 1.0, Vec3(0.0, 1.0, 0.0) }, // Green sphere
{ Vec3(2, -1, 6), 1.0, Vec3(0.0, 0.0, 1.0) } // Blue sphere
};
int numSpheres = sizeof(spheres) / sizeof(Sphere);
Vec3 lightPos(5, 5, 0);
unsigned char* d_framebuffer;
unsigned char* h_framebuffer = new unsigned char[WIDTH * HEIGHT * 3];
Sphere* d_spheres;
cudaMalloc(&d_framebuffer, WIDTH * HEIGHT * 3);
cudaMalloc(&d_spheres, numSpheres * sizeof(Sphere));
cudaMemcpy(d_spheres, spheres, numSpheres * sizeof(Sphere), cudaMemcpyHostToDevice);
dim3 threadsPerBlock(16, 16);
dim3 numBlocks((WIDTH + threadsPerBlock.x - 1) / threadsPerBlock.x,
(HEIGHT + threadsPerBlock.y - 1) / threadsPerBlock.y);
renderKernel<<<numBlocks, threadsPerBlock>>>(d_framebuffer, d_spheres, numSpheres, lightPos);
cudaDeviceSynchronize();
cudaMemcpy(h_framebuffer, d_framebuffer, WIDTH * HEIGHT * 3, cudaMemcpyDeviceToHost);
saveImage("output.ppm", h_framebuffer, WIDTH, HEIGHT);
cudaFree(d_framebuffer);
cudaFree(d_spheres);
delete[] h_framebuffer;
std::cout << "High-resolution image saved as output.ppm" << std::endl;
return 0;
}

23
objs/sphere.h Normal file
View File

@ -0,0 +1,23 @@
#pragma once
#include <cuda_runtime.h>
#include <cmath>
#include "linalg/linalg.h"
struct Sphere {
Vec3 center;
double radius;
Vec3 color;
__device__ bool intersect(const Vec3& rayOrigin, const Vec3& rayDir, double& t) const {
Vec3 oc = rayOrigin - center;
double b = oc.dot(rayDir);
double c = oc.dot(oc) - radius * radius;
double h = b * b - c;
if (h < 0.0) return false;
h = sqrt(h);
t = -b - h;
return true;
}
};

46
tests/test_heavy.cu
View File

@ -1,43 +1,16 @@
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <iostream>
#include <fstream>
#include <cmath>
#include "linalg/linalg.h"
#include "sphere.h"
#include "img/handler.h"
#define WIDTH 3840
#define HEIGHT 2160
#define SAMPLES_PER_PIXEL 8
struct Vec3 {
double x, y, z;
__host__ __device__ Vec3() : x(0), y(0), z(0) {}
__host__ __device__ Vec3(double x, double y, double z) : x(x), y(y), z(z) {}
__host__ __device__ Vec3 operator+(const Vec3& b) const { return Vec3(x + b.x, y + b.y, z + b.z); }
__host__ __device__ Vec3 operator-(const Vec3& b) const { return Vec3(x - b.x, y - b.y, z - b.z); }
__host__ __device__ Vec3 operator*(double b) const { return Vec3(x * b, y * b, z * b); }
__host__ __device__ Vec3 operator-() const { return Vec3(-x, -y, -z); }
__host__ __device__ double dot(const Vec3& b) const { return x * b.x + y * b.y + z * b.z; }
__host__ __device__ Vec3 normalize() const { double len = sqrt(x * x + y * y + z * z); return Vec3(x / len, y / len, z / len); }
};
struct Sphere {
Vec3 center;
double radius;
Vec3 color;
__device__ bool intersect(const Vec3& rayOrigin, const Vec3& rayDir, double& t) const {
Vec3 oc = rayOrigin - center;
double b = oc.dot(rayDir);
double c = oc.dot(oc) - radius * radius;
double h = b * b - c;
if (h < 0.0) return false;
h = sqrt(h);
t = -b - h;
return true;
}
};
__device__ Vec3 phongShading(const Vec3& point, const Vec3& normal, const Vec3& lightDir, const Vec3& viewDir, const Vec3& color) {
double ambientStrength = 0.1;
@ -93,14 +66,7 @@ __global__ void renderKernel(unsigned char* framebuffer, Sphere* spheres, int nu
framebuffer[pixelIndex + 2] = static_cast<unsigned char>(fmin(color.z, 1.0) * 255);
}
void saveImage(const char* filename, unsigned char* framebuffer) {
std::ofstream imageFile(filename, std::ios::out | std::ios::binary);
imageFile << "P6\n" << WIDTH << " " << HEIGHT << "\n255\n";
for (int i = 0; i < WIDTH * HEIGHT * 3; i++) {
imageFile << framebuffer[i];
}
imageFile.close();
}
int main() {
Sphere spheres[] = {
@ -125,7 +91,7 @@ int main() {
cudaDeviceSynchronize();
cudaMemcpy(h_framebuffer, d_framebuffer, WIDTH * HEIGHT * 3, cudaMemcpyDeviceToHost);
saveImage("output.ppm", h_framebuffer);
saveImage("output.ppm", h_framebuffer, WIDTH, HEIGHT);
cudaFree(d_framebuffer);
cudaFree(d_spheres);