diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
index 6acbfd2024e46..820ce503385ae 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -9188,7 +9188,8 @@ void SPIRVTranslator::ConstructJob(Compilation &C, const JobAction &JA,
       ExtArg += ",+SPV_INTEL_token_type"
                 ",+SPV_INTEL_bfloat16_conversion"
                 ",+SPV_INTEL_joint_matrix"
-                ",+SPV_INTEL_hw_thread_queries";
+                ",+SPV_INTEL_hw_thread_queries"
+                ",+SPV_KHR_uniform_group_instructions";
     TranslatorArgs.push_back(TCArgs.MakeArgString(ExtArg));
   }
   for (auto I : Inputs) {
diff --git a/clang/lib/Sema/SPIRVBuiltins.td b/clang/lib/Sema/SPIRVBuiltins.td
index 30872777e3102..7eab0e1926ced 100644
--- a/clang/lib/Sema/SPIRVBuiltins.td
+++ b/clang/lib/Sema/SPIRVBuiltins.td
@@ -929,13 +929,13 @@ foreach name = ["GroupBroadcast"] in {
   }
 }
 
-foreach name = ["GroupIAdd", "GroupNonUniformIMul", "GroupNonUniformBitwiseOr",
-                "GroupNonUniformBitwiseXor", "GroupNonUniformBitwiseAnd"] in {
+foreach name = ["GroupIAdd", "GroupIMulKHR", "GroupBitwiseOrKHR",
+                "GroupBitwiseXorKHR", "GroupBitwiseAndKHR"] in {
   def : SPVBuiltin<name, [AIGenTypeN, UInt, UInt, AIGenTypeN], Attr.Convergent>;
 }
 
 foreach name = ["GroupFAdd", "GroupFMin", "GroupFMax",
-                "GroupNonUniformFMul"] in {
+                "GroupFMulKHR"] in {
   def : SPVBuiltin<name, [FGenTypeN, UInt, UInt, FGenTypeN], Attr.Convergent>;
 }
 
diff --git a/clang/test/Driver/sycl-spirv-ext.c b/clang/test/Driver/sycl-spirv-ext.c
index 7f0def659e969..fc854a5aaa484 100644
--- a/clang/test/Driver/sycl-spirv-ext.c
+++ b/clang/test/Driver/sycl-spirv-ext.c
@@ -47,7 +47,8 @@
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_token_type
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_bfloat16_conversion
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_joint_matrix
-// CHECK-DEFAULT-SAME:,+SPV_INTEL_hw_thread_queries"
+// CHECK-DEFAULT-SAME:,+SPV_INTEL_hw_thread_queries
+// CHECK-DEFAULT-SAME:,+SPV_KHR_uniform_group_instructions"
 // CHECK-FPGA-HW: llvm-spirv{{.*}}"-spirv-ext=-all
 // CHECK-FPGA-HW-SAME:,+SPV_EXT_shader_atomic_float_add
 // CHECK-FPGA-HW-SAME:,+SPV_EXT_shader_atomic_float_min_max
diff --git a/sycl/include/sycl/ext/oneapi/functional.hpp b/sycl/include/sycl/ext/oneapi/functional.hpp
index 46a5a18fd27ac..18b0c23cac74c 100644
--- a/sycl/include/sycl/ext/oneapi/functional.hpp
+++ b/sycl/include/sycl/ext/oneapi/functional.hpp
@@ -82,16 +82,16 @@ __SYCL_CALC_OVERLOAD(GroupOpISigned, IAdd, sycl::plus<T>)
 __SYCL_CALC_OVERLOAD(GroupOpIUnsigned, IAdd, sycl::plus<T>)
 __SYCL_CALC_OVERLOAD(GroupOpFP, FAdd, sycl::plus<T>)
 
-__SYCL_CALC_OVERLOAD(GroupOpISigned, NonUniformIMul, sycl::multiplies<T>)
-__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, NonUniformIMul, sycl::multiplies<T>)
-__SYCL_CALC_OVERLOAD(GroupOpFP, NonUniformFMul, sycl::multiplies<T>)
-
-__SYCL_CALC_OVERLOAD(GroupOpISigned, NonUniformBitwiseOr, sycl::bit_or<T>)
-__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, NonUniformBitwiseOr, sycl::bit_or<T>)
-__SYCL_CALC_OVERLOAD(GroupOpISigned, NonUniformBitwiseXor, sycl::bit_xor<T>)
-__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, NonUniformBitwiseXor, sycl::bit_xor<T>)
-__SYCL_CALC_OVERLOAD(GroupOpISigned, NonUniformBitwiseAnd, sycl::bit_and<T>)
-__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, NonUniformBitwiseAnd, sycl::bit_and<T>)
+__SYCL_CALC_OVERLOAD(GroupOpISigned, IMulKHR, sycl::multiplies<T>)
+__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, IMulKHR, sycl::multiplies<T>)
+__SYCL_CALC_OVERLOAD(GroupOpFP, FMulKHR, sycl::multiplies<T>)
+
+__SYCL_CALC_OVERLOAD(GroupOpISigned, BitwiseOrKHR, sycl::bit_or<T>)
+__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, BitwiseOrKHR, sycl::bit_or<T>)
+__SYCL_CALC_OVERLOAD(GroupOpISigned, BitwiseXorKHR, sycl::bit_xor<T>)
+__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, BitwiseXorKHR, sycl::bit_xor<T>)
+__SYCL_CALC_OVERLOAD(GroupOpISigned, BitwiseAndKHR, sycl::bit_and<T>)
+__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, BitwiseAndKHR, sycl::bit_and<T>)
 
 #undef __SYCL_CALC_OVERLOAD
 
diff --git a/sycl/test/check_device_code/group_operations.cpp b/sycl/test/check_device_code/group_operations.cpp
index e5ea55bf8f1b1..f0689c436fe48 100644
--- a/sycl/test/check_device_code/group_operations.cpp
+++ b/sycl/test/check_device_code/group_operations.cpp
@@ -54,24 +54,24 @@ template <typename G> [[gnu::always_inline]] void test(G g) {
 SYCL_EXTERNAL void test_group(group<> g) { test(g); }
 
 // int8_t
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjji(i32 2, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjji(i32 2, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjji(i32 2, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjji(i32 2, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjji(i32 2, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjji(i32 2, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjji(i32 2, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjji(i32 2, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjji(i32 2, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjji(i32 2, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjji(i32 2, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjji(i32 2, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjji(i32 2, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjji(i32 2, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjji(i32 2, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjji(i32 2, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjji(i32 2, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjji(i32 2, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxjji(i32 2, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxjji(i32 2, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxjji(i32 2, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMinjji(i32 2, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMinjji(i32 2, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMinjji(i32 2, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljji(i32 2, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljji(i32 2, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljji(i32 2, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjji(i32 2, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjji(i32 2, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjji(i32 2, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjji(i32 2, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjji(i32 2, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjji(i32 2, i32 2, i32 1)
@@ -80,24 +80,24 @@ SYCL_EXTERNAL void test_group(group<> g) { test(g); }
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastji{{m|y}}(i32 2, i32 1, i64 2)
 
 // uint8_t
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjjj(i32 2, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjjj(i32 2, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjjj(i32 2, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjjj(i32 2, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjjj(i32 2, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjjj(i32 2, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjjj(i32 2, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjjj(i32 2, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjjj(i32 2, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjjj(i32 2, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjjj(i32 2, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjjj(i32 2, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjjj(i32 2, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjjj(i32 2, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjjj(i32 2, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjjj(i32 2, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjjj(i32 2, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjjj(i32 2, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxjjj(i32 2, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxjjj(i32 2, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxjjj(i32 2, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMinjjj(i32 2, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMinjjj(i32 2, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMinjjj(i32 2, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljjj(i32 2, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljjj(i32 2, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljjj(i32 2, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjjj(i32 2, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjjj(i32 2, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjjj(i32 2, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjjj(i32 2, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjjj(i32 2, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjjj(i32 2, i32 2, i32 1)
@@ -106,24 +106,24 @@ SYCL_EXTERNAL void test_group(group<> g) { test(g); }
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjj{{m|y}}(i32 2, i32 1, i64 2)
 
 // int16_t
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjji(i32 2, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjji(i32 2, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjji(i32 2, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjji(i32 2, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjji(i32 2, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjji(i32 2, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjji(i32 2, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjji(i32 2, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjji(i32 2, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjji(i32 2, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjji(i32 2, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjji(i32 2, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjji(i32 2, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjji(i32 2, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjji(i32 2, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjji(i32 2, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjji(i32 2, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjji(i32 2, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxjji(i32 2, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxjji(i32 2, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxjji(i32 2, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMinjji(i32 2, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMinjji(i32 2, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMinjji(i32 2, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljji(i32 2, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljji(i32 2, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljji(i32 2, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjji(i32 2, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjji(i32 2, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjji(i32 2, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjji(i32 2, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjji(i32 2, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjji(i32 2, i32 2, i32 1)
@@ -132,24 +132,24 @@ SYCL_EXTERNAL void test_group(group<> g) { test(g); }
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastji{{m|y}}(i32 2, i32 1, i64 2)
 
 // uint16_t
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjjj(i32 2, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjjj(i32 2, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjjj(i32 2, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjjj(i32 2, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjjj(i32 2, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjjj(i32 2, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjjj(i32 2, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjjj(i32 2, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjjj(i32 2, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjjj(i32 2, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjjj(i32 2, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjjj(i32 2, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjjj(i32 2, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjjj(i32 2, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjjj(i32 2, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjjj(i32 2, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjjj(i32 2, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjjj(i32 2, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxjjj(i32 2, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxjjj(i32 2, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxjjj(i32 2, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMinjjj(i32 2, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMinjjj(i32 2, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMinjjj(i32 2, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljjj(i32 2, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljjj(i32 2, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljjj(i32 2, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjjj(i32 2, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjjj(i32 2, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjjj(i32 2, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjjj(i32 2, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjjj(i32 2, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjjj(i32 2, i32 2, i32 1)
@@ -158,24 +158,24 @@ SYCL_EXTERNAL void test_group(group<> g) { test(g); }
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjj{{m|y}}(i32 2, i32 1, i64 2)
 
 // int32_t
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjji(i32 2, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjji(i32 2, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjji(i32 2, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjji(i32 2, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjji(i32 2, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjji(i32 2, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjji(i32 2, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjji(i32 2, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjji(i32 2, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjji(i32 2, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjji(i32 2, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjji(i32 2, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjji(i32 2, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjji(i32 2, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjji(i32 2, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjji(i32 2, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjji(i32 2, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjji(i32 2, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxjji(i32 2, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxjji(i32 2, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxjji(i32 2, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMinjji(i32 2, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMinjji(i32 2, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMinjji(i32 2, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljji(i32 2, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljji(i32 2, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljji(i32 2, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjji(i32 2, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjji(i32 2, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjji(i32 2, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjji(i32 2, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjji(i32 2, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjji(i32 2, i32 2, i32 1)
@@ -184,24 +184,24 @@ SYCL_EXTERNAL void test_group(group<> g) { test(g); }
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastji{{m|y}}(i32 2, i32 1, i64 2)
 
 // uint32_t
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjjj(i32 2, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjjj(i32 2, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjjj(i32 2, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjjj(i32 2, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjjj(i32 2, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjjj(i32 2, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjjj(i32 2, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjjj(i32 2, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjjj(i32 2, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjjj(i32 2, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjjj(i32 2, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjjj(i32 2, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjjj(i32 2, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjjj(i32 2, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjjj(i32 2, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjjj(i32 2, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjjj(i32 2, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjjj(i32 2, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxjjj(i32 2, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxjjj(i32 2, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxjjj(i32 2, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMinjjj(i32 2, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMinjjj(i32 2, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMinjjj(i32 2, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljjj(i32 2, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljjj(i32 2, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljjj(i32 2, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjjj(i32 2, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjjj(i32 2, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjjj(i32 2, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjjj(i32 2, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjjj(i32 2, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjjj(i32 2, i32 2, i32 1)
@@ -210,24 +210,24 @@ SYCL_EXTERNAL void test_group(group<> g) { test(g); }
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjj{{m|y}}(i32 2, i32 1, i64 2)
 
 // int64_t (Linux: long, Windows: long long)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjj{{l|x}}(i32 2, i32 0, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjj{{l|x}}(i32 2, i32 1, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjj{{l|x}}(i32 2, i32 2, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjj{{l|x}}(i32 2, i32 0, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjj{{l|x}}(i32 2, i32 1, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjj{{l|x}}(i32 2, i32 2, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjj{{l|x}}(i32 2, i32 0, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjj{{l|x}}(i32 2, i32 1, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjj{{l|x}}(i32 2, i32 2, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndKHRjj{{l|x}}(i32 2, i32 0, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndKHRjj{{l|x}}(i32 2, i32 1, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndKHRjj{{l|x}}(i32 2, i32 2, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrKHRjj{{l|x}}(i32 2, i32 0, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrKHRjj{{l|x}}(i32 2, i32 1, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrKHRjj{{l|x}}(i32 2, i32 2, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorKHRjj{{l|x}}(i32 2, i32 0, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorKHRjj{{l|x}}(i32 2, i32 1, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorKHRjj{{l|x}}(i32 2, i32 2, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupSMaxjj{{l|x}}(i32 2, i32 0, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupSMaxjj{{l|x}}(i32 2, i32 1, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupSMaxjj{{l|x}}(i32 2, i32 2, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupSMinjj{{l|x}}(i32 2, i32 0, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupSMinjj{{l|x}}(i32 2, i32 1, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupSMinjj{{l|x}}(i32 2, i32 2, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformIMuljj{{l|x}}(i32 2, i32 0, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformIMuljj{{l|x}}(i32 2, i32 1, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformIMuljj{{l|x}}(i32 2, i32 2, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulKHRjj{{l|x}}(i32 2, i32 0, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulKHRjj{{l|x}}(i32 2, i32 1, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulKHRjj{{l|x}}(i32 2, i32 2, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIAddjj{{l|x}}(i32 2, i32 0, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIAddjj{{l|x}}(i32 2, i32 1, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIAddjj{{l|x}}(i32 2, i32 2, i64 1)
@@ -236,24 +236,24 @@ SYCL_EXTERNAL void test_group(group<> g) { test(g); }
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBroadcastj{{l|x}}{{m|y}}(i32 2, i64 1, i64 2)
 
 // uint64_t (Linux: unsigned long, Windows: unsigned long long)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjj{{m|y}}(i32 2, i32 0, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjj{{m|y}}(i32 2, i32 1, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjj{{m|y}}(i32 2, i32 2, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjj{{m|y}}(i32 2, i32 0, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjj{{m|y}}(i32 2, i32 1, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjj{{m|y}}(i32 2, i32 2, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjj{{m|y}}(i32 2, i32 0, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjj{{m|y}}(i32 2, i32 1, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjj{{m|y}}(i32 2, i32 2, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndKHRjj{{m|y}}(i32 2, i32 0, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndKHRjj{{m|y}}(i32 2, i32 1, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndKHRjj{{m|y}}(i32 2, i32 2, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrKHRjj{{m|y}}(i32 2, i32 0, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrKHRjj{{m|y}}(i32 2, i32 1, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrKHRjj{{m|y}}(i32 2, i32 2, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorKHRjj{{m|y}}(i32 2, i32 0, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorKHRjj{{m|y}}(i32 2, i32 1, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorKHRjj{{m|y}}(i32 2, i32 2, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupUMaxjj{{m|y}}(i32 2, i32 0, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupUMaxjj{{m|y}}(i32 2, i32 1, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupUMaxjj{{m|y}}(i32 2, i32 2, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupUMinjj{{m|y}}(i32 2, i32 0, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupUMinjj{{m|y}}(i32 2, i32 1, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupUMinjj{{m|y}}(i32 2, i32 2, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformIMuljj{{m|y}}(i32 2, i32 0, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformIMuljj{{m|y}}(i32 2, i32 1, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformIMuljj{{m|y}}(i32 2, i32 2, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulKHRjj{{m|y}}(i32 2, i32 0, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulKHRjj{{m|y}}(i32 2, i32 1, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulKHRjj{{m|y}}(i32 2, i32 2, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIAddjj{{m|y}}(i32 2, i32 0, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIAddjj{{m|y}}(i32 2, i32 1, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIAddjj{{m|y}}(i32 2, i32 2, i64 1)
@@ -268,9 +268,9 @@ SYCL_EXTERNAL void test_group(group<> g) { test(g); }
 // CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMinjjDF16_(i32 2, i32 0, half 0xH3C00)
 // CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMinjjDF16_(i32 2, i32 1, half 0xH3C00)
 // CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMinjjDF16_(i32 2, i32 2, half 0xH3C00)
-// CHECK: call spir_func half @_Z[[#]]__spirv_GroupNonUniformFMuljjDF16_(i32 2, i32 0, half 0xH3C00)
-// CHECK: call spir_func half @_Z[[#]]__spirv_GroupNonUniformFMuljjDF16_(i32 2, i32 1, half 0xH3C00)
-// CHECK: call spir_func half @_Z[[#]]__spirv_GroupNonUniformFMuljjDF16_(i32 2, i32 2, half 0xH3C00)
+// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMulKHRjjDF16_(i32 2, i32 0, half 0xH3C00)
+// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMulKHRjjDF16_(i32 2, i32 1, half 0xH3C00)
+// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMulKHRjjDF16_(i32 2, i32 2, half 0xH3C00)
 // CHECK: call spir_func half @_Z[[#]]__spirv_GroupFAddjjDF16_(i32 2, i32 0, half 0xH3C00)
 // CHECK: call spir_func half @_Z[[#]]__spirv_GroupFAddjjDF16_(i32 2, i32 1, half 0xH3C00)
 // CHECK: call spir_func half @_Z[[#]]__spirv_GroupFAddjjDF16_(i32 2, i32 2, half 0xH3C00)
@@ -285,9 +285,9 @@ SYCL_EXTERNAL void test_group(group<> g) { test(g); }
 // CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMinjjf(i32 2, i32 0, float 1.000000e+00)
 // CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMinjjf(i32 2, i32 1, float 1.000000e+00)
 // CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMinjjf(i32 2, i32 2, float 1.000000e+00)
-// CHECK: call spir_func float @_Z[[#]]__spirv_GroupNonUniformFMuljjf(i32 2, i32 0, float 1.000000e+00)
-// CHECK: call spir_func float @_Z[[#]]__spirv_GroupNonUniformFMuljjf(i32 2, i32 1, float 1.000000e+00)
-// CHECK: call spir_func float @_Z[[#]]__spirv_GroupNonUniformFMuljjf(i32 2, i32 2, float 1.000000e+00)
+// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMulKHRjjf(i32 2, i32 0, float 1.000000e+00)
+// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMulKHRjjf(i32 2, i32 1, float 1.000000e+00)
+// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMulKHRjjf(i32 2, i32 2, float 1.000000e+00)
 // CHECK: call spir_func float @_Z[[#]]__spirv_GroupFAddjjf(i32 2, i32 0, float 1.000000e+00)
 // CHECK: call spir_func float @_Z[[#]]__spirv_GroupFAddjjf(i32 2, i32 1, float 1.000000e+00)
 // CHECK: call spir_func float @_Z[[#]]__spirv_GroupFAddjjf(i32 2, i32 2, float 1.000000e+00)
@@ -302,9 +302,9 @@ SYCL_EXTERNAL void test_group(group<> g) { test(g); }
 // CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMinjjd(i32 2, i32 0, double 1.000000e+00)
 // CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMinjjd(i32 2, i32 1, double 1.000000e+00)
 // CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMinjjd(i32 2, i32 2, double 1.000000e+00)
-// CHECK: call spir_func double @_Z[[#]]__spirv_GroupNonUniformFMuljjd(i32 2, i32 0, double 1.000000e+00)
-// CHECK: call spir_func double @_Z[[#]]__spirv_GroupNonUniformFMuljjd(i32 2, i32 1, double 1.000000e+00)
-// CHECK: call spir_func double @_Z[[#]]__spirv_GroupNonUniformFMuljjd(i32 2, i32 2, double 1.000000e+00)
+// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMulKHRjjd(i32 2, i32 0, double 1.000000e+00)
+// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMulKHRjjd(i32 2, i32 1, double 1.000000e+00)
+// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMulKHRjjd(i32 2, i32 2, double 1.000000e+00)
 // CHECK: call spir_func double @_Z[[#]]__spirv_GroupFAddjjd(i32 2, i32 0, double 1.000000e+00)
 // CHECK: call spir_func double @_Z[[#]]__spirv_GroupFAddjjd(i32 2, i32 1, double 1.000000e+00)
 // CHECK: call spir_func double @_Z[[#]]__spirv_GroupFAddjjd(i32 2, i32 2, double 1.000000e+00)
@@ -315,24 +315,24 @@ SYCL_EXTERNAL void test_group(group<> g) { test(g); }
 SYCL_EXTERNAL void test_sub_group(sub_group g) { test(g); }
 
 // int8_t
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjji(i32 3, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjji(i32 3, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjji(i32 3, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjji(i32 3, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjji(i32 3, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjji(i32 3, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjji(i32 3, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjji(i32 3, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjji(i32 3, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjji(i32 3, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjji(i32 3, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjji(i32 3, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjji(i32 3, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjji(i32 3, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjji(i32 3, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjji(i32 3, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjji(i32 3, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjji(i32 3, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxjji(i32 3, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxjji(i32 3, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxjji(i32 3, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMinjji(i32 3, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMinjji(i32 3, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMinjji(i32 3, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljji(i32 3, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljji(i32 3, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljji(i32 3, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjji(i32 3, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjji(i32 3, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjji(i32 3, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjji(i32 3, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjji(i32 3, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjji(i32 3, i32 2, i32 1)
@@ -341,24 +341,24 @@ SYCL_EXTERNAL void test_sub_group(sub_group g) { test(g); }
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjij(i32 3, i32 1, i32 2)
 
 // uint8_t
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjjj(i32 3, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjjj(i32 3, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjjj(i32 3, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjjj(i32 3, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjjj(i32 3, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjjj(i32 3, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjjj(i32 3, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjjj(i32 3, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjjj(i32 3, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjjj(i32 3, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjjj(i32 3, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjjj(i32 3, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjjj(i32 3, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjjj(i32 3, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjjj(i32 3, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjjj(i32 3, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjjj(i32 3, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjjj(i32 3, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxjjj(i32 3, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxjjj(i32 3, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxjjj(i32 3, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMinjjj(i32 3, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMinjjj(i32 3, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMinjjj(i32 3, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljjj(i32 3, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljjj(i32 3, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljjj(i32 3, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjjj(i32 3, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjjj(i32 3, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjjj(i32 3, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjjj(i32 3, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjjj(i32 3, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjjj(i32 3, i32 2, i32 1)
@@ -367,24 +367,24 @@ SYCL_EXTERNAL void test_sub_group(sub_group g) { test(g); }
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjjj(i32 3, i32 1, i32 2)
 
 // int16_t
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjji(i32 3, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjji(i32 3, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjji(i32 3, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjji(i32 3, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjji(i32 3, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjji(i32 3, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjji(i32 3, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjji(i32 3, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjji(i32 3, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjji(i32 3, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjji(i32 3, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjji(i32 3, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjji(i32 3, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjji(i32 3, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjji(i32 3, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjji(i32 3, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjji(i32 3, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjji(i32 3, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxjji(i32 3, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxjji(i32 3, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxjji(i32 3, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMinjji(i32 3, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMinjji(i32 3, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMinjji(i32 3, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljji(i32 3, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljji(i32 3, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljji(i32 3, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjji(i32 3, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjji(i32 3, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjji(i32 3, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjji(i32 3, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjji(i32 3, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjji(i32 3, i32 2, i32 1)
@@ -393,24 +393,24 @@ SYCL_EXTERNAL void test_sub_group(sub_group g) { test(g); }
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjij(i32 3, i32 1, i32 2)
 
 // uint16_t
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjjj(i32 3, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjjj(i32 3, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjjj(i32 3, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjjj(i32 3, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjjj(i32 3, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjjj(i32 3, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjjj(i32 3, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjjj(i32 3, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjjj(i32 3, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjjj(i32 3, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjjj(i32 3, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjjj(i32 3, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjjj(i32 3, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjjj(i32 3, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjjj(i32 3, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjjj(i32 3, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjjj(i32 3, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjjj(i32 3, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxjjj(i32 3, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxjjj(i32 3, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxjjj(i32 3, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMinjjj(i32 3, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMinjjj(i32 3, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMinjjj(i32 3, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljjj(i32 3, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljjj(i32 3, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljjj(i32 3, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjjj(i32 3, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjjj(i32 3, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjjj(i32 3, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjjj(i32 3, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjjj(i32 3, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjjj(i32 3, i32 2, i32 1)
@@ -419,24 +419,24 @@ SYCL_EXTERNAL void test_sub_group(sub_group g) { test(g); }
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjjj(i32 3, i32 1, i32 2)
 
 // int32_t
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjji(i32 3, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjji(i32 3, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjji(i32 3, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjji(i32 3, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjji(i32 3, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjji(i32 3, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjji(i32 3, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjji(i32 3, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjji(i32 3, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjji(i32 3, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjji(i32 3, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjji(i32 3, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjji(i32 3, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjji(i32 3, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjji(i32 3, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjji(i32 3, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjji(i32 3, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjji(i32 3, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxjji(i32 3, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxjji(i32 3, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxjji(i32 3, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMinjji(i32 3, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMinjji(i32 3, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMinjji(i32 3, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljji(i32 3, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljji(i32 3, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljji(i32 3, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjji(i32 3, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjji(i32 3, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjji(i32 3, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjji(i32 3, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjji(i32 3, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjji(i32 3, i32 2, i32 1)
@@ -445,24 +445,24 @@ SYCL_EXTERNAL void test_sub_group(sub_group g) { test(g); }
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjij(i32 3, i32 1, i32 2)
 
 // uint32_t
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjjj(i32 3, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjjj(i32 3, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjjj(i32 3, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjjj(i32 3, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjjj(i32 3, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjjj(i32 3, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjjj(i32 3, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjjj(i32 3, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjjj(i32 3, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjjj(i32 3, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjjj(i32 3, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRjjj(i32 3, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjjj(i32 3, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjjj(i32 3, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRjjj(i32 3, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjjj(i32 3, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjjj(i32 3, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRjjj(i32 3, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxjjj(i32 3, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxjjj(i32 3, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxjjj(i32 3, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMinjjj(i32 3, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMinjjj(i32 3, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMinjjj(i32 3, i32 2, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljjj(i32 3, i32 0, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljjj(i32 3, i32 1, i32 1)
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupNonUniformIMuljjj(i32 3, i32 2, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjjj(i32 3, i32 0, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjjj(i32 3, i32 1, i32 1)
+// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRjjj(i32 3, i32 2, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjjj(i32 3, i32 0, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjjj(i32 3, i32 1, i32 1)
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddjjj(i32 3, i32 2, i32 1)
@@ -471,24 +471,24 @@ SYCL_EXTERNAL void test_sub_group(sub_group g) { test(g); }
 // CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjjj(i32 3, i32 1, i32 2)
 
 // int64_t (Linux: long, Windows: long long)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjj{{l|x}}(i32 3, i32 0, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjj{{l|x}}(i32 3, i32 1, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjj{{l|x}}(i32 3, i32 2, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjj{{l|x}}(i32 3, i32 0, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjj{{l|x}}(i32 3, i32 1, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjj{{l|x}}(i32 3, i32 2, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjj{{l|x}}(i32 3, i32 0, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjj{{l|x}}(i32 3, i32 1, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjj{{l|x}}(i32 3, i32 2, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndKHRjj{{l|x}}(i32 3, i32 0, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndKHRjj{{l|x}}(i32 3, i32 1, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndKHRjj{{l|x}}(i32 3, i32 2, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrKHRjj{{l|x}}(i32 3, i32 0, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrKHRjj{{l|x}}(i32 3, i32 1, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrKHRjj{{l|x}}(i32 3, i32 2, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorKHRjj{{l|x}}(i32 3, i32 0, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorKHRjj{{l|x}}(i32 3, i32 1, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorKHRjj{{l|x}}(i32 3, i32 2, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupSMaxjj{{l|x}}(i32 3, i32 0, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupSMaxjj{{l|x}}(i32 3, i32 1, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupSMaxjj{{l|x}}(i32 3, i32 2, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupSMinjj{{l|x}}(i32 3, i32 0, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupSMinjj{{l|x}}(i32 3, i32 1, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupSMinjj{{l|x}}(i32 3, i32 2, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformIMuljj{{l|x}}(i32 3, i32 0, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformIMuljj{{l|x}}(i32 3, i32 1, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformIMuljj{{l|x}}(i32 3, i32 2, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulKHRjj{{l|x}}(i32 3, i32 0, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulKHRjj{{l|x}}(i32 3, i32 1, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulKHRjj{{l|x}}(i32 3, i32 2, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIAddjj{{l|x}}(i32 3, i32 0, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIAddjj{{l|x}}(i32 3, i32 1, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIAddjj{{l|x}}(i32 3, i32 2, i64 1)
@@ -497,24 +497,24 @@ SYCL_EXTERNAL void test_sub_group(sub_group g) { test(g); }
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBroadcastj{{l|x}}j(i32 3, i64 1, i32 2)
 
 // uint64_t (Linux: unsigned long, Windows: unsigned long long)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjj{{m|y}}(i32 3, i32 0, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjj{{m|y}}(i32 3, i32 1, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseAndjj{{m|y}}(i32 3, i32 2, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjj{{m|y}}(i32 3, i32 0, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjj{{m|y}}(i32 3, i32 1, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseOrjj{{m|y}}(i32 3, i32 2, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjj{{m|y}}(i32 3, i32 0, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjj{{m|y}}(i32 3, i32 1, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformBitwiseXorjj{{m|y}}(i32 3, i32 2, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndKHRjj{{m|y}}(i32 3, i32 0, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndKHRjj{{m|y}}(i32 3, i32 1, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndKHRjj{{m|y}}(i32 3, i32 2, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrKHRjj{{m|y}}(i32 3, i32 0, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrKHRjj{{m|y}}(i32 3, i32 1, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrKHRjj{{m|y}}(i32 3, i32 2, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorKHRjj{{m|y}}(i32 3, i32 0, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorKHRjj{{m|y}}(i32 3, i32 1, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorKHRjj{{m|y}}(i32 3, i32 2, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupUMaxjj{{m|y}}(i32 3, i32 0, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupUMaxjj{{m|y}}(i32 3, i32 1, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupUMaxjj{{m|y}}(i32 3, i32 2, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupUMinjj{{m|y}}(i32 3, i32 0, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupUMinjj{{m|y}}(i32 3, i32 1, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupUMinjj{{m|y}}(i32 3, i32 2, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformIMuljj{{m|y}}(i32 3, i32 0, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformIMuljj{{m|y}}(i32 3, i32 1, i64 1)
-// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupNonUniformIMuljj{{m|y}}(i32 3, i32 2, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulKHRjj{{m|y}}(i32 3, i32 0, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulKHRjj{{m|y}}(i32 3, i32 1, i64 1)
+// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulKHRjj{{m|y}}(i32 3, i32 2, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIAddjj{{m|y}}(i32 3, i32 0, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIAddjj{{m|y}}(i32 3, i32 1, i64 1)
 // CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIAddjj{{m|y}}(i32 3, i32 2, i64 1)
@@ -529,9 +529,9 @@ SYCL_EXTERNAL void test_sub_group(sub_group g) { test(g); }
 // CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMinjjDF16_(i32 3, i32 0, half 0xH3C00)
 // CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMinjjDF16_(i32 3, i32 1, half 0xH3C00)
 // CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMinjjDF16_(i32 3, i32 2, half 0xH3C00)
-// CHECK: call spir_func half @_Z[[#]]__spirv_GroupNonUniformFMuljjDF16_(i32 3, i32 0, half 0xH3C00)
-// CHECK: call spir_func half @_Z[[#]]__spirv_GroupNonUniformFMuljjDF16_(i32 3, i32 1, half 0xH3C00)
-// CHECK: call spir_func half @_Z[[#]]__spirv_GroupNonUniformFMuljjDF16_(i32 3, i32 2, half 0xH3C00)
+// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMulKHRjjDF16_(i32 3, i32 0, half 0xH3C00)
+// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMulKHRjjDF16_(i32 3, i32 1, half 0xH3C00)
+// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMulKHRjjDF16_(i32 3, i32 2, half 0xH3C00)
 // CHECK: call spir_func half @_Z[[#]]__spirv_GroupFAddjjDF16_(i32 3, i32 0, half 0xH3C00)
 // CHECK: call spir_func half @_Z[[#]]__spirv_GroupFAddjjDF16_(i32 3, i32 1, half 0xH3C00)
 // CHECK: call spir_func half @_Z[[#]]__spirv_GroupFAddjjDF16_(i32 3, i32 2, half 0xH3C00)
@@ -546,9 +546,9 @@ SYCL_EXTERNAL void test_sub_group(sub_group g) { test(g); }
 // CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMinjjf(i32 3, i32 0, float 1.000000e+00)
 // CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMinjjf(i32 3, i32 1, float 1.000000e+00)
 // CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMinjjf(i32 3, i32 2, float 1.000000e+00)
-// CHECK: call spir_func float @_Z[[#]]__spirv_GroupNonUniformFMuljjf(i32 3, i32 0, float 1.000000e+00)
-// CHECK: call spir_func float @_Z[[#]]__spirv_GroupNonUniformFMuljjf(i32 3, i32 1, float 1.000000e+00)
-// CHECK: call spir_func float @_Z[[#]]__spirv_GroupNonUniformFMuljjf(i32 3, i32 2, float 1.000000e+00)
+// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMulKHRjjf(i32 3, i32 0, float 1.000000e+00)
+// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMulKHRjjf(i32 3, i32 1, float 1.000000e+00)
+// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMulKHRjjf(i32 3, i32 2, float 1.000000e+00)
 // CHECK: call spir_func float @_Z[[#]]__spirv_GroupFAddjjf(i32 3, i32 0, float 1.000000e+00)
 // CHECK: call spir_func float @_Z[[#]]__spirv_GroupFAddjjf(i32 3, i32 1, float 1.000000e+00)
 // CHECK: call spir_func float @_Z[[#]]__spirv_GroupFAddjjf(i32 3, i32 2, float 1.000000e+00)
@@ -563,9 +563,9 @@ SYCL_EXTERNAL void test_sub_group(sub_group g) { test(g); }
 // CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMinjjd(i32 3, i32 0, double 1.000000e+00)
 // CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMinjjd(i32 3, i32 1, double 1.000000e+00)
 // CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMinjjd(i32 3, i32 2, double 1.000000e+00)
-// CHECK: call spir_func double @_Z[[#]]__spirv_GroupNonUniformFMuljjd(i32 3, i32 0, double 1.000000e+00)
-// CHECK: call spir_func double @_Z[[#]]__spirv_GroupNonUniformFMuljjd(i32 3, i32 1, double 1.000000e+00)
-// CHECK: call spir_func double @_Z[[#]]__spirv_GroupNonUniformFMuljjd(i32 3, i32 2, double 1.000000e+00)
+// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMulKHRjjd(i32 3, i32 0, double 1.000000e+00)
+// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMulKHRjjd(i32 3, i32 1, double 1.000000e+00)
+// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMulKHRjjd(i32 3, i32 2, double 1.000000e+00)
 // CHECK: call spir_func double @_Z[[#]]__spirv_GroupFAddjjd(i32 3, i32 0, double 1.000000e+00)
 // CHECK: call spir_func double @_Z[[#]]__spirv_GroupFAddjjd(i32 3, i32 1, double 1.000000e+00)
 // CHECK: call spir_func double @_Z[[#]]__spirv_GroupFAddjjd(i32 3, i32 2, double 1.000000e+00)