Added more tests for function arguments

8 files changed, 180 insertions, 15 deletions

diff --git a/kernels/compiler_function_argument.cl b/kernels/compiler_function_argument.cl
new file mode 100644
index 00000000..fe6de28f
--- /dev/null
+++ b/kernels/compiler_function_argument.cl
@@ -0,0 +1,7 @@
+__kernel void
+compiler_function_argument(__global int *dst, int value)
+{
+  int id = (int)get_global_id(0);
+  dst[id] = value;
+}
+
diff --git a/kernels/compiler_function_argument0.cl b/kernels/compiler_function_argument0.cl
new file mode 100644
index 00000000..6bc2e921
--- /dev/null
+++ b/kernels/compiler_function_argument0.cl
@@ -0,0 +1,7 @@
+__kernel void
+compiler_function_argument0(__global int *dst, short value)
+{
+  int id = (int)get_global_id(0);
+  dst[id] = value;
+}
+
diff --git a/kernels/compiler_function_argument1.cl b/kernels/compiler_function_argument1.cl
new file mode 100644
index 00000000..8842b0b4
--- /dev/null
+++ b/kernels/compiler_function_argument1.cl
@@ -0,0 +1,7 @@
+__kernel void
+compiler_function_argument1(__global int *dst, char value, short value0, int value1)
+{
+  int id = (int)get_global_id(0);
+  dst[id] = value + value0 + value1;
+}
+
diff --git a/src/cl_command_queue_gen7.c b/src/cl_command_queue_gen7.c
index 5da418cf..0ed0fe7a 100644
--- a/src/cl_command_queue_gen7.c
+++ b/src/cl_command_queue_gen7.c
@@ -102,7 +102,6 @@ error:
 
 static void
 cl_curbe_fill(cl_kernel ker,
-              char *curbe,
               const size_t *global_wk_off,
               const size_t *global_wk_sz,
               const size_t *local_wk_sz)
@@ -110,7 +109,7 @@ cl_curbe_fill(cl_kernel ker,
   int32_t offset;
 #define UPLOAD(ENUM, VALUE) \
   if ((offset = gbe_kernel_get_curbe_offset(ker->opaque, ENUM, 0)) >= 0) \
-    *((uint32_t *) (curbe + offset)) = VALUE;
+    *((uint32_t *) (ker->curbe + offset)) = VALUE;
   UPLOAD(GBE_CURBE_LOCAL_SIZE_X, local_wk_sz[0]);
   UPLOAD(GBE_CURBE_LOCAL_SIZE_Y, local_wk_sz[1]);
   UPLOAD(GBE_CURBE_LOCAL_SIZE_Z, local_wk_sz[2]);
@@ -125,12 +124,12 @@ cl_curbe_fill(cl_kernel ker,
   UPLOAD(GBE_CURBE_GROUP_NUM_Z, global_wk_sz[2]/local_wk_sz[2]);
 #undef UPLOAD
 
-  /* Write identity for the stack pointer. This will make the stack pointer
-   * computations faster in the kernel
+  /* Write identity for the stack pointer. This is required by the stack pointer
+   * computation in the kernel
    */
   if ((offset = gbe_kernel_get_curbe_offset(ker->opaque, GBE_CURBE_STACK_POINTER, 0)) >= 0) {
     const uint32_t simd_sz = gbe_kernel_get_simd_width(ker->opaque);
-    uint32_t *stackptr = (uint32_t *) (curbe + offset);
+    uint32_t *stackptr = (uint32_t *) (ker->curbe + offset);
     int32_t i;
     for (i = 0; i < simd_sz; ++i) stackptr[i] = i;
   }
@@ -170,7 +169,6 @@ cl_command_queue_ND_range_gen7(cl_command_queue queue,
 {
   cl_context ctx = queue->ctx;
   cl_gpgpu gpgpu = queue->gpgpu;
-  char *curbe = NULL;        /* Does not include per-thread local IDs */
   char *final_curbe = NULL;  /* Includes them and one sub-buffer per group */
   cl_gpgpu_kernel kernel;
   const uint32_t simd_sz = cl_kernel_get_simd_width(ker);
@@ -187,8 +185,8 @@ cl_command_queue_ND_range_gen7(cl_command_queue queue,
   kernel.slm_sz = 0;
 
   /* Curbe step 1: fill the constant buffer data shared by all threads */
-  curbe = alloca(ker->curbe_sz);
-  cl_curbe_fill(ker, curbe, global_wk_off, global_wk_sz, local_wk_sz);
+  if (ker->curbe)
+    cl_curbe_fill(ker, global_wk_off, global_wk_sz, local_wk_sz);
 
   /* Compute the number of HW threads we need */
   TRY (cl_kernel_work_group_sz, ker, local_wk_sz, 3, &local_sz);
@@ -209,12 +207,14 @@ cl_command_queue_ND_range_gen7(cl_command_queue queue,
   cl_gpgpu_states_setup(gpgpu, &kernel);
 
   /* Curbe step 2. Give the localID and upload it to video memory */
-  TRY_ALLOC (final_curbe, (char*) alloca(thread_n * cst_sz));
-  if (curbe)
-    for (i = 0; i < thread_n; ++i)
-      memcpy(final_curbe + cst_sz * i, curbe, cst_sz);
-  TRY (cl_set_varying_payload, ker, final_curbe, local_wk_sz, simd_sz, cst_sz, thread_n);
-  cl_gpgpu_upload_constants(gpgpu, final_curbe, thread_n*cst_sz);
+  if (ker->curbe) {
+    assert(cst_sz > 0);
+    TRY_ALLOC (final_curbe, (char*) alloca(thread_n * cst_sz));
+      for (i = 0; i < thread_n; ++i)
+        memcpy(final_curbe + cst_sz * i, ker->curbe, cst_sz);
+    TRY (cl_set_varying_payload, ker, final_curbe, local_wk_sz, simd_sz, cst_sz, thread_n);
+    cl_gpgpu_upload_constants(gpgpu, final_curbe, thread_n*cst_sz);
+  }
 
   /* Start a new batch buffer */
   batch_sz = cl_kernel_compute_batch_sz(ker);
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index 0fd5e823..f8b3805d 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -29,7 +29,10 @@ ADD_LIBRARY(utests SHARED
             compiler_lower_return1.cpp
             compiler_lower_return2.cpp
             compiler_array.cpp
-            compiler_array0.cpp)
+            compiler_array0.cpp
+            compiler_function_argument.cpp
+            compiler_function_argument0.cpp
+            compiler_function_argument1.cpp)
 TARGET_LINK_LIBRARIES(utests cl m)
 
 ADD_EXECUTABLE(run utest_run.cpp)
diff --git a/utests/compiler_function_argument.cpp b/utests/compiler_function_argument.cpp
new file mode 100644
index 00000000..29775ceb
--- /dev/null
+++ b/utests/compiler_function_argument.cpp
@@ -0,0 +1,46 @@
+/* 
+ * Copyright © 2012 Intel Corporation
+ *
+ * This library is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation; either
+ * version 2 of the License, or (at your option) any later version.
+ *
+ * This library is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with this library. If not, see <http://www.gnu.org/licenses/>.
+ *
+ * Author: Benjamin Segovia <benjamin.segovia@intel.com>
+ */
+
+#include "utest_helper.hpp"
+
+void compiler_function_argument(void)
+{
+  const size_t n = 2048;
+  const int value = 34;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_function_argument");
+  OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, sizeof(int), &value);
+
+  // Run the kernel
+  globals[0] = n;
+  locals[0] = 16;
+  OCL_NDRANGE(1);
+  OCL_MAP_BUFFER(0);
+
+  // Check results
+  for (uint32_t i = 0; i < n; ++i)
+    OCL_ASSERT(((int*)buf_data[0])[i] == value);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_function_argument);
+
+
diff --git a/utests/compiler_function_argument0.cpp b/utests/compiler_function_argument0.cpp
new file mode 100644
index 00000000..3aac9cb7
--- /dev/null
+++ b/utests/compiler_function_argument0.cpp
@@ -0,0 +1,45 @@
+/* 
+ * Copyright © 2012 Intel Corporation
+ *
+ * This library is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation; either
+ * version 2 of the License, or (at your option) any later version.
+ *
+ * This library is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with this library. If not, see <http://www.gnu.org/licenses/>.
+ *
+ * Author: Benjamin Segovia <benjamin.segovia@intel.com>
+ */
+
+#include "utest_helper.hpp"
+
+void compiler_function_argument0(void)
+{
+  const size_t n = 2048;
+  const short value = 34;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_function_argument0");
+  OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, sizeof(short), &value);
+
+  // Run the kernel
+  globals[0] = n;
+  locals[0] = 16;
+  OCL_NDRANGE(1);
+  OCL_MAP_BUFFER(0);
+
+  // Check results
+  for (uint32_t i = 0; i < n; ++i)
+    OCL_ASSERT(((int*)buf_data[0])[i] == value);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_function_argument0);
+
diff --git a/utests/compiler_function_argument1.cpp b/utests/compiler_function_argument1.cpp
new file mode 100644
index 00000000..dff46a81
--- /dev/null
+++ b/utests/compiler_function_argument1.cpp
@@ -0,0 +1,50 @@
+/* 
+ * Copyright © 2012 Intel Corporation
+ *
+ * This library is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation; either
+ * version 2 of the License, or (at your option) any later version.
+ *
+ * This library is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with this library. If not, see <http://www.gnu.org/licenses/>.
+ *
+ * Author: Benjamin Segovia <benjamin.segovia@intel.com>
+ */
+
+#include "utest_helper.hpp"
+
+void compiler_function_argument1(void)
+{
+  const size_t n = 2048;
+  const char value = 34;
+  const short value0 = 31;
+  const int value1 = 3;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_function_argument1");
+  OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, sizeof(char), &value);
+  OCL_SET_ARG(2, sizeof(short), &value0);
+  OCL_SET_ARG(3, sizeof(int), &value1);
+
+  // Run the kernel
+  globals[0] = n;
+  locals[0] = 16;
+  OCL_NDRANGE(1);
+  OCL_MAP_BUFFER(0);
+
+  // Check results
+  for (uint32_t i = 0; i < n; ++i)
+    OCL_ASSERT(((int*)buf_data[0])[i] == value + value0 + value1);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_function_argument1);
+
+