bgirgis/OpenCL-examples
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

update example01 with different number of global and local threads -- now first trial is slower... why

Browse Source
This commit is contained in:
Dakota St. Laurent
2015-06-16 18:49:02 -04:00
parent fd6b886c67
commit bd11706eac
2 changed files with 20 additions and 12 deletions
Download Patch File Download Diff File Expand all files Collapse all files

7
README.md
View File

@@ -41,3 +41,10 @@ so far, the third version (each thread doing only one element) is the fastest, a
- figure out how OpenCL manages memory (when are buffers cleared on the GPU?) - figure out how OpenCL manages memory (when are buffers cleared on the GPU?)
- figure out how to view the OpenCL assembly code if possible (is warp divergence happening?) - figure out how to view the OpenCL assembly code if possible (is warp divergence happening?)
## Some Notes
From the [guide on programming OpenCL for NVIDIA](http://www.nvidia.com/content/cudazone/download/OpenCL/NVIDIA_OpenCL_ProgrammingGuide.pdf):
- **CUDA streaming multiprocessor** corresponds to an OpenCL compute unit
- **CUDA thread** corresponds to an OpenCL work-item
- **CUDA thread block** corresponds to an OpenCL work-group

25
example01.cpp
View File

@@ -73,13 +73,13 @@ int main() {
// timing will be done on the host. // timing will be done on the host.
" void kernel looped_add(global const int* v1, global const int* v2, global int* v3, " " void kernel looped_add(global const int* v1, global const int* v2, global int* v3, "
" global const int* constants) {" " global const int* constants) {"
" int ID, Nthreads, n, k, ratio, start, stop;" " int ID, NUM_GLOBAL_WITEMS, n, k, ratio, start, stop;"
" ID = get_global_id(0);" " ID = get_global_id(0);"
" Nthreads = get_global_size(0);" " NUM_GLOBAL_WITEMS = get_global_size(0);"
" n = constants[0];" // size of vectors " n = constants[0];" // size of vectors
" k = constants[1];" // number of loop iterations " k = constants[1];" // number of loop iterations
"" ""
" ratio = (n / Nthreads);" // elements per thread " ratio = (n / NUM_GLOBAL_WITEMS);" // elements per thread
" start = ratio * ID;" " start = ratio * ID;"
" stop = ratio * (ID+1);" " stop = ratio * (ID+1);"
"" ""
@@ -92,12 +92,12 @@ int main() {
"" ""
" void kernel add(global const int* v1, global const int* v2, global int* v3, " " void kernel add(global const int* v1, global const int* v2, global int* v3, "
" global const int* constants) {" " global const int* constants) {"
" int ID, Nthreads, n, ratio, start, stop;" " int ID, NUM_GLOBAL_WITEMS, n, ratio, start, stop;"
" ID = get_global_id(0);" " ID = get_global_id(0);"
" Nthreads = get_global_size(0);" " NUM_GLOBAL_WITEMS = get_global_size(0);"
" n = constants[0];" " n = constants[0];"
"" ""
" ratio = (n / Nthreads);" " ratio = (n / NUM_GLOBAL_WITEMS);"
" start = ratio * ID;" " start = ratio * ID;"
" stop = ratio * (ID+1);" " stop = ratio * (ID+1);"
"" ""
@@ -120,10 +120,9 @@ int main() {
exit(1); exit(1);
} }
int n, k, Nthreads; const int n = 131072; // size of vectors (32 * 512 * 8)
n = 128000; // size of vectors const int k = 1000; // number of loop iterations
k = 1000; // number of loop iterations const int NUM_GLOBAL_WITEMS = 32 * 512;
Nthreads = 128;
int constants[2] = {n, k}; int constants[2] = {n, k};
// run the CPU code // run the CPU code
@@ -133,7 +132,7 @@ int main() {
// adds the same two vectors multiple times -- i.e. it's equivalent to the // adds the same two vectors multiple times -- i.e. it's equivalent to the
// host (CPU) code above, but with some necessary overhead. // host (CPU) code above, but with some necessary overhead.
cl::CommandQueue queue(context, default_device); cl::CommandQueue queue(context, default_device);
cl::KernelFunctor add(cl::Kernel(program, "add"), queue, cl::NullRange, cl::NDRange(Nthreads), cl::NullRange); cl::KernelFunctor add(cl::Kernel(program, "add"), queue, cl::NullRange, cl::NDRange(NUM_GLOBAL_WITEMS), cl::NDRange(32));
cl::Kernel looped_add = cl::Kernel(program, "looped_add"); cl::Kernel looped_add = cl::Kernel(program, "looped_add");
cl::Kernel add_single = cl::Kernel(program, "add_single"); cl::Kernel add_single = cl::Kernel(program, "add_single");
@@ -166,7 +165,9 @@ int main() {
looped_add.setArg(1, buffer_B); looped_add.setArg(1, buffer_B);
looped_add.setArg(2, buffer_C); looped_add.setArg(2, buffer_C);
looped_add.setArg(3, buffer_constants); looped_add.setArg(3, buffer_constants);
queue.enqueueNDRangeKernel(looped_add, cl::NullRange, 32, cl::NullRange); queue.enqueueNDRangeKernel(looped_add, cl::NullRange, // kernel, offset
cl::NDRange(NUM_GLOBAL_WITEMS), // global number of work items
cl::NDRange(32)); // local number (per group)
// read result from GPU to here; including for the sake of timing // read result from GPU to here; including for the sake of timing
queue.enqueueReadBuffer(buffer_C, CL_TRUE, 0, sizeof(int)*n, C); queue.enqueueReadBuffer(buffer_C, CL_TRUE, 0, sizeof(int)*n, C);