From cceb75396fe11773f0a8c947be6b158baf10e6fe Mon Sep 17 00:00:00 2001 From: Steven Robertson Date: Mon, 30 Aug 2010 14:45:44 -0400 Subject: [PATCH] Before I rip out tempita and start a DSL --- cuburnlib/cuda.py | 3 +-- cuburnlib/device_code.py | 49 ++++++++++++++++++++++++++++++++-------- main.py | 38 ++++++++++--------------------- 3 files changed, 53 insertions(+), 37 deletions(-) diff --git a/cuburnlib/cuda.py b/cuburnlib/cuda.py index dd1f994..c094c42 100644 --- a/cuburnlib/cuda.py +++ b/cuburnlib/cuda.py @@ -35,7 +35,6 @@ class LaunchContext(object): def __init__(self, entries, block=(1,1,1), grid=(1,1), seed=None, tests=False): self.devinfo = pycuda.tools.DeviceData() - self.stream = cuda.Stream() self.entry_types = entries self.block, self.grid, self.build_tests = block, grid, tests self.rand = np.random.mtrand.RandomState(seed) @@ -73,7 +72,7 @@ class LaunchContext(object): inst = self.ptx.instances[test_type] print "Running test: %s... " % inst.name try: - self.stream.synchronize() + cuda.Context.synchronize() if inst.call(self): print "Test %s passed." % inst.name else: diff --git a/cuburnlib/device_code.py b/cuburnlib/device_code.py index a6e714d..321f0fc 100644 --- a/cuburnlib/device_code.py +++ b/cuburnlib/device_code.py @@ -8,6 +8,7 @@ from cuburnlib.ptx import PTXFragment, PTXEntryPoint, PTXTest class MWCRNG(PTXFragment): def __init__(self): + self.threads_ready = 0 if not os.path.isfile('primes.bin'): raise EnvironmentError('primes.bin not found') @@ -57,24 +58,27 @@ class MWCRNG(PTXFragment): """ def set_up(self, ctx): + if self.threads_ready >= ctx.threads: + return # Load raw big-endian u32 multipliers from primes.bin. with open('primes.bin') as primefp: dt = np.dtype(np.uint32).newbyteorder('B') mults = np.frombuffer(primefp.read(), dtype=dt) + stream = cuda.Stream() # Randomness in choosing multipliers is good, but larger multipliers # have longer periods, which is also good. This is a compromise. - # TODO: fix mutability, enable shuffle here - # TODO: prevent period-1 random generators - #ctx.rand.shuffle(mults[:ctx.threads*4]) + mults = np.array(mults[:ctx.threads*4]) + ctx.rand.shuffle(mults) # Copy multipliers and seeds to the device multdp, multl = ctx.mod.get_global('mwc_rng_mults') - # TODO: get async to work - #cuda.memcpy_htod_async(multdp, mults.tostring()[:multl], ctx.stream) - cuda.memcpy_htod(multdp, mults.tostring()[:multl]) + cuda.memcpy_htod_async(multdp, mults.tostring()[:multl]) + # Intentionally excludes both 0 and (2^32-1), as they can lead to + # degenerate sequences of period 0 + states = np.array(ctx.rand.randint(1, 0xffffffff, size=2*ctx.threads), + dtype=np.uint32) statedp, statel = ctx.mod.get_global('mwc_rng_state') - #cuda.memcpy_htod_async(statedp, ctx.rand.bytes(statel), ctx.stream) - - cuda.memcpy_htod(statedp, ctx.rand.bytes(statel)) + cuda.memcpy_htod_async(statedp, states.tostring()) + self.threads_ready = ctx.threads def tests(self, ctx): return [MWCRNGTest] @@ -151,4 +155,31 @@ loopstart: return False return True +class CameraCoordTransform(PTXFragment): + # This is here until I get the device stream packer going, or decide on + # how to handle C struct addressing if we go for unpacked structures + prelude = ".global .u32 camera_coords[8];" + + def _cam_coord_xf(self, x, y, dreg): + """ + Given `.f32 x, y`, a coordinate in IFS space, writes the integer + offset from the start of the sampling lattice into `.u32 dreg`. + """ + + return """{ + .pred is_badval; + // TODO: This will change when data streaming is done + .reg .u32 camera_coord_address; + mov.u32 camera_coord_address, camera_coords; + // TODO: see if preloading everything hurts register count + .reg .f32 width_scale, width_upper_bound, height_scale, height_upper_bound; + ldu.v4.f32 {width_scale, width_upper_bound, + height_scale, height_upper_bound}, + [camera_coord_address+0]; + .reg .f32 x_xf, y_xf; + mad.rz.f32 x_xf, x, width_scale""" + # TODO unfinished + + + diff --git a/main.py b/main.py index aa29c57..8bc8636 100644 --- a/main.py +++ b/main.py @@ -11,39 +11,25 @@ import os import sys -import time -import ctypes -import struct - -# These imports are order-sensitive! -import pyglet -import pyglet.gl as gl -gl.get_current_context() +from ctypes import * import numpy as np -from fr0stlib import pyflam3 -from cuburnlib.device_code import MWCRNGTest -from cuburnlib.cuda import LaunchContext - -class CUGenome(pyflam3.Genome): - def _render(self, frame, trans): - obuf = (ctypes.c_ubyte * ((3+trans)*self.width*self.height))() - stats = pyflam3.RenderStats() - pyflam3.flam3_render(ctypes.byref(frame), obuf, - pyflam3.flam3_field_both, - trans+3, trans, ctypes.byref(stats)) - return obuf, stats, frame - +#from cuburnlib.device_code import MWCRNGTest +#from cuburnlib.cuda import LaunchContext +from fr0stlib.pyflam3 import * +from fr0stlib.pyflam3._flam3 import * +from cuburnlib.render import * def main(genome_path): - ctx = LaunchContext([MWCRNGTest], block=(256,1,1), grid=(64,1), tests=True) - ctx.compile(True) - ctx.run_tests() + #ctx = LaunchContext([MWCRNGTest], block=(256,1,1), grid=(64,1), tests=True) + #ctx.compile(True) + #ctx.run_tests() with open(genome_path) as fp: - genome = CUGenome.from_string(fp.read())[0] - + genomes = Genome.from_string(fp.read()) + render = Render(genomes) + render.render_frame() #genome.width, genome.height = 512, 512 #genome.sample_density = 1000