[Beignet] [PATCH 2/2] change built-in function name from get_sub_group_id to get_sub_group_local_id

Guo Yejun yejun.guo at intel.com
Tue Dec 22 16:17:48 PST 2015


Fix bug at https://bugs.freedesktop.org/show_bug.cgi?id=93469

The fucntion is mapped to OP_SIMD_ID which returns the SIMD lane ID.
However, the SIMD lane ID is the equivalent of get_sub_group_local_id().

contributor: Georg Kolling <georg.kolling at gmail.com>
Signed-off-by: Guo Yejun <yejun.guo at intel.com>
---
 backend/src/libocl/tmpl/ocl_simd.tmpl.h    |  2 +-
 backend/src/llvm/llvm_gen_ocl_function.hxx |  2 +-
 kernels/compiler_get_sub_group_id.cl       |  8 --------
 kernels/compiler_get_sub_group_local_id.cl |  8 ++++++++
 kernels/compiler_sub_group_shuffle.cl      |  4 ++--
 utests/CMakeLists.txt                      |  2 +-
 utests/compiler_get_sub_group_id.cpp       | 33 ------------------------------
 utests/compiler_get_sub_group_local_id.cpp | 33 ++++++++++++++++++++++++++++++
 8 files changed, 46 insertions(+), 46 deletions(-)
 delete mode 100644 kernels/compiler_get_sub_group_id.cl
 create mode 100644 kernels/compiler_get_sub_group_local_id.cl
 delete mode 100644 utests/compiler_get_sub_group_id.cpp
 create mode 100644 utests/compiler_get_sub_group_local_id.cpp

diff --git a/backend/src/libocl/tmpl/ocl_simd.tmpl.h b/backend/src/libocl/tmpl/ocl_simd.tmpl.h
index 4055070..9d9404b 100644
--- a/backend/src/libocl/tmpl/ocl_simd.tmpl.h
+++ b/backend/src/libocl/tmpl/ocl_simd.tmpl.h
@@ -27,7 +27,7 @@ int sub_group_any(int);
 int sub_group_all(int);
 
 uint get_max_sub_group_size(void);
-uint get_sub_group_id(void);
+uint get_sub_group_local_id(void);
 
 OVERLOADABLE float intel_sub_group_shuffle(float x, uint c);
 OVERLOADABLE int intel_sub_group_shuffle(int x, uint c);
diff --git a/backend/src/llvm/llvm_gen_ocl_function.hxx b/backend/src/llvm/llvm_gen_ocl_function.hxx
index 046e1ae..e3d89a3 100644
--- a/backend/src/llvm/llvm_gen_ocl_function.hxx
+++ b/backend/src/llvm/llvm_gen_ocl_function.hxx
@@ -162,7 +162,7 @@ DECL_LLVM_GEN_FUNCTION(SAT_CONV_F16_TO_U32, _Z16convert_uint_satDh)
 DECL_LLVM_GEN_FUNCTION(SIMD_ANY, sub_group_any)
 DECL_LLVM_GEN_FUNCTION(SIMD_ALL, sub_group_all)
 DECL_LLVM_GEN_FUNCTION(SIMD_SIZE, get_max_sub_group_size)
-DECL_LLVM_GEN_FUNCTION(SIMD_ID, get_sub_group_id)
+DECL_LLVM_GEN_FUNCTION(SIMD_ID, get_sub_group_local_id)
 DECL_LLVM_GEN_FUNCTION(SIMD_SHUFFLE, intel_sub_group_shuffle)
 
 DECL_LLVM_GEN_FUNCTION(READ_TM, __gen_ocl_read_tm)
diff --git a/kernels/compiler_get_sub_group_id.cl b/kernels/compiler_get_sub_group_id.cl
deleted file mode 100644
index afaa2a6..0000000
--- a/kernels/compiler_get_sub_group_id.cl
+++ /dev/null
@@ -1,8 +0,0 @@
-__kernel void compiler_get_sub_group_id(global int *dst)
-{
-  int i = get_global_id(0);
-  if (i == 0)
-    dst[0] = get_max_sub_group_size();
-
-  dst[i+1] = get_sub_group_id();
-}
diff --git a/kernels/compiler_get_sub_group_local_id.cl b/kernels/compiler_get_sub_group_local_id.cl
new file mode 100644
index 0000000..0a28285
--- /dev/null
+++ b/kernels/compiler_get_sub_group_local_id.cl
@@ -0,0 +1,8 @@
+__kernel void compiler_get_sub_group_local_id(global int *dst)
+{
+  int i = get_global_id(0);
+  if (i == 0)
+    dst[0] = get_max_sub_group_size();
+
+  dst[i+1] = get_sub_group_local_id();
+}
diff --git a/kernels/compiler_sub_group_shuffle.cl b/kernels/compiler_sub_group_shuffle.cl
index a171faa..322da74 100644
--- a/kernels/compiler_sub_group_shuffle.cl
+++ b/kernels/compiler_sub_group_shuffle.cl
@@ -6,8 +6,8 @@ __kernel void compiler_sub_group_shuffle(global int *dst, int c)
   dst++;
 
   int from = i;
-  int j = get_max_sub_group_size() - get_sub_group_id() - 1;
-  int o0 = get_sub_group_id();
+  int j = get_max_sub_group_size() - get_sub_group_local_id() - 1;
+  int o0 = get_sub_group_local_id();
   int o1 = intel_sub_group_shuffle(from, c);
   int o2 = intel_sub_group_shuffle(from, 5);
   int o3 = intel_sub_group_shuffle(from, j);
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index 2c6aea4..db62e38 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -221,7 +221,7 @@ set (utests_sources
   runtime_alloc_host_ptr_buffer.cpp
   runtime_use_host_ptr_image.cpp
   compiler_get_max_sub_group_size.cpp
-  compiler_get_sub_group_id.cpp
+  compiler_get_sub_group_local_id.cpp
   compiler_sub_group_shuffle.cpp
   builtin_global_linear_id.cpp
   builtin_local_linear_id.cpp
diff --git a/utests/compiler_get_sub_group_id.cpp b/utests/compiler_get_sub_group_id.cpp
deleted file mode 100644
index 0d88d29..0000000
--- a/utests/compiler_get_sub_group_id.cpp
+++ /dev/null
@@ -1,33 +0,0 @@
-#include "utest_helper.hpp"
-
-void compiler_get_sub_group_id(void)
-{
-  const size_t n = 256;
-
-  // Setup kernel and buffers
-  OCL_CREATE_KERNEL("compiler_get_sub_group_id");
-  OCL_CREATE_BUFFER(buf[0], 0, (n+1) * sizeof(int), NULL);
-  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
-
-  globals[0] = n;
-  locals[0] = 16;
-
-  OCL_MAP_BUFFER(0);
-  for (int32_t i = 0; i < (int32_t) (n+1); ++i)
-    ((int*)buf_data[0])[i] = -1;
-  OCL_UNMAP_BUFFER(0);
-
-  // Run the kernel on GPU
-  OCL_NDRANGE(1);
-
-  // Compare
-  OCL_MAP_BUFFER(0);
-  int* dst = (int *)buf_data[0];
-  OCL_ASSERT(8 == dst[0] || 16 == dst[0]);
-  for (int32_t i = 1; i < (int32_t) n; ++i){
-    OCL_ASSERT((i-1) % dst[0] == dst[i]);
-  }
-  OCL_UNMAP_BUFFER(0);
-}
-
-MAKE_UTEST_FROM_FUNCTION(compiler_get_sub_group_id);
diff --git a/utests/compiler_get_sub_group_local_id.cpp b/utests/compiler_get_sub_group_local_id.cpp
new file mode 100644
index 0000000..2df4e9b
--- /dev/null
+++ b/utests/compiler_get_sub_group_local_id.cpp
@@ -0,0 +1,33 @@
+#include "utest_helper.hpp"
+
+void compiler_get_sub_group_local_id(void)
+{
+  const size_t n = 256;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_get_sub_group_local_id");
+  OCL_CREATE_BUFFER(buf[0], 0, (n+1) * sizeof(int), NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+
+  globals[0] = n;
+  locals[0] = 16;
+
+  OCL_MAP_BUFFER(0);
+  for (int32_t i = 0; i < (int32_t) (n+1); ++i)
+    ((int*)buf_data[0])[i] = -1;
+  OCL_UNMAP_BUFFER(0);
+
+  // Run the kernel on GPU
+  OCL_NDRANGE(1);
+
+  // Compare
+  OCL_MAP_BUFFER(0);
+  int* dst = (int *)buf_data[0];
+  OCL_ASSERT(8 == dst[0] || 16 == dst[0]);
+  for (int32_t i = 1; i < (int32_t) n; ++i){
+    OCL_ASSERT((i-1) % dst[0] == dst[i]);
+  }
+  OCL_UNMAP_BUFFER(0);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_get_sub_group_local_id);
-- 
1.9.1



More information about the Beignet mailing list