diff --git a/README.md b/README.md index e9158ea..be868ee 100644 --- a/README.md +++ b/README.md @@ -6,6 +6,7 @@ This currently runs on OS X, and I'm using local header files instead of global ``` git clone git@github.com:SaintDako/OpenCL-examples.git +cd OpenCL-examples mkdir CL curl https://www.khronos.org/registry/cl/api/1.2/cl.hpp -o CL/cl.hpp ``` diff --git a/example00/main.cpp b/example00/main.cpp index c41a496..17c6e82 100644 --- a/example00/main.cpp +++ b/example00/main.cpp @@ -1,5 +1,5 @@ #include -#include "CL/cl.hpp" +#include "../CL/cl.hpp" int main() { // get all platforms (drivers), e.g. NVIDIA diff --git a/example01/README.md b/example01/README.md index 0e0f2b4..230871b 100644 --- a/example01/README.md +++ b/example01/README.md @@ -4,15 +4,17 @@ This example compares the timings of adding vectors on the CPU versus adding vec ## Compiling ``` -clang++ -std=c++0x -framework OpenCL version01.cpp -o version01.out +clang++ -std=c++0x -framework OpenCL main.cpp -o main.out ``` To ignore deprecation warnings, add the flag `-Wno-deprecated-declarations`. +Run from this directory, as a relative path is used for the OpenCL header file (for now). + ## About -The code runs the following implementations of adding large vectors (131072 elements; 8 * 32 * 512). The vectors are added together 1000 times. +The code runs the following implementations of adding large vectors (131072 elements; 8 * 32 * 512). The vectors are added together 10000 times. - CPU -- GPU, where 32 * 512 threads are spawned and each thread thus gets 8 elements to calculate -- GPU, same as before but each iteration involves writing the buffers (to demonstrate overhead) -- GPU, where 8 * 32 * 512 threads are spawned - one for each element +- GPU, where 1024 threads are spawned and each thread thus gets 128 elements to calculate; there are two implementations of this: + - (Version 1) each thread gets 128 sequential elements (thread 0 gets 0-127, 1 gets 128-255, ...) + - (Version 2) each thread gets 128 elements, but coalescing happens (thread 0 gets 0,128,256..., thread 1 gets 1,129,257...) diff --git a/example01/main.cpp b/example01/main.cpp index 6be5adf..b523c7a 100644 --- a/example01/main.cpp +++ b/example01/main.cpp @@ -2,6 +2,7 @@ #include #include "../CL/cl.hpp" +#define NUM_GLOBAL_WITEMS 1024 void compareResults (double CPUtime, double GPUtime, int trial) { double time_ratio = (CPUtime / GPUtime); @@ -39,6 +40,30 @@ double timeAddVectorsCPU(int n, int k) { } +void warmup(cl::Context &context, cl::CommandQueue &queue, + cl::Kernel &add, int A[], int B[], int n) { + int C[n]; + // allocate space + cl::Buffer buffer_A(context, CL_MEM_READ_WRITE, sizeof(int) * n); + cl::Buffer buffer_B(context, CL_MEM_READ_WRITE, sizeof(int) * n); + cl::Buffer buffer_C(context, CL_MEM_READ_WRITE, sizeof(int) * n); + + // push write commands to queue + queue.enqueueWriteBuffer(buffer_A, CL_TRUE, 0, sizeof(int)*n, A); + queue.enqueueWriteBuffer(buffer_B, CL_TRUE, 0, sizeof(int)*n, B); + + // RUN ZE KERNEL + add.setArg(1, buffer_B); + add.setArg(0, buffer_A); + add.setArg(2, buffer_C); + for (int i=0; i<5; i++) + queue.enqueueNDRangeKernel(add, cl::NullRange, cl::NDRange(NUM_GLOBAL_WITEMS), cl::NDRange(32)); + + queue.enqueueReadBuffer(buffer_C, CL_TRUE, 0, sizeof(int)*n, C); + queue.finish(); +} + + int main(int argc, char* argv[]) { bool verbose; @@ -47,10 +72,9 @@ int main(int argc, char* argv[]) { else verbose = false; - const int n = 131072; // size of vectors (32 * 512 * 8) - const int k = 1000; // number of loop iterations - const int NUM_GLOBAL_WITEMS = 32 * 512; // number of threads for versions 1, 2 - int constants[2] = {n, k}; + const int n = 8*32*512; // size of vectors + const int k = 10000; // number of loop iterations + // const int NUM_GLOBAL_WITEMS = 1024; // number of threads // get all platforms (drivers), e.g. NVIDIA std::vector all_platforms; @@ -80,15 +104,17 @@ int main(int argc, char* argv[]) { // calculates for each element; C = A + B std::string kernel_code= - // is equivalent to the host's "time_add_vectors" function, except the - // timing will be done on the host. - " void kernel add_looped(global const int* v1, global const int* v2, global int* v3, " - " global const int* constants) {" - " int ID, NUM_GLOBAL_WITEMS, n, k, ratio, start, stop;" + " void kernel add(global const int* v1, global const int* v2, global int* v3) {" + " int ID;" + " ID = get_global_id(0);" + " v3[ID] = v1[ID] + v2[ID];" + " }" + "" + " void kernel add_looped_1(global const int* v1, global const int* v2, global int* v3, " + " const int n, const int k) {" + " int ID, NUM_GLOBAL_WITEMS, ratio, start, stop;" " ID = get_global_id(0);" " NUM_GLOBAL_WITEMS = get_global_size(0);" - " n = constants[0];" // size of vectors - " k = constants[1];" // number of loop iterations "" " ratio = (n / NUM_GLOBAL_WITEMS);" // elements per thread " start = ratio * ID;" @@ -101,43 +127,25 @@ int main(int argc, char* argv[]) { " }" " }" "" - " void kernel add(global const int* v1, global const int* v2, global int* v3, " - " global const int* constants) {" - " int ID, NUM_GLOBAL_WITEMS, n, ratio, start, stop;" + " void kernel add_looped_2(global const int* v1, global const int* v2, global int* v3," + " const int n, const int k) {" + " int ID, NUM_GLOBAL_WITEMS, step;" " ID = get_global_id(0);" " NUM_GLOBAL_WITEMS = get_global_size(0);" - " n = constants[0];" + " step = (n / NUM_GLOBAL_WITEMS);" "" - " ratio = (n / NUM_GLOBAL_WITEMS);" - " start = ratio * ID;" - " stop = ratio * (ID+1);" - "" - " for (int i=start; i