diff --git a/Makefile b/Makefile index b0da7f1..13cc5a4 100644 --- a/Makefile +++ b/Makefile @@ -1,26 +1,25 @@ CXX = gcc # clFFT lib & inc -CLFFT_LIB = -lOpenCL -L./clFFT/build/package/lib64 -lclFFT +CLFFT_LIB = -lOpenCL -L./usr/local/lib64 -lclFFT CLFFT_INCLUDE = -I./clFFT/build/package/include # standard math library CXXFLAGS = -c $(CLFFT_INCLUDE) -LDFLAGS = -lm $(CLFFT_LIB) +LDFLAGS = -lm $(CLFFT_LIB) -lfftw3 -lm EXE = Example all: ex04 ex05 -# entire process ex04: example04/build/main.o @if [ ! -d "./example04/bin" ]; then mkdir ./example04/bin; fi $(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) diff --git a/README.md b/README.md index 6499001..8203aff 100644 --- a/README.md +++ b/README.md @@ -22,8 +22,8 @@ gcc main.c -o main.out -lOpenCL For examples 04 and 05, you can run ```bash -make ex04 -make ex05 +make ex04 # executable is ./example04/bin/Example +make ex05 # executable is ./example05/bin/Example make # makes both! ``` @@ -54,10 +54,17 @@ 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-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. +An example of the CLFFT library for an in-place complex-planar transform. There is also Python code to check the answer; FFTW code will be added later, probably. + +- clFFT is required; installation instructions can be found inside example04/README.md +- for Python, numpy and scipy are required ## 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. +Another CLFFT example where an in-place real transform is 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` +- for Python, numpy and scipy are required ## 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 index 9786151..9aac985 100644 --- a/example04/README.md +++ b/example04/README.md @@ -10,17 +10,10 @@ mkdir build cd build cmake ../src make -make install +sudo make install +export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/usr/local/lib64 ``` -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 @@ -30,3 +23,4 @@ make ex04 ``` and it should print out a vector! :hamburger: + diff --git a/example05/main.c b/example05/main.c index 3954d7a..9f7dd98 100644 --- a/example05/main.c +++ b/example05/main.c @@ -2,6 +2,7 @@ #include #include #include +#include const char *kernelSource = "#pragma OPENCL EXTENSION cl_khr_fp64 : enable \n" \ @@ -45,7 +46,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 = 128; + unsigned int N = 2048; unsigned int N_pad = 2*roundUpToNearest( (N+2)/2, 32 ); size_t N_bytes = N_pad * sizeof(double); @@ -66,14 +67,13 @@ int main( int argc, char* argv[] ) { clfftSetup(&fftSetup); // host version of v - double *h_v; // real & imaginary parts + double *h_v; h_v = (double*) malloc(N_bytes); - // initialize v on host + // initialize v on host (GPU and CPU) int i; - for (i = 0; i < N; i++) { + for (i = 0; i < N; i++) h_v[i] = i; - } // global & local number of threads size_t globalSize, localSize; @@ -121,29 +121,67 @@ int main( int argc, char* argv[] ) { 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 + // FFT data, multiply elements, 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); + 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 epsilon) + arrays_equal = 0; + } + + if (arrays_equal) + printf("Arrays are equal!\n"); + else + printf("Arrays are NOT equal!\n"); + + // release FFT stuff + fftw_free(V); clfftDestroyPlan( &planHandleForward ); clfftDestroyPlan( &planHandleBackward ); clfftTeardown(); @@ -156,6 +194,7 @@ int main( int argc, char* argv[] ) { clReleaseContext(context); //release host memory + free(v); free(h_v); return 0; diff --git a/example05/main.py b/example05/main.py index 3bb9813..51d80c7 100644 --- a/example05/main.py +++ b/example05/main.py @@ -1,7 +1,7 @@ import numpy as np import scipy.fftpack as fft -N = 128 +N = 2048 v = np.arange(N) v_fft = fft.fft(v)