Fix bugs exposed in CI

 - implement fallback for CUDA BE;
 - fix assert function name on Windows.

Author: vladimirlaz

sycl/include/CL/sycl/ONEAPI/sub_group.hpp

@@ -241,6 +241,9 @@ struct sub_group {
       std::is_same<typename detail::remove_AS<T>::type, T>::value, T>
   load(T *src) const {
 
+#ifdef __NVPTX__
+    return src[get_local_id()[0]];
+#else // __NVPTX__
     auto l = __spirv_GenericCastToPtrExplicit_ToLocal<T>(
         src, __spv::StorageClass::Workgroup);
     if (l)
@@ -258,6 +261,7 @@ struct sub_group {
 
     // Fallback for other address spaces to be mapped to global
     return load(__spirv_PtrCastToGeneric<T>(src));
+#endif // __NVPTX__
   }
 #else  //__SYCL_DEVICE_ONLY__
   template <typename T> T load(T *src) const {
@@ -375,6 +379,9 @@ struct sub_group {
       std::is_same<typename detail::remove_AS<T>::type, T>::value>
   store(T *dst, const typename detail::remove_AS<T>::type &x) const {
 
+#ifdef __NVPTX__
+    dst[get_local_id()[0]] = x;
+#else // __NVPTX__
     auto l = __spirv_GenericCastToPtrExplicit_ToLocal<T>(
         dst, __spv::StorageClass::Workgroup);
     if (l) {
@@ -396,6 +403,7 @@ struct sub_group {
 
     // Fallback for other address spaces to be mapped to global
     store(__spirv_PtrCastToGeneric<T>(dst), x);
+#endif // __NVPTX__
   }
 #else  //__SYCL_DEVICE_ONLY__
   template <typename T> void store(T *dst, const T &x) const {

sycl/test/extensions/sub_group_as.cpp

@@ -61,7 +61,7 @@ int main(int argc, char *argv[]) {
             // CHECK: call spir_func i8 addrspace(1)* @_Z41__spirv_GenericCastToPtrExplicit_ToGlobalPKvN5__spv12StorageClass4FlagE(i8 addrspace(4)*
             // CHECK: call spir_func i32 @_Z30__spirv_SubgroupBlockReadINTELIjET_PU3AS1Kj(i32 addrspace(1)*
             // CHECK: call spir_func i8* @_Z42__spirv_GenericCastToPtrExplicit_ToPrivatePKvN5__spv12StorageClass4FlagE(i8 addrspace(4)*
-            // CHECK: call spir_func void @__assert_fail
+            // CHECK: call spir_func void {{.*}}assert
             // CHECK: call spir_func i8 addrspace(4)* @_Z24__spirv_PtrCastToGenericPKv(i8 addrspace(4)*
             // CHECK: call spir_func i32 @_Z30__spirv_SubgroupBlockReadINTELIjET_PU3AS1Kj(i32 addrspace(1)*
             // Global address space
@@ -78,7 +78,7 @@ int main(int argc, char *argv[]) {
             // CHECK: call spir_func i8 addrspace(1)* @_Z41__spirv_GenericCastToPtrExplicit_ToGlobalPKvN5__spv12StorageClass4FlagE(i8 addrspace(4)*
             // CHECK: call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELIjEvPU3AS1jT_(i32 addrspace(1)*
             // CHECK: call spir_func i8* @_Z42__spirv_GenericCastToPtrExplicit_ToPrivatePKvN5__spv12StorageClass4FlagE(i8 addrspace(4)*
-            // CHECK: call spir_func void @__assert_fail
+            // CHECK: call spir_func void {{.*}}assert
             // CHECK: call spir_func i8 addrspace(4)* @_Z24__spirv_PtrCastToGenericPKv(i8 addrspace(4)*
             // CHECK: call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELIjEvPU3AS1jT_(i32 addrspace(1)*
             sg.store(&global[i], x + y + z);