Merge from 'sycl' to 'sycl-web' (#1)

codeplaysoftware · Nov 11, 2020 · 341d71d · 341d71d
2 parents 0f3c3af + 6e3f244
commit 341d71d
Show file tree

Hide file tree

Showing 4 changed files with 12 additions and 12 deletions.
diff --git a/libclc/ptx-nvidiacl/libspirv/group/collectives.cl b/libclc/ptx-nvidiacl/libspirv/group/collectives.cl
@@ -385,6 +385,10 @@ long __clc__3d_to_linear_local_id(ulong3 id) {
       uint scope, TYPE x, ulong3 local_id) {                                   \
     ulong linear_local_id = __clc__3d_to_linear_local_id(local_id);            \
     return __spirv_GroupBroadcast(scope, x, linear_local_id);                  \
+  }                                                                            \
+  _CLC_DEF _CLC_OVERLOAD _CLC_CONVERGENT TYPE __spirv_GroupBroadcast(          \
+      uint scope, TYPE x, uint local_id) {                                     \
+    return __spirv_GroupBroadcast(scope, x, (ulong)local_id);                  \
   }
 __CLC_GROUP_BROADCAST(char);
 __CLC_GROUP_BROADCAST(uchar);
@@ -411,6 +415,10 @@ _CLC_DECL _CLC_CONVERGENT half
 _Z17__spirv_GroupBroadcastjDF16_Dv3_m(uint scope, half x, ulong3 local_id) {
   return __spirv_GroupBroadcast(scope, x, local_id);
 }
+_CLC_DECL _CLC_CONVERGENT half
+_Z22__spirv_GroupBroadcastjDF16_j(uint scope, half x, uint local_id) {
+  return __spirv_GroupBroadcast(scope, x, (ulong)local_id);
+}
 
 #undef __CLC_GROUP_BROADCAST
 

diff --git a/sycl/test/on-device/sub_group/broadcast.cpp b/sycl/test/on-device/sub_group/broadcast.cpp
@@ -1,6 +1,3 @@
-// XFAIL: cuda
-// CUDA compilation and runtime do not yet support sub-groups.
-
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: env SYCL_DEVICE_TYPE=HOST %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
@@ -19,7 +16,7 @@
 
 int main() {
   queue Queue;
-  if (!core_sg_supported(Queue.get_device())) {
+  if (Queue.get_device().is_host()) {
     std::cout << "Skipping test\n";
     return 0;
   }

diff --git a/sycl/test/on-device/sub_group/broadcast_fp16.cpp b/sycl/test/on-device/sub_group/broadcast_fp16.cpp
@@ -1,6 +1,3 @@
-// XFAIL: cuda
-// CUDA compilation and runtime do not yet support sub-groups.
-
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 
@@ -16,10 +13,11 @@
 
 int main() {
   queue Queue;
-  if (!core_sg_supported(Queue.get_device())) {
+  if (Queue.get_device().is_host()) {
     std::cout << "Skipping test\n";
     return 0;
   }
   check<cl::sycl::half>(Queue);
+  std::cout << "Test passed." << std::endl;
   return 0;
 }
diff --git a/sycl/test/on-device/sub_group/broadcast_fp64.cpp b/sycl/test/on-device/sub_group/broadcast_fp64.cpp
@@ -1,6 +1,3 @@
-// XFAIL: cuda
-// CUDA compilation and runtime do not yet support sub-groups.
-
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: env SYCL_DEVICE_TYPE=HOST %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
@@ -19,7 +16,7 @@
 
 int main() {
   queue Queue;
-  if (!core_sg_supported(Queue.get_device())) {
+  if (Queue.get_device().is_host()) {
     std::cout << "Skipping test\n";
     return 0;
   }