From f3956a422e42fc76f2eb2f2c2473f5bd188501d9 Mon Sep 17 00:00:00 2001
From: Pradeep <pradeep@arrayfire.com>
Date: Fri, 6 Nov 2015 13:13:53 -0500
Subject: [PATCH] More fixes for Windows CUDA examples

For some stupid reason, plot3, plotting and surface CUDA examples
are not compiling on WINDOWS stating CUDA kernel has no visiblity
to the global constant variables.
---
 examples/cuda/plot3.cu    |  4 ++--
 examples/cuda/plotting.cu | 12 ++++++------
 examples/cuda/surface.cu  | 18 ++++++++++--------
 3 files changed, 18 insertions(+), 16 deletions(-)

diff --git a/examples/cuda/plot3.cu b/examples/cuda/plot3.cu
index 84239b16..2b40682f 100644
--- a/examples/cuda/plot3.cu
+++ b/examples/cuda/plot3.cu
@@ -88,7 +88,7 @@ int main(void)
 
 
 __global__
-void gen_curve(float t, float dx, float* out)
+void gen_curve(float t, float dx, float* out, const float ZMIN, const size_t ZSIZE)
 {
     int offset = blockIdx.x * blockDim.x  + threadIdx.x;
 
@@ -110,5 +110,5 @@ void kernel(float t, float dx, float* dev_out)
     static const dim3 threads(1024);
     dim3 blocks(divup(ZSIZE, 1024));
 
-    gen_curve<<< blocks, threads >>>(t, dx, dev_out);
+    gen_curve<<< blocks, threads >>>(t, dx, dev_out, ZMIN, ZSIZE);
 }
diff --git a/examples/cuda/plotting.cu b/examples/cuda/plotting.cu
index 203317b9..74a56a29 100644
--- a/examples/cuda/plotting.cu
+++ b/examples/cuda/plotting.cu
@@ -10,10 +10,10 @@ const unsigned DIMY = 800;
 const unsigned WIN_ROWS = 2;
 const unsigned WIN_COLS = 2;
 
-const float    dx = 0.1;
-const float    FRANGE_START = 0.f;
-const float    FRANGE_END = 2 * 3.141592f;
-const size_t   DATA_SIZE = ( FRANGE_END - FRANGE_START ) / dx;
+static const float    dx = 0.1;
+static const float    FRANGE_START = 0.f;
+static const float    FRANGE_END = 2 * 3.141592f;
+static const size_t   DATA_SIZE = ( FRANGE_END - FRANGE_START ) / dx;
 
 void kernel(float* dev_out);
 
@@ -97,7 +97,7 @@ int main(void)
 
 
 __global__
-void simple_sinf(float* out)
+void simple_sinf(float* out, const size_t DATA_SIZE, const float dx)
 {
     int x = blockIdx.x * blockDim.x  + threadIdx.x;
 
@@ -112,5 +112,5 @@ void kernel(float* dev_out)
     static const dim3 threads(DATA_SIZE);
     dim3 blocks(1);
 
-    simple_sinf<<< blocks, threads >>>(dev_out);
+	simple_sinf << < blocks, threads >> >(dev_out, DATA_SIZE, dx);
 }
diff --git a/examples/cuda/surface.cu b/examples/cuda/surface.cu
index 3904102f..4ea89d03 100644
--- a/examples/cuda/surface.cu
+++ b/examples/cuda/surface.cu
@@ -8,10 +8,10 @@
 const unsigned DIMX = 1000;
 const unsigned DIMY = 800;
 
-static const float XMIN = -1.0f;
-static const float XMAX = 2.f;
-static const float YMIN = -1.0f;
-static const float YMAX = 1.f;
+const float XMIN = -1.0f;
+const float XMAX = 2.f;
+const float YMIN = -1.0f;
+const float YMAX = 1.f;
 
 const float DX = 0.01;
 const size_t XSIZE = (XMAX-XMIN)/DX+1;
@@ -86,13 +86,15 @@ int main(void)
 
 
 __global__
-void sincos_surf(float t, float dx, float* out)
+void sincos_surf(float t, float dx, float* out,
+				 const float XMIN, const float YMIN,
+				 const size_t XSIZE, const size_t YSIZE)
 {
     int i = blockIdx.x * blockDim.x  + threadIdx.x;
     int j = blockIdx.y * blockDim.y  + threadIdx.y;
 
-    float x=XMIN+i*dx;
-    float y=YMIN+j*dx;
+    float x= ::XMIN + i*dx;
+    float y= ::YMIN + j*dx;
     if (i<XSIZE && j<YSIZE) {
         int offset = j + i * YSIZE;
         out[ 3 * offset     ] = x;
@@ -112,5 +114,5 @@ void kernel(float t, float dx, float* dev_out)
     dim3 blocks(divup(XSIZE, threads.x),
                 divup(YSIZE, threads.y));
 
-    sincos_surf<<< blocks, threads >>>(t, dx, dev_out);
+    sincos_surf<<< blocks, threads >>>(t, dx, dev_out, XMIN, YMIN, XSIZE, YSIZE);
 }