diff options
| -rwxr-xr-x | benchmark.py | 12 | ||||
| -rwxr-xr-x | camera.py | 1 | ||||
| -rw-r--r-- | gpu.py | 12 | ||||
| -rw-r--r-- | solids/pmts.py | 2 | ||||
| -rw-r--r-- | src/render.cu | 14 |
5 files changed, 19 insertions, 22 deletions
diff --git a/benchmark.py b/benchmark.py index 9f90c7b..5a344bf 100755 --- a/benchmark.py +++ b/benchmark.py @@ -152,23 +152,23 @@ if __name__ == '__main__': from chroma import detectors import gc - lbne = detectors.build_lbne_200kton() + lbne = detectors.lbne() lbne.build(bits=11) - gpu_instance = gpu.GPU() + context = gpu.create_cuda_context() gpu_geometry = gpu.GPUGeometry(lbne) - gpu_instance.print_mem_info() print '%s ray intersections/sec.' % \ tools.ufloat_to_str(intersect(gpu_geometry)) + # run garbage collection since there is a reference loop + # in the GPUArray class. gc.collect() - gpu_instance.print_mem_info() print '%s photons loaded/sec.' % tools.ufloat_to_str(load_photons()) gc.collect() - gpu_instance.print_mem_info() print '%s photons propagated/sec.' % \ tools.ufloat_to_str(propagate(gpu_geometry)) gc.collect() - gpu_instance.print_mem_info() print '%s 100 MeV events histogrammed/s' % \ tools.ufloat_to_str(pdf(gpu_geometry, max(lbne.pmtids))) + + context.pop() @@ -540,6 +540,7 @@ class Camera(multiprocessing.Process): elif event.key == K_g: self.green_magenta = not self.green_magenta + self.update() elif event.key == K_F12: self.screenshot() @@ -16,7 +16,6 @@ from pycuda import gpuarray as ga import chroma.src from chroma.tools import timeit, profile_if_possible from chroma.geometry import standard_wavelengths -from chroma.color import map_to_color from chroma import event cuda.init() @@ -27,7 +26,7 @@ cuda_options = ('--use_fast_math',)#, '--ptxas-options=-v'] @pycuda.tools.context_dependent_memoize def get_cu_module(name, options=None, include_source_directory=True): """Returns a pycuda.compiler.SourceModule object from a CUDA source file - located in the chroma src directory at src/[name].cu.""" + located in the chroma src directory at src/[name].""" if options is None: options = [] elif isinstance(options, tuple): @@ -46,7 +45,10 @@ def get_cu_module(name, options=None, include_source_directory=True): return pycuda.compiler.SourceModule(source, options=options, no_extern_c=True) +@pycuda.tools.memoize def get_cu_source(name): + """Get the source code for a CUDA source file located in the chroma src + directory at src/[name].""" srcdir = os.path.dirname(os.path.abspath(chroma.src.__file__)) with open('%s/%s' % (srcdir, name)) as f: source = f.read() @@ -166,8 +168,6 @@ class GPUPhotons(object): self.last_hit_triangles[:nphotons].set(photons.last_hit_triangles.astype(np.int32)) self.flags[:nphotons].set(photons.flags.astype(np.uint32)) - #cuda_options = ('--use_fast_math', '-w')#, '--ptxas-options=-v'] - module = get_cu_module('propagate.cu', options=cuda_options) self.gpu_funcs = GPUFuncs(module) @@ -775,8 +775,8 @@ class GPUPDF(object): return hitcount, pdf_value, pdf_value * pdf_frac_uncert -def create_context(device_id=None): - """Initialize and return a GPU context on the specified device. +def create_cuda_context(device_id=None): + """Initialize and return a CUDA context on the specified device. If device_id is None, the default device is used.""" try: cuda.mem_get_info() diff --git a/solids/pmts.py b/solids/pmts.py index 4c8d288..43aed3f 100644 --- a/solids/pmts.py +++ b/solids/pmts.py @@ -38,7 +38,7 @@ def build_pmt_shell(filename, outer_material=water, nsteps=16): # convert mm -> m profile /= 1000.0 - return Solid(rotate_extrude(profile[:,0], profile[:,1], nsteps), glass, outer_material, color=0x99ffffff) + return Solid(rotate_extrude(profile[:,0], profile[:,1], nsteps), glass, outer_material, color=0xeeffffff) def build_pmt(filename, glass_thickness, outer_material=water, nsteps=16): profile = read_csv(filename) diff --git a/src/render.cu b/src/render.cu index c1e8ea1..d9ce1b1 100644 --- a/src/render.cu +++ b/src/render.cu @@ -142,11 +142,7 @@ render(int nthreads, float3 *_origin, float3 *_direction, Geometry *geometry, float fg = 0.0f; float fb = 0.0f; for (i=0; i < n; i++) { - float alpha; - if (i < alpha_depth-1) - alpha = color_a[i].w; - else - alpha = 1.0; + float alpha = color_a[i].w; fr += scale*color_a[i].x*alpha; fg += scale*color_a[i].y*alpha; @@ -158,10 +154,10 @@ render(int nthreads, float3 *_origin, float3 *_direction, Geometry *geometry, if (n < alpha_depth) a = floorf(255*(1.0f-scale)); else - a = 255; - unsigned int r = floorf(fr); - unsigned int g = floorf(fg); - unsigned int b = floorf(fb); + a = 255; + unsigned int r = floorf(fr/(1.0f-scale)); + unsigned int g = floorf(fg/(1.0f-scale)); + unsigned int b = floorf(fb/(1.0f-scale)); pixels[id] = a << 24 | r << 16 | g << 8 | b; } |