aaron/ispc
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

handwired CDP launch

Browse Source
This commit is contained in:
Evghenii
2013-11-12 11:25:43 +01:00
parent 4e5299a9bf
commit f445a470df
2 changed files with 10 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files

5
examples_cuda/aobench/ao.cu
View File

@@ -417,6 +417,11 @@ void ao_ispc_tasks(
const int nby = ntiley; const int nby = ntiley;
const int nbz = 1; const int nbz = 1;
const dim3 blocks (nbx, nby, nbz); const dim3 blocks (nbx, nby, nbz);
if (threadIdx.x == 0)
{
printf(" --- using CDP -- \n");
ao_task<<<blocks, 128>>>(w,h,nsubsamples,image); ao_task<<<blocks, 128>>>(w,h,nsubsamples,image);
} }
cudaDeviceSynchronize();
}
#endif #endif

6
examples_cuda/aobench/ao_cu.cpp
View File

@@ -340,8 +340,10 @@ extern "C"
const char * module = &module_str[0]; const char * module = &module_str[0];
#endif #endif
CUmodule cudaModule = loadModule(module); CUmodule cudaModule = loadModule(module);
CUfunction cudaFunction = getFunction(cudaModule, func_name); // CUfunction cudaFunction = getFunction(cudaModule, func_name);
deviceLaunch(cudaFunction, countx, county, countz, func_args); // deviceLaunch(cudaFunction, countx, county, countz, func_args);
CUfunction cudaFunction = getFunction(cudaModule, "ao_ispc_tasks");
deviceLaunch(cudaFunction, 1, 1, 1, func_args);
unloadModule(cudaModule); unloadModule(cudaModule);
} }
void CUDASync(void *handle) void CUDASync(void *handle)