diff options
| author | Tor Aamodt <[email protected]> | 2010-07-18 13:58:22 -0800 |
|---|---|---|
| committer | Tor Aamodt <[email protected]> | 2010-07-18 13:58:22 -0800 |
| commit | 316dfac44cb01925017e495f471c84cce3174937 (patch) | |
| tree | 12fd456535b580fc287ed564e3a7c4a5279d428e /benchmarks/CUDA/template/template_kernel.cu | |
| parent | de305b7d6871110ad72037fb5e0a274c047838f3 (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/template/template_kernel.cu')
| -rw-r--r-- | benchmarks/CUDA/template/template_kernel.cu | 8 |
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]); |
