From 093200a449ea38952de52012e324036c106e294b Mon Sep 17 00:00:00 2001 From: Brett Weiland Date: Mon, 27 May 2024 20:56:59 -0500 Subject: boutta switch away from templates --- .gdb_history | 2 ++ .vscode/launch.json | 15 ++++++++++ build/indigo_worlds | Bin 1030064 -> 0 bytes build/kernel.o | Bin 17216 -> 0 bytes build/main.o | Bin 32792 -> 32912 bytes camera.cuh | 81 ++++++++++++++++++++++++++++++++++++++++++++++++++++ common.cuh | 43 ++++++++++++++++++++++++++++ entity.cuh | 31 ++++++++++++++++++++ kernel.cu | 17 +++-------- kernel.cuh | 5 ++++ main.cu | 14 +++------ makefile | 6 ++-- render_object.cuh | 12 ++++++++ scene.cuh | 32 +++++++++++++++++++++ scene.h | 12 ++++++++ sphere.cuh | 17 +++++++++++ 16 files changed, 261 insertions(+), 26 deletions(-) create mode 100644 .vscode/launch.json delete mode 100755 build/indigo_worlds delete mode 100644 build/kernel.o create mode 100644 camera.cuh create mode 100644 common.cuh create mode 100644 entity.cuh create mode 100644 kernel.cuh create mode 100644 render_object.cuh create mode 100644 scene.cuh create mode 100644 scene.h create mode 100644 sphere.cuh diff --git a/.gdb_history b/.gdb_history index 9cf0bf1..8e2a810 100644 --- a/.gdb_history +++ b/.gdb_history @@ -54,3 +54,5 @@ next print image hexdump 0x55555561a140 quit +exit() +exit diff --git a/.vscode/launch.json b/.vscode/launch.json new file mode 100644 index 0000000..f4ee283 --- /dev/null +++ b/.vscode/launch.json @@ -0,0 +1,15 @@ +{ + // Use IntelliSense to learn about possible attributes. + // Hover to view descriptions of existing attributes. + // For more information, visit: https://go.microsoft.com/fwlink/?linkid=830387 + "version": "0.2.0", + "configurations": [ + { + "name": "CUDA C++: Launch", + "type": "cuda-gdb", + "request": "launch", + "program": "${workspaceFolder}/build/indigo_worlds" + } + ] +} + diff --git a/build/indigo_worlds b/build/indigo_worlds deleted file mode 100755 index d99fd92..0000000 Binary files a/build/indigo_worlds and /dev/null differ diff --git a/build/kernel.o b/build/kernel.o deleted file mode 100644 index 549a102..0000000 Binary files a/build/kernel.o and /dev/null differ diff --git a/build/main.o b/build/main.o index 1102ea1..737de98 100644 Binary files a/build/main.o and b/build/main.o differ diff --git a/camera.cuh b/camera.cuh new file mode 100644 index 0000000..61944fe --- /dev/null +++ b/camera.cuh @@ -0,0 +1,81 @@ +#ifndef CAMERA_H +#define CAMERA_H + +#include "entity.cuh" +#include "common.cuh" +#include + + +//template class scene; + +//I am soooo high lol +template +class camera : public entity { + using T3 = typename vect_t3::vect_t; + using T2 = typename vect_t2::vect_t; + public: + __device__ void render(); + __device__ camera(scene *pscene, const T fov = 1, const T3 pos = vect_create(0), const T3 rot = vect_create(0)) + : pscene(pscene), fov(fov), entity(pos, rot, vect_create(0)) {}; + + //__device__ ~camera(); + private: + T fov; + T2 size; + int steps = 100; + T clip_min = .1; + T clip_max = 100; + scene *pscene; +}; + +/** +//later we'll make scenes objects, rn im lazy (TODO) +template __device__ void camera::render() { + //TODO *really* need to clean this up once you get further + //extra dimentions is extra math + //either generisize float3 or stop using this fucking template nonsense + const uint3 unnormalized_uv = ((blockDim * blockIdx) + threadIdx); + const unsigned int img_index = (unnormalized_uv.x + (unnormalized_uv.y * blockDim.x * gridDim.x)) * 4; + const T3 uv = ((2 * vect_create(unnormalized_uv)) / vect_create(gridDim * blockDim)) - 1; + const T3 ray_direction = normalize(vect_create(uv.x, uv.y, 1)); + T dist; + T total_dist = 0; + T3 ray; + int i; + + + //if(img_index == 640) { printf("%f, %f, %f\n", uv.x, uv.y, uv.z); } + + T min_dist = clip_max; + + render_object **objs = pscene->get_objs(); + for(i = 0; i < steps; i++) { + ray = this->pos_ + (total_dist * ray_direction); + //gyagh memory lookups + for(unsigned int oi = 0; objs[oi] != NULL; oi++) { + dist = object.distance_estimator(ray); + } + if((dist < clip_min)) { + //image[img_index] = 0xff; + break; + } + if((dist > clip_max)) { + //image[img_index+2] = 0xff; + break; + } + total_dist += dist; + } + + + + + //image[img_index] = 0x00; + //image[img_index+1] = 0x00; + //image[img_index+2] = p; + //image[img_index+3] = 0xff; + +} +**/ + + +#endif diff --git a/common.cuh b/common.cuh new file mode 100644 index 0000000..d5a9cdf --- /dev/null +++ b/common.cuh @@ -0,0 +1,43 @@ +#ifndef COMMON_H +#define COMMON_H + +#include "include/helper_math.h" + +/** + +template class vect_t2; +template class vect_t3; +template class vect_t4; + +//this feels so hacky... idk why people are so scared of metaprogramming +template <> class vect_t2 { public: using vect_t = double2; }; +template <> class vect_t3 { public: using vect_t = double3; }; +template <> class vect_t4 { public: using vect_t = double4; }; + +template <> class vect_t2 { public: using vect_t = float2; }; +template <> class vect_t3 { public: using vect_t = float3; }; +template <> class vect_t4 { public: using vect_t = float4; }; + + +template __device__ T vect_create(X x); +template __device__ T vect_create(X x, Y y, Z z); + +//I have no fucking clue if this is right, check me later ig +template __device__ inline float3 vect_create(X x) { return make_float3(x); } + +template __device__ inline float3 vect_create(X x, Y y, Z z) { return make_float3(x, y, z); } +**/ + +/** I'm not sure weather float or double percision is nessesary. I was using +templates, but this changes the structure of my entire project in unwanted +ways, so I'm switching over to typedefs. **/ + +typedef float2 vect_t2; +typedef float3 vect_t3; +typedef float4 vect_t4; +typedef float T; + +#define vect1to3(x) (make_float3(x)) +#define make_vect(x, y, z) (make_float3(x, y, z)) + +#endif diff --git a/entity.cuh b/entity.cuh new file mode 100644 index 0000000..737ef8b --- /dev/null +++ b/entity.cuh @@ -0,0 +1,31 @@ +#ifndef ENTITY_H +#define ENTITY_H +#include "common.cuh" + +//we could make a template to allow double percision, but start with float +//idk how nessesary it is yet so I'll go ahead. +//I know I needed it for zoomin far into the mandelbrot ig, so it's not +//out of the question +template class entity { + using T3 = typename vect_t3::vect_t; + public: + __device__ entity() : pos_(vect_create(0)), rot_(vect_create(0)), scale_(vect_create(0)) {}; + __device__ entity(const T3 pos, const T3 rot, const T3 scale) : pos_(pos), rot_(rot), scale_(scale) {}; + __device__ entity(const float3 pos) : pos_(pos), rot_(vect_create(0)), scale_(vect_create(0)) {}; + + + T3 get_pos() const { return pos_; } + T3 get_rot() const { return rot_; } + T3 get_scale() const { return scale_; } + + __device__ void set_pos(const T3 pos) { pos_ = pos; } + __device__ void set_rot(const T3 rot) { rot_ = rot; } + __device__ void set_scale(const T3 scale) { scale_ = scale; } + + protected: + T3 pos_; + T3 rot_; + T3 scale_; + +}; +#endif diff --git a/kernel.cu b/kernel.cu index 6bed13c..c6895dc 100644 --- a/kernel.cu +++ b/kernel.cu @@ -1,18 +1,9 @@ #include #include #include -#include "include/helper_math.h" +#include "scene.cuh" -__global__ void test_image(uint8_t *image) { - int2 unnormalized_coordinates = make_int2(blockDim.x, blockDim.y) * make_int2(blockIdx.x, blockIdx.y) + make_int2(threadIdx.x, threadIdx.y); - int2 img_res = make_int2(blockDim.x, blockDim.y) * make_int2(gridDim.x, gridDim.y); //can move if needed - size_t img_index = ((unnormalized_coordinates.y * img_res.y) + unnormalized_coordinates.x) * 4; - - //vec3 col = 0.5 + 0.5*cos(iTime+uv.xyx+vec3(0,2,4)); - - - image[img_index] = 0x00; - image[img_index+1] = 0x00; - image[img_index+2] = 0xff; - image[img_index+3] = 0xff; +__global__ void render(uint8_t *image) { + scene scene; + //scene.render(image); } diff --git a/kernel.cuh b/kernel.cuh new file mode 100644 index 0000000..2c5256e --- /dev/null +++ b/kernel.cuh @@ -0,0 +1,5 @@ +#ifndef KERNEL_H +#define KERNEL_H +#include +__global__ void render(uint8_t *image); +#endif diff --git a/main.cu b/main.cu index 4225b09..4aa6cbb 100644 --- a/main.cu +++ b/main.cu @@ -3,10 +3,9 @@ #include #include +#include "kernel.cuh" #include "raylib.h" -__global__ void test_image(uint8_t *image); - int main() { //bluuuuugh i'll figure out occupancy later, this res are easy //calculated manually for gtx1060 with 20 SM, 1024 threads/SM @@ -17,7 +16,8 @@ int main() { uint8_t *image_d; Color texture_data[res_x * res_y]; - + SetTargetFPS(10); + //see if GPU is connected (my egpu is finicky) { @@ -36,16 +36,10 @@ int main() { Image image = GenImageColor(res_x, res_y, BLUE); Texture tex = LoadTextureFromImage(image); - - - - - //if(!IsWindowFullscreen()) ToggleFullscreen(); - while(!WindowShouldClose()) { //cuda stuff cudaMalloc((void **)&image_d, res_x * res_y * sizeof(Color)); - test_image<<>>(image_d); + render<<>>(image_d); cudaDeviceSynchronize(); cudaMemcpy(texture_data, (void **)image_d, res_x * res_y * sizeof(Color), cudaMemcpyDeviceToHost); diff --git a/makefile b/makefile index b733c07..7f6bddb 100644 --- a/makefile +++ b/makefile @@ -2,14 +2,14 @@ LIBS = -lraylib -lGL -lm -lpthread -ldl -lrt -lX11 $CC = gcc INC = -I /opt/cuda/include make: - nvcc $(LIBS) $(INC) --debug -c main.cu -o build/main.o + nvcc $(LIBS) $(INC) -O0 --debug -c main.cu -o build/main.o nvcc --device-debug --compile kernel.cu -o build/kernel.o - nvcc $(LIBS) -o build/indigo_worlds build/main.o build/kernel.o + nvcc $(LIBS) -O0 -o build/indigo_worlds build/main.o build/kernel.o run: build/indigo_worlds clean: - rm -r build/* + rm -rf build/* diff --git a/render_object.cuh b/render_object.cuh new file mode 100644 index 0000000..0063d94 --- /dev/null +++ b/render_object.cuh @@ -0,0 +1,12 @@ +#ifndef RENDER_OBJECT_H +#define RENDER_OBJECT_H +#include "entity.cuh" + +template class render_object : public entity { + using T3 = typename vect_t3::vect_t; + using entity::entity; + public: + virtual __device__ T distance_estimator(T3 point) const = 0; +}; + +#endif diff --git a/scene.cuh b/scene.cuh new file mode 100644 index 0000000..31bb99f --- /dev/null +++ b/scene.cuh @@ -0,0 +1,32 @@ +#ifndef SCENE_H +#define SCENE_H + +#include "common.cuh" +#include "sphere.cuh" +#include "render_object.cuh" +#include "include/helper_math.h" + +template class camera; + +//when we get animations with multiple scenes, we'll make this a virtual function +//with array of DE objects and cam +template +class scene { + using T3 = typename vect_t3::vect_t; + public: + //__device__ void render(uint8_t *image) { cam.render(); }; + __device__ render_object **get_objs() { return objs; } + __device__ render_object **get_image() { return image; } + private: + camera cam = camera(); + sphere sp1 = sphere(vect_create(0,0.4,-5)); + sphere sp2 = sphere(vect_create(0,-0.4,-5)); + protected: + //idk why I need to specify the size... why can't the compiler figure that out? + render_object *objs[3] = {&sp1, &sp2, NULL}; + uint8_t *image; +}; + +#include "camera.cuh" + +#endif diff --git a/scene.h b/scene.h new file mode 100644 index 0000000..05ce2af --- /dev/null +++ b/scene.h @@ -0,0 +1,12 @@ +#ifndef SCENE_H +#include SCENE_H +#include "camera.cuh" +#include "render_object.cuh" + +//for now we only neeed once scene, later we'll expand this to just be a virtual template +template class scene { + + +} + +#endif \ No newline at end of file diff --git a/sphere.cuh b/sphere.cuh new file mode 100644 index 0000000..33c77f2 --- /dev/null +++ b/sphere.cuh @@ -0,0 +1,17 @@ +#ifndef SPHERE_H +#define SPHERE_H +#include "render_object.cuh" +template class sphere : public render_object { + using render_object::render_object; + using T3 = typename vect_t3::vect_t; + public: + __device__ T distance_estimator(T3 point) const; + private: + T r_ = 1; +}; + +template __device__ T sphere::distance_estimator(T3 point) const { + return length(point) - r_; +} + +#endif -- cgit v1.2.3