From 3129ddfbf920434637c66ceb16b256ae9eccfc5f Mon Sep 17 00:00:00 2001
From: "Dakota St. Laurent" <d.h.stlaurent@gmail.com>
Date: Mon, 3 Aug 2015 06:50:31 -0400
Subject: [PATCH] add example03 for using fp64 extension

---
 .gitignore       |   5 +++
 example03/main.c | 109 +++++++++++++++++++++++++++++++++++++++++++++++
 2 files changed, 114 insertions(+)
 create mode 100644 example03/main.c

diff --git a/.gitignore b/.gitignore
index b4e66d2..8606aaa 100644
--- a/.gitignore
+++ b/.gitignore
@@ -1,5 +1,10 @@
 # compiled files
 *.out
+**/bin
+**/build
+
+# clFFT library
+clFFT
 
 # THIS IS MY SWAP
 # (vim swap files)
diff --git a/example03/main.c b/example03/main.c
new file mode 100644
index 0000000..04f4baf
--- /dev/null
+++ b/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;
+}