summaryrefslogtreecommitdiff
path: root/benchmarks/CUDA
diff options
context:
space:
mode:
authorTor Aamodt <[email protected]>2010-07-18 13:58:22 -0800
committerTor Aamodt <[email protected]>2010-07-18 13:58:22 -0800
commit316dfac44cb01925017e495f471c84cce3174937 (patch)
tree12fd456535b580fc287ed564e3a7c4a5279d428e /benchmarks/CUDA
parentde305b7d6871110ad72037fb5e0a274c047838f3 (diff)
continuing refactoring
modify template to actually do a function call (for testing real function call, not printf) [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6884]
Diffstat (limited to 'benchmarks/CUDA')
-rw-r--r--benchmarks/CUDA/template/template_kernel.cu8
1 files changed, 7 insertions, 1 deletions
diff --git a/benchmarks/CUDA/template/template_kernel.cu b/benchmarks/CUDA/template/template_kernel.cu
index 4ced538..14610d8 100644
--- a/benchmarks/CUDA/template/template_kernel.cu
+++ b/benchmarks/CUDA/template/template_kernel.cu
@@ -30,6 +30,12 @@
//! @param g_idata input data in global memory
//! @param g_odata output data in global memory
////////////////////////////////////////////////////////////////////////////////
+
+__device__ float foo( float bar )
+{
+ return 2.0f*bar;
+}
+
__global__ void
testKernel( float* g_idata, float* g_odata)
{
@@ -45,7 +51,7 @@ testKernel( float* g_idata, float* g_odata)
// read in input data from global memory
// use the bank checker macro to check for bank conflicts during host
// emulation
- SDATA(tid) = g_idata[tid];
+ SDATA(tid) = foo(g_idata[tid])/2.0;
__syncthreads();
printf("thread tid=%u reads %f from g_idata[]\n", tid, sdata[tid]);