diff --git a/Makefile b/Makefile index 13cc5a4..52dc8f9 100644 --- a/Makefile +++ b/Makefile @@ -1,14 +1,19 @@ CXX = gcc # clFFT lib & inc -CLFFT_LIB = -lOpenCL -L./usr/local/lib64 -lclFFT -CLFFT_INCLUDE = -I./clFFT/build/package/include +CLFFT_LIB = -L/usr/local/cuda-7.0/targets/x86_64-linux/lib -lOpenCL -L./usr/local/lib64 -lclFFT # standard math library CXXFLAGS = -c $(CLFFT_INCLUDE) LDFLAGS = -lm $(CLFFT_LIB) -lfftw3 -lm EXE = Example +# ignore warnings when compiling if warn=0 +ifeq ($(warn), 0) + CXXFLAGS += -w +endif + + all: ex04 ex05 ex04: example04/build/main.o @@ -24,7 +29,6 @@ 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 $@ diff --git a/README.md b/README.md index 8203aff..c865c0f 100644 --- a/README.md +++ b/README.md @@ -60,7 +60,7 @@ An example of the CLFFT library for an in-place complex-planar transform. There - for Python, numpy and scipy are required ## example 05 -Another CLFFT example where an in-place real transform is performed. There's also FFTW code and Python code for checking the answer. +Another CLFFT example where an in-place real transform and an out-of-place real transform are performed. There's also FFTW code and Python code for checking the answer. - clFFT is required; installation instructions can be found inside example04/README.md - FFTW is required; installation is as simple as extracting FFTW's tar file, then running `./configure && sudo make && sudo make install` diff --git a/example05/main.c b/example05/main.c index 9f7dd98..6b3b980 100644 --- a/example05/main.c +++ b/example05/main.c @@ -3,8 +3,9 @@ #include #include #include - -const char *kernelSource = + + +static 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" \ @@ -27,6 +28,23 @@ int roundUpToNearest(int x, int n) { return x + (n - x_rem); } + +void checkIfArraysEqual(double *h_v, double *v, int N, double epsilon) { + int arrays_equal = 1; + int i; + + for (i=0; i epsilon) + arrays_equal = 0; + } + + if (arrays_equal) + printf("Arrays are equal!\n"); + else + printf("Arrays are NOT equal!\n"); +} + int main( int argc, char* argv[] ) { /* This setup is a bit tricky. Since we're doing a real transform, CLFFT @@ -46,7 +64,7 @@ int main( int argc, char* argv[] ) { * by 2 (192). The kernel will operate on zeros, but it should be faster * than the scenario with warp divergence. */ - unsigned int N = 2048; + unsigned int N = 4096; unsigned int N_pad = 2*roundUpToNearest( (N+2)/2, 32 ); size_t N_bytes = N_pad * sizeof(double); @@ -75,6 +93,37 @@ int main( int argc, char* argv[] ) { for (i = 0; i < N; i++) h_v[i] = i; + // CPU TRANSFORM ---------------------------------------------------------- + double *v; + fftw_complex *V; + int N_COMPLEX = N/2 + 1; + int REAL = 0; + int IMAG = 1; + + v = (double*) malloc(N * sizeof(double)); + V = (fftw_complex*) malloc(N_COMPLEX * sizeof(fftw_complex)); + + fftw_plan fft = fftw_plan_dft_r2c_1d(N, v, V, FFTW_MEASURE); + fftw_plan ifft = fftw_plan_dft_c2r_1d(N, V, v, FFTW_MEASURE); + + // initialize v here because otherwise fftw_execute will run before we + // initialize the plan... for some reason. + for (i=0; i epsilon) - arrays_equal = 0; - } - if (arrays_equal) - printf("Arrays are equal!\n"); - else - printf("Arrays are NOT equal!\n"); + clfftDestroyPlan( &planHandleForward ); + clfftDestroyPlan( &planHandleBackward ); + printf("Testing in-place real transform... "); + checkIfArraysEqual(h_v, v, N, 0.0); + + + + // REAL OUT-OF-PLACE TRANSFORM -------------------------------------------- + // reset array + d_v = clCreateBuffer(context, CL_MEM_READ_WRITE, N_bytes, NULL, NULL); + d_V = clCreateBuffer(context, CL_MEM_READ_WRITE, N_bytes, NULL, NULL); + + cl_mem inputBuffers[1] = {0}, outputBuffers[1] = {0}; + inputBuffers[0] = d_v; + outputBuffers[0] = d_V; + + err = clEnqueueWriteBuffer(queue, d_v, CL_TRUE, 0, N_bytes, h_v, 0, NULL, NULL); + + clfftCreateDefaultPlan(&planHandleForward, context, dim, clLengths); + clfftSetPlanPrecision(planHandleForward, CLFFT_DOUBLE); + clfftSetLayout(planHandleForward, CLFFT_REAL, CLFFT_HERMITIAN_INTERLEAVED); + clfftSetResultLocation(planHandleForward, CLFFT_OUTOFPLACE); + clfftBakePlan(planHandleForward, 1, &queue, NULL, NULL); + + clfftCreateDefaultPlan(&planHandleBackward, context, dim, clLengths); + clfftSetPlanPrecision(planHandleBackward, CLFFT_DOUBLE); + clfftSetLayout(planHandleBackward, CLFFT_HERMITIAN_INTERLEAVED, CLFFT_REAL); + clfftSetResultLocation(planHandleBackward, CLFFT_OUTOFPLACE); + clfftBakePlan(planHandleBackward, 1, &queue, NULL, NULL); + + clfftEnqueueTransform(planHandleForward, CLFFT_FORWARD, 1, &queue, 0, NULL, NULL, &inputBuffers, &outputBuffers, 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, &inputBuffers, &outputBuffers, NULL); + clFinish(queue); + clEnqueueReadBuffer(queue, d_v, CL_TRUE, 0, N_bytes, h_v, 0, NULL, NULL ); + clFinish(queue); + + printf("Testing out-of-place transform... "); + checkIfArraysEqual(h_v, v, N, 0.0); + // release FFT stuff fftw_free(V);