rename __gen_ocl_get_simd_id/size to get_sub_group_id/size

Signed-off-by: Guo Yejun <yejun.guo@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Signed-off-by: Zhigang Gong <zhigang.gong@intel.com>

10 files changed, 31 insertions, 26 deletions

diff --git a/backend/src/libocl/tmpl/ocl_simd.tmpl.h b/backend/src/libocl/tmpl/ocl_simd.tmpl.h
index 620e329a..14e5750c 100644
--- a/backend/src/libocl/tmpl/ocl_simd.tmpl.h
+++ b/backend/src/libocl/tmpl/ocl_simd.tmpl.h
@@ -24,5 +24,9 @@
 // SIMD level function
 /////////////////////////////////////////////////////////////////////////////
 
-uint __gen_ocl_get_simd_size(void);
-uint __gen_ocl_get_simd_id(void);
+uint get_sub_group_size(void);
+uint get_sub_group_id(void);
+
+OVERLOADABLE float intel_sub_group_shuffle(float x, uint c);
+OVERLOADABLE int intel_sub_group_shuffle(int x, uint c);
+OVERLOADABLE uint intel_sub_group_shuffle(uint x, uint c);
diff --git a/backend/src/llvm/llvm_gen_ocl_function.hxx b/backend/src/llvm/llvm_gen_ocl_function.hxx
index e2bffdea..a0e0b945 100644
--- a/backend/src/llvm/llvm_gen_ocl_function.hxx
+++ b/backend/src/llvm/llvm_gen_ocl_function.hxx
@@ -154,8 +154,9 @@ DECL_LLVM_GEN_FUNCTION(CONV_F32_TO_F16, __gen_ocl_f32to16)
 // SIMD level function for internal usage
 DECL_LLVM_GEN_FUNCTION(SIMD_ANY, __gen_ocl_simd_any)
 DECL_LLVM_GEN_FUNCTION(SIMD_ALL, __gen_ocl_simd_all)
-DECL_LLVM_GEN_FUNCTION(SIMD_SIZE, __gen_ocl_get_simd_size)
-DECL_LLVM_GEN_FUNCTION(SIMD_ID, __gen_ocl_get_simd_id)
+DECL_LLVM_GEN_FUNCTION(SIMD_SIZE, get_sub_group_size)
+DECL_LLVM_GEN_FUNCTION(SIMD_ID, get_sub_group_id)
+DECL_LLVM_GEN_FUNCTION(SIMD_SHUFFLE, intel_sub_group_shuffle)
 
 DECL_LLVM_GEN_FUNCTION(READ_TM, __gen_ocl_read_tm)
 DECL_LLVM_GEN_FUNCTION(REGION, __gen_ocl_region)
diff --git a/kernels/compiler_get_simd_id.cl b/kernels/compiler_get_simd_id.cl
deleted file mode 100644
index dfe625a8..00000000
--- a/kernels/compiler_get_simd_id.cl
+++ /dev/null
@@ -1,8 +0,0 @@
-__kernel void compiler_get_simd_id(global int *dst)
-{
-  int i = get_global_id(0);
-  if (i == 0)
-    dst[0] = __gen_ocl_get_simd_size();
-
-  dst[i+1] = __gen_ocl_get_simd_id();
-}
diff --git a/kernels/compiler_get_simd_size.cl b/kernels/compiler_get_simd_size.cl
deleted file mode 100644
index 6e303a35..00000000
--- a/kernels/compiler_get_simd_size.cl
+++ /dev/null
@@ -1,5 +0,0 @@
-__kernel void compiler_get_simd_size(global int *dst)
-{
-  int i = get_global_id(0);
-  dst[i] = __gen_ocl_get_simd_size();
-}
diff --git a/kernels/compiler_get_sub_group_id.cl b/kernels/compiler_get_sub_group_id.cl
new file mode 100644
index 00000000..10033ff0
--- /dev/null
+++ b/kernels/compiler_get_sub_group_id.cl
@@ -0,0 +1,8 @@
+__kernel void compiler_get_sub_group_id(global int *dst)
+{
+  int i = get_global_id(0);
+  if (i == 0)
+    dst[0] = get_sub_group_size();
+
+  dst[i+1] = get_sub_group_id();
+}
diff --git a/kernels/compiler_get_sub_group_size.cl b/kernels/compiler_get_sub_group_size.cl
new file mode 100644
index 00000000..4d5e3ebc
--- /dev/null
+++ b/kernels/compiler_get_sub_group_size.cl
@@ -0,0 +1,5 @@
+__kernel void compiler_get_sub_group_size(global int *dst)
+{
+  int i = get_global_id(0);
+  dst[i] = get_sub_group_size();
+}
diff --git a/src/cl_command_queue_gen7.c b/src/cl_command_queue_gen7.c
index e27a2112..89f39b3c 100644
--- a/src/cl_command_queue_gen7.c
+++ b/src/cl_command_queue_gen7.c
@@ -210,7 +210,7 @@ cl_curbe_fill(cl_kernel ker,
   UPLOAD(GBE_CURBE_WORK_DIM, work_dim);
 #undef UPLOAD
 
-  /* __gen_ocl_get_simd_id needs it */
+  /* get_sub_group_id needs it */
   if ((offset = interp_kernel_get_curbe_offset(ker->opaque, GBE_CURBE_LANE_ID, 0)) >= 0) {
     const uint32_t simd_sz = interp_kernel_get_simd_width(ker->opaque);
     uint32_t *laneid = (uint32_t *) (ker->curbe + offset);
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index dcb33850..7c0c596c 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -209,8 +209,8 @@ set (utests_sources
   vload_bench.cpp
   runtime_use_host_ptr_buffer.cpp
   runtime_alloc_host_ptr_buffer.cpp
-  compiler_get_simd_size.cpp
-  compiler_get_simd_id.cpp)
+  compiler_get_sub_group_size.cpp
+  compiler_get_sub_group_id.cpp)
 
 if (LLVM_VERSION_NODOT VERSION_GREATER 34)
   SET(utests_sources
diff --git a/utests/compiler_get_simd_id.cpp b/utests/compiler_get_sub_group_id.cpp
index ad10bf75..0d88d293 100644
--- a/utests/compiler_get_simd_id.cpp
+++ b/utests/compiler_get_sub_group_id.cpp
@@ -1,11 +1,11 @@
 #include "utest_helper.hpp"
 
-void compiler_get_simd_id(void)
+void compiler_get_sub_group_id(void)
 {
   const size_t n = 256;
 
   // Setup kernel and buffers
-  OCL_CREATE_KERNEL("compiler_get_simd_id");
+  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]);
 
@@ -30,4 +30,4 @@ void compiler_get_simd_id(void)
   OCL_UNMAP_BUFFER(0);
 }
 
-MAKE_UTEST_FROM_FUNCTION(compiler_get_simd_id);
+MAKE_UTEST_FROM_FUNCTION(compiler_get_sub_group_id);
diff --git a/utests/compiler_get_simd_size.cpp b/utests/compiler_get_sub_group_size.cpp
index ea70cd96..20339d72 100644
--- a/utests/compiler_get_simd_size.cpp
+++ b/utests/compiler_get_sub_group_size.cpp
@@ -1,11 +1,11 @@
 #include "utest_helper.hpp"
 
-void compiler_get_simd_size(void)
+void compiler_get_sub_group_size(void)
 {
   const size_t n = 256;
 
   // Setup kernel and buffers
-  OCL_CREATE_KERNEL("compiler_get_simd_size");
+  OCL_CREATE_KERNEL("compiler_get_sub_group_size");
   OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(int), NULL);
   OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
 
@@ -29,4 +29,4 @@ void compiler_get_simd_size(void)
   OCL_UNMAP_BUFFER(0);
 }
 
-MAKE_UTEST_FROM_FUNCTION(compiler_get_simd_size);
+MAKE_UTEST_FROM_FUNCTION(compiler_get_sub_group_size);