diff options
Diffstat (limited to 'gpu.py')
| -rw-r--r-- | gpu.py | 73 |
1 files changed, 0 insertions, 73 deletions
@@ -1,73 +0,0 @@ -import os -import numpy as np -from pycuda import autoinit -from pycuda.compiler import SourceModule -import pycuda.driver as cuda -from pycuda import gpuarray -import layout - -float3 = gpuarray.vec.float3 -float4 = gpuarray.vec.float4 - -source = open(layout.source + '/kernel.cu').read() - -def make_vector(arr, dtype=float3): - if len(arr.shape) != 2 or arr.shape[-1] != 3: - raise Exception('shape mismatch') - - v = np.empty(arr.shape[0], dtype) - v['x'] = arr[:,0] - v['y'] = arr[:,1] - v['z'] = arr[:,2] - - return v - -class GPU(object): - """ - Object to handle all of the texture allocation/referencing when loading - a geometry onto the GPU. - """ - def __init__(self): - print 'device %s' % autoinit.device.name() - - self.module = SourceModule(source, options=['-I' + layout.source], - no_extern_c=True, cache_dir=False) - - self.get_function = self.module.get_function - - def load_geometry(self, geometry): - """ - Load all the textures from `geometry` onto the GPU and return a list - of texture references. - """ - self.mesh_vec = make_vector(geometry.mesh.reshape(geometry.mesh.shape[0]*3,3), float4) - self.lower_bounds_vec = make_vector(geometry.lower_bounds, float4) - self.upper_bounds_vec = make_vector(geometry.upper_bounds, float4) - self.uchild_map = geometry.child_map.astype(np.uint32) - self.uchild_len = geometry.child_len.astype(np.uint32) - - self.mesh_gpu = cuda.to_device(self.mesh_vec) - self.lower_bounds_gpu = cuda.to_device(self.lower_bounds_vec) - self.upper_bounds_gpu = cuda.to_device(self.upper_bounds_vec) - self.child_map_gpu = cuda.to_device(self.uchild_map) - self.child_len_gpu = cuda.to_device(self.uchild_len) - - self.mesh_tex = self.module.get_texref('mesh') - self.lower_bounds_tex = self.module.get_texref('lower_bounds') - self.upper_bounds_tex = self.module.get_texref('upper_bounds') - self.child_map_tex = self.module.get_texref('child_map_arr') - self.child_len_tex = self.module.get_texref('child_len_arr') - - self.mesh_tex.set_address(self.mesh_gpu, self.mesh_vec.nbytes) - self.lower_bounds_tex.set_address(self.lower_bounds_gpu, self.lower_bounds_vec.nbytes) - self.upper_bounds_tex.set_address(self.upper_bounds_gpu, self.upper_bounds_vec.nbytes) - self.child_map_tex.set_address(self.child_map_gpu, self.uchild_map.nbytes) - self.child_len_tex.set_address(self.child_len_gpu, self.uchild_len.nbytes) - - self.mesh_tex.set_format(cuda.array_format.FLOAT, 4) - self.lower_bounds_tex.set_format(cuda.array_format.FLOAT, 4) - self.upper_bounds_tex.set_format(cuda.array_format.FLOAT, 4) - self.child_map_tex.set_format(cuda.array_format.UNSIGNED_INT32, 1) - self.child_len_tex.set_format(cuda.array_format.UNSIGNED_INT32, 1) - - return [self.mesh_tex, self.lower_bounds_tex, self.upper_bounds_tex, self.child_map_tex, self.child_len_tex] |