load material/surface index lookup arrays to the gpu and bind them to textures. also, forgot to include the gpu code for material/surface structures

3 files changed, 61 insertions, 0 deletions

diff --git a/geometry.py b/geometry.py
index 0569dcd..938001e 100644
--- a/geometry.py
+++ b/geometry.py
@@ -231,6 +231,26 @@ class Geometry(object):
 
             set_surface(i, surface.absorption_gpu, surface.reflection_diffuse_gpu, surface.reflection_specular_gpu)
 
+        material1_index_tex = module.get_texref('material1_index')
+        material2_index_tex = module.get_texref('material2_index')
+        surface1_index_tex = module.get_texref('surface1_index')
+        surface2_index_tex = module.get_texref('surface2_index')
+
+        material1_index_gpu = cuda.to_device(self.material1_index)
+        material2_index_gpu = cuda.to_device(self.material2_index)
+        surface1_index_gpu = cuda.to_device(self.surface1_index)
+        surface2_index_gpu = cuda.to_device(self.surface2_index)
+
+        material1_index_tex.set_address(material1_index_gpu, self.material1_index.nbytes)
+        material2_index_tex.set_address(material2_index_gpu, self.material2_index.nbytes)
+        surface1_index_tex.set_address(surface1_index_gpu, self.surface1_index.nbytes)
+        surface2_index_tex.set_address(surface2_index_gpu, self.surface2_index.nbytes)
+
+        material1_index_tex.set_format(cuda.array_format.SIGNED_INT32, 1)
+        material2_index_tex.set_format(cuda.array_format.SIGNED_INT32, 1)
+        surface1_index_tex.set_format(cuda.array_format.SIGNED_INT32, 1)
+        surface2_index_tex.set_format(cuda.array_format.SIGNED_INT32, 1)
+
         vertices = np.empty(len(self.mesh.vertices), dtype=gpuarray.vec.float4)
         vertices['x'] = self.mesh.vertices[:,0]
         vertices['y'] = self.mesh.vertices[:,1]
diff --git a/src/kernel.cu b/src/kernel.cu
index c56bbe7..4d85f6e 100644
--- a/src/kernel.cu
+++ b/src/kernel.cu
@@ -14,6 +14,12 @@
 texture<float4, 1, cudaReadModeElementType> vertices;
 texture<uint4, 1, cudaReadModeElementType> triangles;
 
+/* material/surface index lookup for each triangle */
+texture<int, 1, cudaReadModeElementType> material1_index;
+texture<int, 1, cudaReadModeElementType> material2_index;
+texture<int, 1, cudaReadModeElementType> surface1_index;
+texture<int, 1, cudaReadModeElementType> surface2_index;
+
 /* lower/upper bounds for the bounding box associated with each node/leaf */
 texture<float4, 1, cudaReadModeElementType> upper_bounds;
 texture<float4, 1, cudaReadModeElementType> lower_bounds;
diff --git a/src/materials.h b/src/materials.h
new file mode 100644
index 0000000..05aa121
--- /dev/null
+++ b/src/materials.h
@@ -0,0 +1,35 @@
+struct Material
+{
+	float *refractive_index;
+	float *absorption_length;
+	float *scattering_length;
+};
+
+struct Surface
+{
+	float *absorption;
+	float *reflection_diffuse;
+	float *reflection_specular;
+};
+
+__device__ struct Material materials[100];
+__device__ struct Surface surfaces[100];
+
+extern "C"
+{
+
+__global__ void set_material(int material_index, float *refractive_index, float *absorption_length, float *scattering_length)
+{
+	materials[material_index].refractive_index = refractive_index;
+	materials[material_index].absorption_length = absorption_length;
+	materials[material_index].scattering_length = scattering_length;
+}
+
+__global__ void set_surface(int surface_index, float *absorption, float *reflection_diffuse, float *reflection_specular)
+{
+	surfaces[surface_index].absorption = absorption;
+	surfaces[surface_index].reflection_diffuse = reflection_diffuse;
+	surfaces[surface_index].reflection_specular = reflection_specular;
+}
+
+} // extern "c"