add windows compiling instructions ... ugh ...

- makefile now uses system-installed version of clFFT

.gitignore

@@ -1,10 +1,11 @@
-# openCL C++ headers
-CL/
-
 # compiled files
 *.out
+*/bin
+*/build
+
+# clFFT library
+clFFT
 
 # THIS IS MY SWAP
 # (vim swap files)
-*.swp
-*.swo
+.*.s??

Makefile

@@ -0,0 +1,40 @@
+CXX = gcc 
+
+# clFFT lib & inc
+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
+	@if [ ! -d "./example04/bin" ]; then mkdir ./example04/bin; fi
+	$(CXX) $< $(LDFLAGS) -o example04/bin/$(EXE)
+
+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)
+
+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/

README.md

@@ -2,24 +2,88 @@
 here is my feeble attempt at learning OpenCL, please don't make fun of me too much :hamburger:
 
 ## Configuration
-This currently runs on OS X, and I'm using local header files instead of global header files because I'm unfamiliar with C++. Deal with it. Run the following in a terminal to set up:
+This code uses OpenCL 1.1 on a NVIDIA GPU.
+
+### Linux
+(Only tested on Ubuntu). For NVIDIA GPUs, I've installed the following packages: `nvidia-346 nvidia-346-dev nvidia-346-uvm nvidia-libopencl1-346 nvidia-modprobe nvidia-opencl-icd-346 nvidia-settings`. Since the `opencl-headers` package in the main repository is for OpenCL 1.2, you can get the OpenCL 1.1 header files from [here](http://packages.ubuntu.com/precise/opencl-headers).
+
+Then to compile the C++ code:
 
 ```
-git clone git@github.com:SaintDako/OpenCL-examples.git
-mkdir CL
-curl https://www.khronos.org/registry/cl/api/1.2/cl.hpp -o CL/cl.hpp
+g++ -std=c++0x main.cpp -o main.out -lOpenCL
 ```
 
+To compile the C code:
+
+```
+gcc main.c -o main.out -lOpenCL
+```
+
+For examples 04 and 05, you can run
+
+```bash
+make ex04  # executable is ./example04/bin/Example
+make ex05  # executable is ./example05/bin/Example
+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:
+
+```
+wget https://www.khronos.org/registry/cl/api/1.1/cl.hpp
+sudo cp cl.hpp /System/Library/Frameworks/OpenCL.framework/Headers/
+```
+
+To compile:
+
+```
+clang++ -std=c++0x -framework OpenCL main.cpp -o main.out
+```
+
+### Windows
+For some reason, the makefile didn't want to work for Windows. I have no idea why.
+
+For example 04, run (inside the directory):
+
+```
+gcc -I/c/Program\ Files/NVIDIA\ GPU\ Computing\ Toolkit/CUDA/v7.5/include -I/c/PATH/TO/CLFFT/include main.c -o main.exe -L/c/Program\ Files/NVIDIA\ GPU\ Computing\ Toolkit/CUDA/v7.5/lib/x64 -lOpenCL -L/c/PATH/TO/CLFFT/lib64/import -lclFFT
+```
+
+where `PATH/TO/CLFFT` is the path to the clFFT library.
+
+For example 05, run (inside the directory):
+
+```
+gcc -I/c/Program\ Files/NVIDIA\ GPU\ Computing\ Toolkit/CUDA/v7.5/include -I/c/PATH/TO/CLFFT/include -I/c/PATH/TO/FFTW main.c -o main.exe -L/c/Program\ Files/NVIDIA\ GPU\ Computing\ Toolkit/CUDA/v7.5/lib/x64 -lOpenCL -L/c/PATH/TO/CLFFT/lib64/import -lclFFT -L/c/PATH/TO/FFTW -lfftw3-3
+```
+
+where `PATH/TO/FFTW` is the path to the FFTW3 library.
+
 ## example 00
 this example is based off of [this example](simpleopencl.blogspot.ca/2013/06/tutorial-simple-start-with-opencl-and-c.html) (example-ception), but it goes a bit further. In the blogspot example, two 10-element vectors are created and a thread is used for each pair of elements. In this example, 10 threads are spawned but two 100-element vectors are used, and it is shown how to split up a specific number of elements per thread.
 
 ## example 01
-See the README in the folder.
+Measures the duration of adding two vectors. See the README in the folder for more details.
 
-## TODO
+## 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.
 
-- 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 03
+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; 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 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`
+- 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):

example00/main.cpp

@@ -1,5 +1,9 @@
 #include <iostream>
-#include "CL/cl.hpp"
+#ifdef __APPLE__
+    #include <OpenCL/cl.hpp>
+#else
+    #include <CL/cl.hpp>
+#endif
 
 int main() {
     // get all platforms (drivers), e.g. NVIDIA

example01/README.md

@@ -1,18 +1,10 @@
 # Example 01
 This example compares the timings of adding vectors on the CPU versus adding vectors on the GPU, the latter of which has different implementations.
 
-## Compiling
-
-```
-clang++ -std=c++0x -framework OpenCL version01.cpp -o version01.out
-```
-
-To ignore deprecation warnings, add the flag `-Wno-deprecated-declarations`.
-
 ## 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...)

example01/main.cpp

@@ -1,7 +1,12 @@
 #include <iostream>
 #include <ctime>
-#include "../CL/cl.hpp"
+#ifdef __APPLE__
+    #include <OpenCL/cl.hpp>
+#else
+    #include <CL/cl.hpp>
+#endif
 
+#define NUM_GLOBAL_WITEMS 1024
 
 void compareResults (double CPUtime, double GPUtime, int trial) {
     double time_ratio = (CPUtime / GPUtime);
@@ -39,6 +44,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 +76,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<cl::Platform> all_platforms;
@@ -80,15 +108,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 +131,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<stop; i++)"
-        "           v3[i] = v1[i] + v2[i];"
+        "       int i,j;"
+        "       for (i=0; i<k; i++) {"
+        "           for (j=ID; j<n; j+=step)"
+        "               v3[j] = v1[j] + v2[j];"
+        "       }"
         "   }"
         ""    
         "   void kernel add_single(global const int* v1, global const int* v2, global int* v3, "
-        "                          global const int* constants) { "
-        "       int k  = constants[1];"
+        "                          const int k) { "
         "       int ID = get_global_id(0);"
         "       for (int i=0; i<k; i++)"
         "           v3[ID] = v1[ID] + v2[ID];"
-        "   }"
-        ""  // same as add_single, but with the overhead (indexing, determining ratio) of versions 01 and 02
-        "   void kernel add_single_overhead(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;"
-        "       ID = get_global_id(0);"
-        "       NUM_GLOBAL_WITEMS = get_global_size(0);"
-        "       n = constants[0];"
-        "       k = constants[1];"
-        ""
-        "       ratio = (n / NUM_GLOBAL_WITEMS);"
-        "       start = ratio * ID;"
-        "       stop  = ratio * (ID+1);"
-        ""
-        "       for (int i=0; i<k; i++)"
-        "           v3[ID] = v1[ID] + v2[ID];"
         "   }";
     sources.push_back({kernel_code.c_str(), kernel_code.length()});
 
@@ -152,20 +164,23 @@ int main(int argc, char* argv[]) {
 
     // set up kernels and vectors for GPU code
     cl::CommandQueue queue(context, default_device);
-    cl::Kernel add_looped = cl::Kernel(program, "add_looped");
     cl::Kernel add          = cl::Kernel(program, "add");
+    cl::Kernel add_looped_1 = cl::Kernel(program, "add_looped_1");
+    cl::Kernel add_looped_2 = cl::Kernel(program, "add_looped_2");
     cl::Kernel add_single   = cl::Kernel(program, "add_single");
-    cl::Kernel add_single_overhead = cl::Kernel(program, "add_single_overhead");
 
     // construct vectors
     int A[n], B[n], C[n];
     for (int i=0; i<n; i++) {
         A[i] = i;
         B[i] = n - i - 1;
-        C[i] = 0;
     }
-    std::clock_t start_time;
 
+    // attempt at warm-up...
+    warmup(context, queue, add, A, B, n);
+    queue.finish();
+
+    std::clock_t start_time;
 
     // VERSION 1 ==========================================
     // start timer
@@ -176,119 +191,51 @@ int main(int argc, char* argv[]) {
     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);
-    cl::Buffer buffer_constants(context, CL_MEM_READ_ONLY, sizeof(int) * 2);
 
     // 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);
-    queue.enqueueWriteBuffer(buffer_constants, CL_TRUE, 0, sizeof(int)*2, constants);
 
     // RUN ZE KERNEL
-    add_looped.setArg(0, buffer_A);
-    add_looped.setArg(1, buffer_B);
-    add_looped.setArg(2, buffer_C);
-    add_looped.setArg(3, buffer_constants);
-    queue.enqueueNDRangeKernel(add_looped, cl::NullRange,      // kernel, offset
+    add_looped_1.setArg(0, buffer_A);
+    add_looped_1.setArg(1, buffer_B);
+    add_looped_1.setArg(2, buffer_C);
+    add_looped_1.setArg(3, n);
+    add_looped_1.setArg(4, k);
+    queue.enqueueNDRangeKernel(add_looped_1, cl::NullRange,  // kernel, offset
             cl::NDRange(NUM_GLOBAL_WITEMS), // global number of work items
             cl::NDRange(32));               // local number (per group)
 
     // read result from GPU to here; including for the sake of timing
     queue.enqueueReadBuffer(buffer_C, CL_TRUE, 0, sizeof(int)*n, C);
-    queue.enqueueBarrier();
+    queue.finish();
     GPUtime1 = (std::clock() - start_time) / (double) CLOCKS_PER_SEC;
 
+
     // VERSION 2 ==========================================
     double GPUtime2;
-    start_time = std::clock();
 
     cl::Buffer buffer_A2(context, CL_MEM_READ_WRITE, sizeof(int)*n);
     cl::Buffer buffer_B2(context, CL_MEM_READ_WRITE, sizeof(int)*n);
     cl::Buffer buffer_C2(context, CL_MEM_READ_WRITE, sizeof(int)*n);
-    cl::Buffer buffer_constants2(context, CL_MEM_READ_ONLY, sizeof(int)*2);
-    for (int i=0; i<k; i++) {
     queue.enqueueWriteBuffer(buffer_A2, CL_TRUE, 0, sizeof(int)*n, A);
     queue.enqueueWriteBuffer(buffer_B2, CL_TRUE, 0, sizeof(int)*n, B);
-        queue.enqueueWriteBuffer(buffer_constants2, CL_TRUE, 0, sizeof(int)*2, constants);
 
-        add_looped.setArg(0, buffer_A2);
-        add_looped.setArg(1, buffer_B2);
-        add_looped.setArg(2, buffer_C2);
-        add_looped.setArg(3, buffer_constants);
-        queue.enqueueNDRangeKernel(add, cl::NullRange, cl::NDRange(NUM_GLOBAL_WITEMS), cl::NDRange(32));
-    }
+    start_time = std::clock();
+    add_looped_2.setArg(0, buffer_A2);
+    add_looped_2.setArg(1, buffer_B2);
+    add_looped_2.setArg(2, buffer_C2);
+    add_looped_2.setArg(3, n);
+    add_looped_2.setArg(4, k);
+    
+    queue.enqueueNDRangeKernel(add_looped_2, cl::NullRange, cl::NDRange(NUM_GLOBAL_WITEMS), cl::NDRange(32));
     queue.enqueueReadBuffer(buffer_C2, CL_TRUE, 0, sizeof(int)*n, C);
-    queue.enqueueBarrier();
+    queue.finish();
     GPUtime2 = (std::clock() - start_time) / (double) CLOCKS_PER_SEC;
 
-    // VERSION 3 ==========================================
-    double GPUtime3;
-    start_time = std::clock();
-
-    cl::Buffer buffer_A3(context, CL_MEM_READ_WRITE, sizeof(int)*n);
-    cl::Buffer buffer_B3(context, CL_MEM_READ_WRITE, sizeof(int)*n);
-    cl::Buffer buffer_C3(context, CL_MEM_READ_WRITE, sizeof(int)*n);
-    cl::Buffer buffer_constants3(context, CL_MEM_READ_ONLY, sizeof(int) * 2);
-    queue.enqueueWriteBuffer(buffer_A3, CL_TRUE, 0, sizeof(int)*n, A);
-    queue.enqueueWriteBuffer(buffer_B3, CL_TRUE, 0, sizeof(int)*n, B);
-    queue.enqueueWriteBuffer(buffer_constants3, CL_TRUE, 0, sizeof(int)*2, constants);
-
-    add_single.setArg(0, buffer_A3);
-    add_single.setArg(1, buffer_B3);
-    add_single.setArg(2, buffer_C3);
-    add_single.setArg(3, buffer_constants3);
-    queue.enqueueNDRangeKernel(add_single, cl::NullRange, cl::NDRange(n), cl::NDRange(32));
-
-    queue.enqueueReadBuffer(buffer_C3, CL_TRUE, 0, sizeof(int)*n, C);
-    queue.enqueueBarrier();
-    GPUtime3 = (std::clock() - start_time) / (double) CLOCKS_PER_SEC;
-
-    // VERSION 4 ==========================================
-    double GPUtime4;
-    start_time = std::clock();
-
-    cl::Buffer buffer_A4(context, CL_MEM_READ_WRITE, sizeof(int)*n);
-    cl::Buffer buffer_B4(context, CL_MEM_READ_WRITE, sizeof(int)*n);
-    cl::Buffer buffer_C4(context, CL_MEM_READ_WRITE, sizeof(int)*n);
-    cl::Buffer buffer_constants4(context, CL_MEM_READ_ONLY, sizeof(int)*2);
-    queue.enqueueWriteBuffer(buffer_A4, CL_TRUE, 0, sizeof(int)*n, A);
-    queue.enqueueWriteBuffer(buffer_B4, CL_TRUE, 0, sizeof(int)*n, B);
-    queue.enqueueWriteBuffer(buffer_constants4, CL_TRUE, 0, sizeof(int)*2, constants);
-
-    add_single_overhead.setArg(0, buffer_A4);
-    add_single_overhead.setArg(1, buffer_B4);
-    add_single_overhead.setArg(2, buffer_C4);
-    add_single_overhead.setArg(3, buffer_constants4);
-    queue.enqueueNDRangeKernel(add_single_overhead, cl::NullRange, cl::NDRange(n), cl::NDRange(32));
-
-    queue.enqueueReadBuffer(buffer_C4, CL_TRUE, 0, sizeof(int)*n, C);
-    queue.enqueueBarrier();
-    GPUtime4 = (std::clock() - start_time) / (double) CLOCKS_PER_SEC;
-
-    // VERSION 5 ==========================================
-    double GPUtime5;
-    start_time = std::clock();
-
-    cl::Buffer buffer_A5(context, CL_MEM_READ_WRITE, sizeof(int) * n);
-    cl::Buffer buffer_B5(context, CL_MEM_READ_WRITE, sizeof(int) * n);
-    cl::Buffer buffer_C5(context, CL_MEM_READ_WRITE, sizeof(int) * n);
-    cl::Buffer buffer_constants5(context, CL_MEM_READ_ONLY, sizeof(int) * 2);
-    queue.enqueueWriteBuffer(buffer_A5, CL_TRUE, 0, sizeof(int)*n, A);
-    queue.enqueueWriteBuffer(buffer_B5, CL_TRUE, 0, sizeof(int)*n, B);
-    queue.enqueueWriteBuffer(buffer_constants5, CL_TRUE, 0, sizeof(int)*2, constants);
-
-    add_looped.setArg(0, buffer_A5);
-    add_looped.setArg(1, buffer_B5);
-    add_looped.setArg(2, buffer_C5);
-    add_looped.setArg(3, buffer_constants5);
-    queue.enqueueNDRangeKernel(add_looped, cl::NullRange, cl::NDRange(n), cl::NDRange(32));
-
-    queue.enqueueReadBuffer(buffer_C5, CL_TRUE, 0, sizeof(int)*n, C);
-    queue.enqueueBarrier();
-    GPUtime5 = (std::clock() - start_time) / (double) CLOCKS_PER_SEC;
-    
     // let's compare!
-    const int NUM_VERSIONS = 5;
-    double GPUtimes[NUM_VERSIONS] = {GPUtime1, GPUtime2, GPUtime3, GPUtime4, GPUtime5};
+    const int NUM_VERSIONS = 2;
+    double GPUtimes[NUM_VERSIONS] = {GPUtime1, GPUtime2};
     if (verbose) {
         for (int i=0; i<NUM_VERSIONS; i++)
             compareResults(CPUtime, GPUtimes[i], i+1);

example02/main.c

@@ -0,0 +1,94 @@
+#include <stdlib.h>
+#include <stdio.h>
+#include <CL/cl.h>
+
+const char *kernel_code =
+    "__kernel void multiply_by(__global int* A, const int c) {"
+    "   A[get_global_id(0)] = c * A[get_global_id(0)];"
+    "}";
+
+
+int factorial(int n) {
+    return (n <= 1) ? 1 : n * factorial(n-1);
+}
+
+
+int main( void ) {
+    // OpenCL related declarations
+    cl_int err;
+    cl_platform_id platform;
+    cl_device_id device;
+    cl_context_properties props[3] = { CL_CONTEXT_PLATFORM, 0, 0 };
+    cl_context ctx;
+    cl_program program;
+    cl_command_queue queue;
+    cl_event event = NULL;
+    cl_kernel k_multiplyby;
+    int i;
+
+    // 
+    const size_t N = 1024; // vector size
+    const int c_max = 5;   // max value to iterate to
+    const int coeff = factorial(c_max);
+    
+    int *A, *B, *C;        // A is initial, B is result, C is expected result
+    A = (int*) malloc(N * sizeof(*A));
+    B = (int*) malloc(N * sizeof(*B));
+    C = (int*) malloc(N * sizeof(*C));
+    for (i=0; i<N; i++) {
+        A[i] = i;
+        C[i] = coeff*i;
+    }
+    cl_mem d_A;  // buffer object for A
+
+    /* Setup OpenCL environment. */
+    err = clGetPlatformIDs( 1, &platform, NULL );
+    err = clGetDeviceIDs( platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL );
+
+    props[1] = (cl_context_properties)platform;
+    ctx = clCreateContext( props, 1, &device, NULL, NULL, &err );
+    queue = clCreateCommandQueue( ctx, device, 0, &err );
+    program = clCreateProgramWithSource(ctx, 1, (const char **) &kernel_code, NULL, &err);
+    err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
+    k_multiplyby = clCreateKernel(program, "multiply_by", &err);
+
+    // initialize buffer with data
+    d_A = clCreateBuffer( ctx, CL_MEM_READ_WRITE, N*sizeof(*A), NULL, &err );
+
+    err = clEnqueueWriteBuffer( queue, d_A, CL_TRUE, 0, N*sizeof(*A), A, 0, NULL, NULL );
+    
+    clSetKernelArg(k_multiplyby, 0, sizeof(cl_mem), &d_A);
+    int c;
+    for (c=2; c<=c_max; c++) {
+        clSetKernelArg(k_multiplyby, 1, sizeof(int), &c);
+        clEnqueueNDRangeKernel(queue, k_multiplyby, 1, NULL, &N, &N, 0, NULL, NULL);
+    }
+    err = clFinish(queue);
+
+    err = clEnqueueReadBuffer( queue, d_A, CL_TRUE, 0, N*sizeof(*B), B, 0, NULL, NULL );
+    err = clFinish(queue);
+
+    int success = 1;
+    for (i=0; i<N; i++) {
+        if (B[i] != C[i]) {
+            success = 0;
+            break;
+        }
+    }
+
+    if (success)
+        printf("Arrays are equal!\n");
+    else
+        printf("Arrays are NOT equal\n");
+
+
+    /* Release OpenCL memory objects. */
+    clReleaseMemObject( d_A );
+    free(A);
+    free(B);
+    free(C);
+    clReleaseCommandQueue( queue );
+    clReleaseContext( ctx );
+
+    return 0;
+}

example02/main.cpp

@@ -0,0 +1,99 @@
+#include <iostream>
+#include <algorithm>
+#include <iterator>
+#ifdef __APPLE__
+    #include <OpenCL/cl.hpp>
+#else
+    #include <CL/cl.hpp>
+#endif
+
+using namespace std;
+using namespace cl;
+
+
+int factorial(int n) {
+    return (n <= 1) ? 1 : n * factorial(n-1);
+}
+
+
+Platform getPlatform() {
+    /* Returns the first platform found. */
+    std::vector<Platform> all_platforms;
+    Platform::get(&all_platforms);
+
+    if (all_platforms.size()==0) {
+        cout << "No platforms found. Check OpenCL installation!\n";
+        exit(1);
+    }
+    return all_platforms[0];
+}
+
+
+Device getDevice(Platform platform, int i, bool display=false) {
+    /* Returns the deviced specified by the index i on platform.
+     * If display is true, then all of the platforms are listed.
+     */
+    std::vector<Device> all_devices;
+    platform.getDevices(CL_DEVICE_TYPE_ALL, &all_devices);
+    if(all_devices.size()==0){
+        cout << "No devices found. Check OpenCL installation!\n";
+        exit(1);
+    }
+
+    if (display) {
+        for (int j=0; j<all_devices.size(); j++)
+            printf("Device %d: %s\n", j, all_devices[j].getInfo<CL_DEVICE_NAME>().c_str());
+    }
+    return all_devices[i];
+}
+
+
+int main() {
+    const int n = 1024;    // size of vectors
+    const int c_max = 5;   // max value to iterate to
+    const int coeff = factorial(c_max);
+
+    int A[n], B[n], C[n];     // A is initial, B is result, C is expected result
+    for (int i=0; i<n; i++) {
+        A[i] = i;
+        C[i] = coeff * i;
+    }
+    Platform default_platform = getPlatform();
+    Device default_device     = getDevice(default_platform, 1);
+    Context context({default_device});
+    Program::Sources sources;
+
+    std::string kernel_code=
+        "void kernel multiply_by(global int* A, const int c) {"
+        "   A[get_global_id(0)] = c * A[get_global_id(0)];"
+        "}";
+    sources.push_back({kernel_code.c_str(), kernel_code.length()});
+
+    Program program(context, sources);
+    if (program.build({default_device}) != CL_SUCCESS) {
+        cout << "Error building: " << program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(default_device) << std::endl;
+        exit(1);
+    }
+    
+    Buffer buffer_A(context, CL_MEM_READ_WRITE, sizeof(int) * n);
+    CommandQueue queue(context, default_device);
+    queue.enqueueWriteBuffer(buffer_A, CL_TRUE, 0, sizeof(int)*n, A);
+
+    Kernel multiply_by = Kernel(program, "multiply_by");
+    multiply_by.setArg(0, buffer_A);
+
+    for (int c=2; c<=c_max; c++) {
+        multiply_by.setArg(1, c);
+        queue.enqueueNDRangeKernel(multiply_by, NullRange, NDRange(n), NDRange(32));
+    }
+
+    queue.enqueueReadBuffer(buffer_A, CL_TRUE, 0, sizeof(int)*n, B);
+    
+    if (std::equal(std::begin(B), std::end(B), std::begin(C)))
+        cout << "Arrays are equal!" << endl;
+    else
+        cout << "Uh-oh, the arrays aren't equal!" << endl;
+
+    return 0;
+}
+

example03/main.c

@@ -0,0 +1,109 @@
+#include <stdio.h>
+#include <stdlib.h>
+#include <math.h>
+#include <CL/cl.h>
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+ 
+const char *kernelSource =
+"#pragma OPENCL EXTENSION cl_khr_fp64 : enable   \n"  \ 
+"__kernel void mult(__global double *v) {  \n" \
+"    int id;                   \n" \
+"    id = get_global_id(0);    \n" \
+"    v[id] = 2*v[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;
+
+    // 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;
+    localSize = 32;
+
+    // show the extensions that are supported
+    /*
+    cl_char extensions[2048] = {0};
+    clGetDeviceInfo(device_id, CL_DEVICE_EXTENSIONS, sizeof(extensions), &extensions, NULL);
+    printf("%s\n", extensions);
+    */
+
+    // 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);
+
+    err  = clSetKernelArg(k_mult, 0, sizeof(cl_mem), &d_v);
+ 
+    err = clEnqueueNDRangeKernel(queue, k_mult, 1, NULL, &globalSize, &localSize, 0, NULL, NULL);
+    clFinish(queue);
+
+    // transfer back
+    clEnqueueReadBuffer(queue, d_v, CL_TRUE, 0, N_bytes, h_v, 0, NULL, NULL );
+    clFinish(queue);
+
+    
+    int correct = 1; 
+    for (i=0; i<N; i++) {
+        if (h_v[i] != (double) 2*i)
+            correct = 0;
+    }
+    if (correct)
+        printf("Array is correct!\n");
+    else
+        printf("Array is incorrect :(\n");
+
+    // release OpenCL resources
+    clReleaseMemObject(d_v);
+    clReleaseProgram(program);
+    clReleaseKernel(k_mult);
+    clReleaseCommandQueue(queue);
+    clReleaseContext(context);
+ 
+    //release host memory
+    free(h_v);
+ 
+    return 0;
+}

example04/README.md

@@ -0,0 +1,26 @@
+# 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
+sudo make install
+export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/usr/local/lib64
+```
+
+## Running it
+In the top-level directory, run
+
+```
+make ex04
+./example04/bin/Example
+```
+
+and it should print out a vector! :hamburger:
+

example04/main.c

@@ -0,0 +1,142 @@
+#include <stdio.h>
+#include <stdlib.h>
+#include <math.h>
+#include <clFFT.h>
+ 
+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);
+
+    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<N; i++)
+        printf("(%f, %f)  ", h_vR[i], h_vI[i]);
+    printf("]\n");
+
+    // release clFFT stuff
+    clfftDestroyPlan( &planHandleForward );
+    clfftDestroyPlan( &planHandleBackward );
+    clfftTeardown();
+ 
+    // release OpenCL resources
+    clReleaseMemObject(d_vR);
+    clReleaseMemObject(d_vI);
+    clReleaseProgram(program);
+    clReleaseKernel(k_mult);
+    clReleaseCommandQueue(queue);
+    clReleaseContext(context);
+ 
+    //release host memory
+    free(h_vR);
+    free(h_vI);
+ 
+    return 0;
+}

example04/main.py

@@ -0,0 +1,13 @@
+import numpy as np
+import scipy.fftpack as fft
+
+N = 128
+v = np.arange(N) + 2*np.arange(N)*1j
+
+v_fft = fft.fft(v)
+
+v_fft_altered = 2 * v_fft
+
+v_final = fft.ifft(v_fft_altered)
+
+print v_final

example05/main.c

@@ -0,0 +1,247 @@
+#include <stdio.h>
+#include <stdlib.h>
+#include <math.h>
+#include <clFFT.h>
+#include <fftw3.h>
+
+
+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" \
+"    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);
+}
+
+
+void checkIfArraysEqual(double *h_v, double *v, int N, double epsilon) {
+    int arrays_equal = 1;
+    int i;
+
+    for (i=0; i<N; i++) {
+        // printf("[%f %f]  ", h_v[i], v[i]);
+        if (abs(v[i] - h_v[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
+     * 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 = 4096;
+    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;
+    h_v = (double*) malloc(N_bytes);
+ 
+    // initialize v on host (GPU and CPU)
+    int i;
+    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<N; i++)
+        v[i] = i;
+
+    fftw_execute(fft);
+    for (i=0; i<N_COMPLEX; i++) {
+        V[i][REAL] = 2 * V[i][REAL];
+        V[i][IMAG] = 4 * V[i][IMAG];
+    }
+    fftw_execute(ifft);
+
+    // scale array as FFTW doesn't automatically do this for back transform
+    for (i=0; i<N; i++)
+        v[i] = v[i]/N;
+
+
+    // GPU STUFF --------------------------------------------------------------
+    // 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;
+    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);
+
+    // REAL IN-PLACE TRANSFORM ------------------------------------------------
+    // 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);
+
+    err  = clSetKernelArg(k_mult, 0, sizeof(cl_mem), &d_v);
+ 
+    // 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);
+    clFinish(queue);
+    clEnqueueReadBuffer(queue, d_v, CL_TRUE, 0, N_bytes, h_v, 0, NULL, NULL );
+    clFinish(queue);
+    
+    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);
+    clfftDestroyPlan( &planHandleForward );
+    clfftDestroyPlan( &planHandleBackward );
+    clfftTeardown();
+ 
+    // release OpenCL resources
+    clReleaseMemObject(d_v);
+    clReleaseProgram(program);
+    clReleaseKernel(k_mult);
+    clReleaseCommandQueue(queue);
+    clReleaseContext(context);
+ 
+    //release host memory
+    free(v);
+    free(h_v);
+ 
+    return 0;
+}

example05/main.py

@@ -0,0 +1,13 @@
+import numpy as np
+import scipy.fftpack as fft
+
+N = 2048 
+v = np.arange(N)
+
+v_fft = fft.fft(v)
+
+v_fft_altered = 2*v_fft.real + 4*v_fft.imag*1j
+
+v_final = fft.ifft(v_fft_altered)
+print v_final
+