summaryrefslogtreecommitdiff
path: root/src/render.cu
diff options
context:
space:
mode:
authorAnthony LaTorre <tlatorre9@gmail.com>2011-09-10 20:40:05 -0400
committerAnthony LaTorre <tlatorre9@gmail.com>2011-09-10 20:40:05 -0400
commit55cec145d72baf58639aad7773e23e28c9ca8676 (patch)
treeec6bb605140bdf4370a686b6e38b5920b93ba168 /src/render.cu
parent622bfafb66fdbd453cfa9fca05033787030d0364 (diff)
downloadchroma-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.cu9
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];