diff --git a/Makefile-04 b/Makefile similarity index 55% rename from Makefile-04 rename to Makefile index 2b629f9..b0da7f1 100644 --- a/Makefile-04 +++ b/Makefile @@ -7,27 +7,31 @@ CLFFT_INCLUDE = -I./clFFT/build/package/include # standard math library CXXFLAGS = -c $(CLFFT_INCLUDE) LDFLAGS = -lm $(CLFFT_LIB) -EXE = Example04 - -all: $(EXE) -compile: example04/build/main.o +EXE = Example +all: ex04 ex05 # entire process -$(EXE): example04/build/main.o +ex04: example04/build/main.o @if [ ! -d "./example04/bin" ]; then mkdir ./example04/bin; fi - $(CXX) $< $(LDFLAGS) -o example04/bin/$@ - -# link only -link: - @if [ ! -d "./example04/bin" ]; then mkdir ./example04/bin; fi - $(CXX) $< $(LDFLAGS) -o example04/bin/$@ + $(CXX) $< $(LDFLAGS) -o example04/bin/$(EXE) # create object file (compile without linking) example04/build/main.o: example04/main.c @if [ ! -d "./example04/build" ]; then mkdir ./example04/build; fi $(CXX) $(CXXFLAGS) $< -o $@ +ex05: example05/build/main.o + @if [ ! -d "./example05/bin" ]; then mkdir ./example05/bin; fi + $(CXX) $< $(LDFLAGS) -o example05/bin/$(EXE) + +# create object file (compile without linking) +example05/build/main.o: example05/main.c + @if [ ! -d "./example05/build" ]; then mkdir ./example05/build; fi + $(CXX) $(CXXFLAGS) $< -o $@ + + # cleaning (remove executables and what not) clean: $(RM) -r ./example04/build/ ./example04/bin/ + $(RM) -r ./example05/build/ ./example05/bin/ diff --git a/README.md b/README.md index 59f40d4..6499001 100644 --- a/README.md +++ b/README.md @@ -19,6 +19,14 @@ To compile the C code: gcc main.c -o main.out -lOpenCL ``` +For examples 04 and 05, you can run + +```bash +make ex04 +make ex05 +make # makes both! +``` + ### OS X OpenCL is installed on OS X by default, but since this code uses the C++ bindings, you'll need to get that too. Get the [official C++ bindings from the OpenCL registr](https://www.khronos.org/registry/cl/api/1.1/cl.hpp) and copy it to the OpenCL framework directory, or do the following: @@ -46,7 +54,10 @@ Demonstrates that one array can be modified several times without having to re-r A simple example using the `cl_khr_fp64` extension which allows for usage of doubles instead of floats. ## example 04 -An example of the CLFFT library for an in-place complex-interleaved transform. +An example of the CLFFT library for an in-place complex-planar transform. There is also Python code to check the answer, which requires numpy / scipy. The C code requires the CLFFT library to be installed in the root of the repository. See more details in the folder's readme. + +## example 05 +Another CLFFT example where an in-place real transform is performed. There's also Python code for checking the answer, which requires numpy / scipy. The C code requires the CLFFT library to be installed in the root of the repository. For instructions on doing this, check out the readme of example04. ## Some Notes From the [guide on programming OpenCL for NVIDIA](http://www.nvidia.com/content/cudazone/download/OpenCL/NVIDIA_OpenCL_ProgrammingGuide.pdf): diff --git a/example04/README.md b/example04/README.md new file mode 100644 index 0000000..9786151 --- /dev/null +++ b/example04/README.md @@ -0,0 +1,32 @@ +# example 04 + +## Installing CLFFT +After cloning the repository, run the following (in the top-level of the directory): + +``` +git clone https://github.com/clMathLibraries/clFFT.git +cd clFFT +mkdir build +cd build +cmake ../src +make +make install +``` + +and then + +``` +export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:`pwd`/clFFT/build/package/lib64 +``` + +(`pwd` should give the path to the top-level of the repository!). + +## Running it +In the top-level directory, run + +``` +make ex04 +./example04/bin/Example +``` + +and it should print out a vector! :hamburger: diff --git a/example04/main.c b/example04/main.c index 97d1572..52ee435 100644 --- a/example04/main.c +++ b/example04/main.c @@ -107,9 +107,7 @@ int main( int argc, char* argv[] ) { clfftEnqueueTransform(planHandleForward, CLFFT_FORWARD, 1, &queue, 0, NULL, NULL, &inputBuffers, NULL, NULL); clFinish(queue); - err = clEnqueueNDRangeKernel(queue, k_mult, 1, NULL, &globalSize, &localSize, 0, NULL, NULL); - if (err != CL_SUCCESS) - printf("oh wtf\n"); + err = clEnqueueNDRangeKernel(queue, k_mult, 1, NULL, &globalSize, &localSize, 0, NULL, NULL); clfftEnqueueTransform(planHandleBackward, CLFFT_BACKWARD, 1, &queue, 0, NULL, NULL, &inputBuffers, NULL, NULL); diff --git a/example05/main.c b/example05/main.c new file mode 100644 index 0000000..3954d7a --- /dev/null +++ b/example05/main.c @@ -0,0 +1,162 @@ +#include +#include +#include +#include + +const char *kernelSource = +"#pragma OPENCL EXTENSION cl_khr_fp64 : enable \n" \ +"__kernel void mult(__global double *v) { \n" \ +" int id, v_re, v_im; \n" \ +" id = get_global_id(0); \n" \ +" v_re = 2*id; \n" \ +" v_im = v_re + 1; \n" \ +" \n" \ +" v[v_re] = 2*v[v_re]; \n" \ +" v[v_im] = 4*v[v_im]; \n" \ +"} \n" \ +"\n" ; + + +int roundUpToNearest(int x, int n) { + /* Rounds x UP to nearest multiple of n. */ + int x_rem = x % n; + if (x_rem == 0) + return x; + + return x + (n - x_rem); +} + + +int main( int argc, char* argv[] ) { + /* This setup is a bit tricky. Since we're doing a real transform, CLFFT + * requires N+2 elements in the array. This is because only N/2 + 1 numbers + * are calculated, and since each number is complex, it requires 2 elements + * for space. + * + * To avoid warp divergence, we want to avoid any conditionals in the + * kernel. Thus we cannot check to see if the thread ID is even or odd to + * act on a real number or imaginary number. To do this, one thread should + * handle one complex number (one real, one imag), i.e. ID_j should handle + * array elements j, j+1. + * + * But we also need the number of global items to be a multiple of 32 (warp + * size). What we can do, for example, N = 128, is pad it by 2 (130), + * divide it by 2 (65), round that UP to the nearest 32 (96), multiply that + * by 2 (192). The kernel will operate on zeros, but it should be faster + * than the scenario with warp divergence. */ + + unsigned int N = 128; + unsigned int N_pad = 2*roundUpToNearest( (N+2)/2, 32 ); + size_t N_bytes = N_pad * sizeof(double); + + // openCL declarations + cl_platform_id platform; + cl_device_id device_id; + cl_context context; + cl_command_queue queue; + cl_program program; + cl_kernel k_mult; + + // clFFT declarations + clfftPlanHandle planHandleForward, planHandleBackward; + clfftDim dim = CLFFT_1D; + size_t clLengths[1] = {N}; + clfftSetupData fftSetup; + clfftInitSetupData(&fftSetup); + clfftSetup(&fftSetup); + + // host version of v + double *h_v; // real & imaginary parts + h_v = (double*) malloc(N_bytes); + + // initialize v on host + int i; + for (i = 0; i < N; i++) { + h_v[i] = i; + } + + // global & local number of threads + size_t globalSize, localSize; + globalSize = N_pad / 2; + localSize = 32; + + // setup OpenCL stuff + cl_int err; + err = clGetPlatformIDs(1, &platform, NULL); + err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device_id, NULL); + context = clCreateContext(0, 1, &device_id, NULL, NULL, &err); + queue = clCreateCommandQueue(context, device_id, 0, &err); + program = clCreateProgramWithSource(context, 1, (const char **) & kernelSource, NULL, &err); + + // Build the program executable + err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL); + if (err != CL_SUCCESS) { + printf("building program failed\n"); + if (err == CL_BUILD_PROGRAM_FAILURE) { + size_t log_size; + clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size); + char *log = (char *) malloc(log_size); + clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, log_size, log, NULL); + printf("%s\n", log); + } + } + k_mult = clCreateKernel(program, "mult", &err); + + // create arrays on host and write them + cl_mem d_v; + d_v = clCreateBuffer(context, CL_MEM_READ_WRITE, N_bytes, NULL, NULL); + err = clEnqueueWriteBuffer(queue, d_v, CL_TRUE, 0, N_bytes, h_v, 0, NULL, NULL); + + // create forward plan and set its params + clfftCreateDefaultPlan(&planHandleForward, context, dim, clLengths); + clfftSetPlanPrecision(planHandleForward, CLFFT_DOUBLE); + clfftSetLayout(planHandleForward, CLFFT_REAL, CLFFT_HERMITIAN_INTERLEAVED); + clfftSetResultLocation(planHandleForward, CLFFT_INPLACE); + clfftBakePlan(planHandleForward, 1, &queue, NULL, NULL); + + // create backward plan and set its params + clfftCreateDefaultPlan(&planHandleBackward, context, dim, clLengths); + clfftSetPlanPrecision(planHandleBackward, CLFFT_DOUBLE); + clfftSetLayout(planHandleBackward, CLFFT_HERMITIAN_INTERLEAVED, CLFFT_REAL); + clfftSetResultLocation(planHandleBackward, CLFFT_INPLACE); + clfftBakePlan(planHandleBackward, 1, &queue, NULL, NULL); + + // set all of ze kernel args... + err = clSetKernelArg(k_mult, 0, sizeof(cl_mem), &d_v); + + // FFT data, apply psi, IFFT data + clfftEnqueueTransform(planHandleForward, CLFFT_FORWARD, 1, &queue, 0, NULL, NULL, &d_v, NULL, NULL); + clFinish(queue); + + err = clEnqueueNDRangeKernel(queue, k_mult, 1, NULL, &globalSize, &localSize, 0, NULL, NULL); + clFinish(queue); + + //clfftEnqueueTransform(planHandleBackward, CLFFT_BACKWARD, 1, &queue, 0, NULL, NULL, &d_v, NULL, NULL); + clFinish(queue); + + // transfer back + clEnqueueReadBuffer(queue, d_v, CL_TRUE, 0, N_bytes, h_v, 0, NULL, NULL ); + clFinish(queue); + + printf("[ "); + for (i=0; i