diff options
| -rw-r--r-- | geometry.py | 20 | ||||
| -rw-r--r-- | src/kernel.cu | 6 | ||||
| -rw-r--r-- | src/materials.h | 35 |
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" |