diff options
author | Anthony LaTorre <tlatorre9@gmail.com> | 2011-09-10 20:40:05 -0400 |
---|---|---|
committer | Anthony LaTorre <tlatorre9@gmail.com> | 2011-09-10 20:40:05 -0400 |
commit | 55cec145d72baf58639aad7773e23e28c9ca8676 (patch) | |
tree | ec6bb605140bdf4370a686b6e38b5920b93ba168 /src/render.cu | |
parent | 622bfafb66fdbd453cfa9fca05033787030d0364 (diff) | |
download | chroma-55cec145d72baf58639aad7773e23e28c9ca8676.tar.gz chroma-55cec145d72baf58639aad7773e23e28c9ca8676.tar.bz2 chroma-55cec145d72baf58639aad7773e23e28c9ca8676.zip |
store geometry struct in shared memory. this increases photon propagation speed from 3.3M -> 3.45!.
Diffstat (limited to 'src/render.cu')
-rw-r--r-- | src/render.cu | 9 |
1 files changed, 9 insertions, 0 deletions
diff --git a/src/render.cu b/src/render.cu index 005950f..d4ed443 100644 --- a/src/render.cu +++ b/src/render.cu @@ -39,11 +39,20 @@ render(int nthreads, float3 *_origin, float3 *_direction, Geometry *g, unsigned int alpha_depth, unsigned int *pixels, float *_dx, unsigned int *dxlen, float4 *_color) { + __shared__ Geometry sg; + + if (threadIdx.x == 0) + sg = *g; + + __syncthreads(); + int id = blockIdx.x*blockDim.x + threadIdx.x; if (id >= nthreads) return; + g = &sg; + float3 origin = _origin[id]; float3 direction = _direction[id]; unsigned int n = dxlen[id]; |