got raycasting laid out

.gitignore

@@ -0,0 +1,2 @@
+old_cuda/*
+src/build/*

3d_fractals_cuda.prf

@@ -0,0 +1 @@
+/home/indigo/.unison/3d_fractals_cuda.prf

entity.cuh

@@ -1,31 +0,0 @@
-#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 T> class entity {
-  using T3 = typename vect_t3<T>::vect_t;
-  public:
-    __device__ entity() : pos_(vect_create<T3>(0)), rot_(vect_create<T3>(0)), scale_(vect_create<T3>(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<T3>(0)), scale_(vect_create<T3>(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

makefile

@@ -1,15 +0,0 @@
-LIBS = -lraylib -lGL -lm -lpthread -ldl -lrt -lX11
-$CC = gcc
-INC = -I /opt/cuda/include
-make:
-	nvcc $(LIBS) $(INC) -O0 --debug -c main.cu -o build/main.o
-	nvcc  --device-debug --compile kernel.cu -o build/kernel.o
-	nvcc $(LIBS) -O0 -o build/indigo_worlds build/main.o build/kernel.o
-
-run:
-	build/indigo_worlds
-
-clean:
-	rm -rf build/*
-
-

render_object.cuh

@@ -1,12 +0,0 @@
-#ifndef RENDER_OBJECT_H
-#define RENDER_OBJECT_H
-#include "entity.cuh"
-
-template<class T> class render_object : public entity<T> {
-  using T3 = typename vect_t3<T>::vect_t;
-  using entity<T>::entity;
-  public:
-    virtual __device__ T distance_estimator(T3 point) const = 0;
-};
-
-#endif

scene.h

@@ -1,12 +0,0 @@
-#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 T> class scene {
-    
-
-}
-
-#endif

sphere.cuh

@@ -1,17 +0,0 @@
-#ifndef SPHERE_H
-#define SPHERE_H
-#include "render_object.cuh"
-template<class T> class sphere : public render_object<T> {
-  using render_object<T>::render_object;
-  using T3 = typename vect_t3<T>::vect_t;
-  public:
-    __device__ T distance_estimator(T3 point) const;
-  private:
-    T r_ = 1;
-};
-
-template <class T> __device__ T sphere<T>::distance_estimator(T3 point) const {
-  return length(point) - r_;
-}
-
-#endif

src/.vscode/launch.json

@@ -0,0 +1,19 @@
+{
+    // 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"
+        },
+        {
+            "name": "CUDA C++: Attach",
+            "type": "cuda-gdb",
+            "request": "attach"
+        }
+    ]
+}

src/camera.cuh

@@ -9,38 +9,32 @@
 //template <class T> class scene;
 
 //I am soooo high lol
-template<class T> 
+class camera : public entity {
-class camera : public entity<T> {
-  using T3 = typename vect_t3<T>::vect_t;
-  using T2 = typename vect_t2<T>::vect_t;
   public:
-    __device__ void render();
+    __device__ camera(scene *pscene, const T fov = 1, const vect3 pos = make_vect3(0), const vect3 rot = make_vect3(0)) 
-    __device__ camera(scene<T> *pscene, const T fov = 1, const T3 pos = vect_create<T3>(0), const T3 rot = vect_create<T3>(0)) 
+      : pscene(pscene), fov(fov), entity(pos, rot, make_vect3(0)) {};
-      : pscene(pscene), fov(fov), entity<T>(pos, rot, vect_create<T3>(0)) {};
-    
-    //__device__ ~camera();
   private:
     T fov;
-    T2 size;
+    float2 size;
     int steps = 100;
     T clip_min = .1;
     T clip_max = 100;
-    scene<T> *pscene;
+    //scene *pscene;
 };
 
 /**
 //later we'll make scenes objects, rn im lazy (TODO)
-template <class T> __device__ void camera<T>::render() {
+template <class T> __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<T3>(unnormalized_uv)) / vect_create<T3>(gridDim * blockDim)) - 1;
+  const vect3 uv = ((2 * make_vect3(unnormalized_uv)) / make_vect3(gridDim * blockDim)) - 1;
-  const T3 ray_direction = normalize(vect_create<T3>(uv.x, uv.y, 1));
+  const vect3 ray_direction = normalize(make_vect3(uv.x, uv.y, 1));
   T dist;
   T total_dist = 0;
-  T3 ray;
+  vect3 ray;
   int i;
 
 
@@ -48,7 +42,7 @@ template <class T> __device__ void camera<T>::render() {
 
   T min_dist = clip_max;
 
-  render_object<T> **objs = pscene->get_objs();
+  render_object **objs = pscene->get_objs();
   for(i = 0; i < steps; i++) {
     ray = this->pos_ + (total_dist * ray_direction);
     //gyagh memory lookups

src/common.cuh

@@ -6,16 +6,16 @@
 /**
 
 template <class T> class vect_t2;
-template <class T> class vect_t3;
+template <class T> class vect3;
 template <class T> class vect_t4;
 
 //this feels so hacky... idk why people are so scared of metaprogramming
 template <> class vect_t2<double> { public: using vect_t = double2; };
-template <> class vect_t3<double> { public: using vect_t = double3; };
+template <> class vect3<double> { public: using vect_t = double3; };
 template <> class vect_t4<double> { public: using vect_t = double4; };
 
 template <> class vect_t2<float> { public: using vect_t = float2; };
-template <> class vect_t3<float> { public: using vect_t = float3; };
+template <> class vect3<float> { public: using vect_t = float3; };
 template <> class vect_t4<float> { public: using vect_t = float4; };
 
 
@@ -32,12 +32,22 @@ template <class float3, class X, class Y, class Z> __device__ inline float3 vect
 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 float2 vect2;
-typedef float3 vect_t3;
+typedef float3 vect3;
-typedef float4 vect_t4;
 typedef float T;
 
-#define vect1to3(x) (make_float3(x))
+#define make_vect3(...) (make_float3(__VA_ARGS__))
-#define make_vect(x, y, z) (make_float3(x, y, z))
+
+//TODO move to color.cuh
+typedef float3 Color;
+#define make_color(...) (make_float3(__VA_ARGS__))
+
+//TODO move to ray.cuh
+struct Ray {
+  Color color;
+  vect3 start;
+  vect3 direction; //MUST BE A UNIT VECTOR
+  int bounces;
+};
 
 #endif 

src/entity.cuh

@@ -0,0 +1,30 @@
+#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
+class Entity {
+  public:
+    __device__ Entity() : pos_(make_vect3(0)), rot_(make_vect3(0)), scale_(make_vect3(0)) {};
+    __device__ Entity(const vect3 pos, const vect3 rot, const vect3 scale) : pos_(pos), rot_(rot), scale_(scale) {};
+    __device__ Entity(const float3 pos) : pos_(pos), rot_(make_vect3(0)), scale_(make_vect3(0)) {};
+
+
+    vect3 get_pos() const { return pos_; }
+    vect3 get_rot() const { return rot_; }
+    vect3 get_scale() const { return scale_; }
+
+    __device__ void set_pos(const vect3 pos) { pos_ = pos; }
+    __device__ void set_rot(const vect3 rot) { rot_ = rot; }
+    __device__ void set_scale(const vect3 scale) { scale_ = scale; }
+
+  protected:
+    vect3 pos_;
+    vect3 rot_;
+    vect3 scale_;
+
+};
+#endif

src/kernel.cu

@@ -4,6 +4,7 @@
 #include "scene.cuh"
 
 __global__ void render(uint8_t *image) {
-  scene<float> scene;
+  Scene scene;
+  scene.debug();
   //scene.render(image);
 }

src/main.cu

@@ -1,10 +1,11 @@
+#include <stdint.h>
 #include <stdio.h>
 #include <stdint.h>
 #include <cuda_runtime.h>
 #include <string.h>
+#include <raylib.h>
 
 #include "kernel.cuh"
-#include "raylib.h"
 
 int main() {
   //bluuuuugh i'll figure out occupancy later, this res are easy
@@ -37,9 +38,13 @@ int main() {
   Texture tex = LoadTextureFromImage(image);
 
   while(!WindowShouldClose()) {
+    cudaError_t err;
     //cuda stuff
     cudaMalloc((void **)&image_d, res_x * res_y * sizeof(Color));
     render<<<blockCount, threadCount>>>(image_d); 
+    if((err = cudaGetLastError()) != cudaSuccess) {
+      printf("kernel did not launch! Error: %s\n", cudaGetErrorString(err));
+    }
     cudaDeviceSynchronize();
     cudaMemcpy(texture_data, (void **)image_d, res_x * res_y * sizeof(Color), cudaMemcpyDeviceToHost);
 
@@ -49,6 +54,5 @@ int main() {
     DrawFPS(0, 0);
     EndDrawing();
   }
-
   return 0;
 }

src/makefile

@@ -0,0 +1,25 @@
+LIBS = -lraylib -lGL -lm -lpthread -ldl -lrt -lX11
+$CC = gcc
+INC = -I /opt/cuda/include
+COMPILED_BIN = build/indigo_worlds
+
+CU_SRCFILES := $(wildcard *.cu)
+CU_OBJFILES := $(patsubst %.cu, %.o, $(CU_SRCFILES))
+
+all: $(CU_OBJFILES)
+	nvcc $(LIBS) -o build/indigo_worlds build/*.o
+
+%.o: %.cu
+	nvcc --device-debug -dc -c $< -o build/$@
+
+run: all
+	build/indigo_worlds
+
+clean:
+	rm -rf build/*
+
+
+#make:
+#	nvcc $(LIBS) $(INC) -O0 --debug -c main.cu -o build/main.o
+#	nvcc --device-debug --compile kernel.cu -o build/kernel.o
+#	nvcc $(LIBS) -O0 -o build/indigo_worlds build/main.o build/kernel.o

src/render_object.cuh

@@ -0,0 +1,14 @@
+#ifndef RENDER_OBJECT_H
+#define RENDER_OBJECT_H
+#include "entity.cuh"
+#include "common.cuh" //TODO color
+
+class Render_object : public Entity {
+  using Entity::Entity;
+  public:
+    virtual __device__ T distance_estimator(vect3 point) const = 0;
+    virtual __device__ Color get_color(struct Ray ray) const = 0;
+
+};
+
+#endif

src/scene.cu

@@ -0,0 +1,14 @@
+#include <stdio.h>
+#include "scene.cuh"
+
+__device__ void Scene::debug() {
+  //const uint3 unnormalized_uv = ((blockDim * blockIdx) + threadIdx);
+  //const unsigned int img_index = (unnormalized_uv.x + (unnormalized_uv.y * blockDim.x * gridDim.x)) * 4;
+}
+
+__device__ Color Scene::raycast(struct Ray ray) {
+  if(ray.bounces++ > max_bounces) return make_color(0);
+  const size_t obj_count = sizeof(objs) / sizeof(objs[0]);
+  for(size_t obj_i = 0; obj_i < obj_count; obj_i++); 
+  return make_color(0);
+}

src/scene.cuh

@@ -5,28 +5,27 @@
 #include "sphere.cuh"
 #include "render_object.cuh"
 #include "include/helper_math.h"
+#include <stdint.h>
 
-template <class T> class camera;
+//template <class T> 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 T>
+class Scene {
-class scene {
+  const int max_bounces = 10;
-  using T3 = typename vect_t3<T>::vect_t;
   public:
     //__device__ void render(uint8_t *image) { cam.render(); };
-    __device__ render_object<T> **get_objs() { return objs; }
+    __device__ Render_object **get_objs() { return objs; }
-    __device__ render_object<T> **get_image() { return image; }
+    __device__ void debug();
+    __device__ Color raycast(struct Ray ray);
   private:
-    camera<T> cam = camera<T>();
+    //camera cam = camera();
-    sphere<T> sp1 = sphere<T>(vect_create<T3>(0,0.4,-5));
+    Sphere sp1 = Sphere(make_vect3(0, .4, -5));
-    sphere<T> sp2 = sphere<T>(vect_create<T3>(0,-0.4,-5));
+    Sphere sp2 = Sphere(make_vect3(0, -0.4,-5));
   protected:
     //idk why I need to specify the size... why can't the compiler figure that out?
-    render_object<T> *objs[3] = {&sp1, &sp2, NULL};
+    Render_object *objs[3] = {&sp1, &sp2, NULL};
     uint8_t *image;
 };
 
-#include "camera.cuh"
-
 #endif

src/sphere.cu

@@ -0,0 +1,9 @@
+#include "sphere.cuh"
+
+__device__ T Sphere::distance_estimator(vect3 point) const {
+  return length(point) - r_;
+}
+
+__device__ Color Sphere::get_color(struct Ray ray) const {
+  return make_color(0);
+}

src/sphere.cuh

@@ -0,0 +1,17 @@
+#ifndef SPHERE_H
+#define SPHERE_H
+
+#include "render_object.cuh"
+#include "common.cuh"
+
+class Sphere : public Render_object {
+  using Render_object::Render_object;
+  public:
+   __device__ T distance_estimator(vect3 point) const;
+   __device__ Color get_color(struct Ray ray) const;
+  private:
+    T r_ = 1;
+};
+
+
+#endif