diff --git a/.gitignore b/.gitignore index 8606aaa..3f7dbfa 100644 --- a/.gitignore +++ b/.gitignore @@ -1,7 +1,7 @@ # compiled files *.out -**/bin -**/build +*/bin +*/build # clFFT library clFFT diff --git a/Makefile-04 b/Makefile-04 new file mode 100644 index 0000000..2b629f9 --- /dev/null +++ b/Makefile-04 @@ -0,0 +1,33 @@ +CXX = gcc + +# clFFT lib & inc +CLFFT_LIB = -lOpenCL -L./clFFT/build/package/lib64 -lclFFT +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 + + +# entire process +$(EXE): 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/$@ + +# 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 $@ + +# cleaning (remove executables and what not) +clean: + $(RM) -r ./example04/build/ ./example04/bin/ diff --git a/README.md b/README.md index a16c5d0..59f40d4 100644 --- a/README.md +++ b/README.md @@ -42,10 +42,11 @@ Measures the duration of adding two vectors. See the README in the folder for mo ## example 02 Demonstrates that one array can be modified several times without having to re-read and re-write data to and from the GPU. -## TODO +## example 03 +A simple example using the `cl_khr_fp64` extension which allows for usage of doubles instead of floats. -- 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?) +## example 04 +An example of the CLFFT library for an in-place complex-interleaved transform. ## 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/main.c b/example04/main.c new file mode 100644 index 0000000..97d1572 --- /dev/null +++ b/example04/main.c @@ -0,0 +1,144 @@ +#include +#include +#include +#include + +const char *kernelSource = +"#pragma OPENCL EXTENSION cl_khr_fp64 : enable \n" \ +"__kernel void mult(__global double *vR, __global double *vI) { \n" \ +" int id; \n" \ +" id = get_global_id(0); \n" \ +" vR[id] = 2*vR[id]; \n" \ +" vI[id] = 2*vI[id]; \n" \ +"} \n" \ +"\n" ; + +int main( int argc, char* argv[] ) { + // problem-related declarations + unsigned int N = 128; + size_t N_bytes = N * 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_vR, *h_vI; // real & imaginary parts + h_vR = (double*) malloc(N_bytes); + h_vI = (double*) malloc(N_bytes); + + // initialize v on host + int i; + for (i = 0; i < N; i++) { + h_vR[i] = i; + h_vI[i] = 2*i; + } + + // global & local number of threads + size_t globalSize, localSize; + globalSize = N; + 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_vR, d_vI; + d_vR = clCreateBuffer(context, CL_MEM_READ_WRITE, N_bytes, NULL, NULL); + d_vI = clCreateBuffer(context, CL_MEM_READ_WRITE, N_bytes, NULL, NULL); + err = clEnqueueWriteBuffer(queue, d_vR, CL_TRUE, 0, N_bytes, h_vR, 0, NULL, NULL); + err |= clEnqueueWriteBuffer(queue, d_vI, CL_TRUE, 0, N_bytes, h_vI, 0, NULL, NULL); + + // create forward plan and set its params + clfftCreateDefaultPlan(&planHandleForward, context, dim, clLengths); + clfftSetPlanPrecision(planHandleForward, CLFFT_DOUBLE); + clfftSetLayout(planHandleForward, CLFFT_COMPLEX_PLANAR, CLFFT_COMPLEX_PLANAR); + 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_COMPLEX_PLANAR, CLFFT_COMPLEX_PLANAR); + 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_vR); + err |= clSetKernelArg(k_mult, 1, sizeof(cl_mem), &d_vI); + + // cl_mem array allows for complex_planar transform + cl_mem inputBuffers[2] = {0, 0}; + inputBuffers[0] = d_vR; + inputBuffers[1] = d_vI; + + // FFT data, apply psi, IFFT data + 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"); + + clfftEnqueueTransform(planHandleBackward, CLFFT_BACKWARD, 1, &queue, 0, NULL, NULL, &inputBuffers, NULL, NULL); + + // transfer back + clEnqueueReadBuffer(queue, d_vR, CL_TRUE, 0, N_bytes, h_vR, 0, NULL, NULL ); + clEnqueueReadBuffer(queue, d_vI, CL_TRUE, 0, N_bytes, h_vI, 0, NULL, NULL ); + clFinish(queue); + + printf("[ "); + for (i=0; i