From 9ed3ef23855a70d151ec707b44bb966bbfb5df1a Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy Date: Tue, 1 Mar 2022 15:50:32 +0300 Subject: [PATCH 1/4] [SYCL] Use uniform group instructions in headers --- clang/lib/Driver/ToolChains/Clang.cpp | 3 +- clang/lib/Sema/SPIRVBuiltins.td | 6 +- sycl/include/sycl/ext/oneapi/functional.hpp | 20 +- .../check_device_code/group_operations.cpp | 420 +++++++++--------- 4 files changed, 225 insertions(+), 224 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 6acbfd2024e46..9f9cf581fa909 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_INTEL_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..15bbc74796e8e 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", "GroupIMulINTEL", "GroupBitwiseOrINTEL", + "GroupBitwiseXorINTEL", "GroupBitwiseAndINTEL"] in { def : SPVBuiltin; } foreach name = ["GroupFAdd", "GroupFMin", "GroupFMax", - "GroupNonUniformFMul"] in { + "GroupFMulINTEL"] in { def : SPVBuiltin; } diff --git a/sycl/include/sycl/ext/oneapi/functional.hpp b/sycl/include/sycl/ext/oneapi/functional.hpp index 46a5a18fd27ac..82b6e720b04d1 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) __SYCL_CALC_OVERLOAD(GroupOpIUnsigned, IAdd, sycl::plus) __SYCL_CALC_OVERLOAD(GroupOpFP, FAdd, sycl::plus) -__SYCL_CALC_OVERLOAD(GroupOpISigned, NonUniformIMul, sycl::multiplies) -__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, NonUniformIMul, sycl::multiplies) -__SYCL_CALC_OVERLOAD(GroupOpFP, NonUniformFMul, sycl::multiplies) - -__SYCL_CALC_OVERLOAD(GroupOpISigned, NonUniformBitwiseOr, sycl::bit_or) -__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, NonUniformBitwiseOr, sycl::bit_or) -__SYCL_CALC_OVERLOAD(GroupOpISigned, NonUniformBitwiseXor, sycl::bit_xor) -__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, NonUniformBitwiseXor, sycl::bit_xor) -__SYCL_CALC_OVERLOAD(GroupOpISigned, NonUniformBitwiseAnd, sycl::bit_and) -__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, NonUniformBitwiseAnd, sycl::bit_and) +__SYCL_CALC_OVERLOAD(GroupOpISigned, IMulINTEL, sycl::multiplies) +__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, IMulINTEL, sycl::multiplies) +__SYCL_CALC_OVERLOAD(GroupOpFP, FMulINTEL, sycl::multiplies) + +__SYCL_CALC_OVERLOAD(GroupOpISigned, BitwiseOrINTEL, sycl::bit_or) +__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, BitwiseOrINTEL, sycl::bit_or) +__SYCL_CALC_OVERLOAD(GroupOpISigned, BitwiseXorINTEL, sycl::bit_xor) +__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, BitwiseXorINTEL, sycl::bit_xor) +__SYCL_CALC_OVERLOAD(GroupOpISigned, BitwiseAndINTEL, sycl::bit_and) +__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, BitwiseAndINTEL, sycl::bit_and) #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..308607549c315 100644 --- a/sycl/test/check_device_code/group_operations.cpp +++ b/sycl/test/check_device_code/group_operations.cpp @@ -54,24 +54,24 @@ template [[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_GroupBitwiseAndINTELjji(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjji(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjji(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(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_GroupIMulINTELjji(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjji(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjji(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_GroupBitwiseAndINTELjjj(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjjj(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjjj(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(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_GroupIMulINTELjjj(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjjj(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjjj(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_GroupBitwiseAndINTELjji(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjji(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjji(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(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_GroupIMulINTELjji(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjji(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjji(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_GroupBitwiseAndINTELjjj(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjjj(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjjj(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(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_GroupIMulINTELjjj(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjjj(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjjj(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_GroupBitwiseAndINTELjji(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjji(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjji(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(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_GroupIMulINTELjji(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjji(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjji(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_GroupBitwiseAndINTELjjj(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjjj(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjjj(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(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_GroupIMulINTELjjj(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjjj(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjjj(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_GroupBitwiseAndINTELjj{{l|x}}(i32 2, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndINTELjj{{l|x}}(i32 2, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndINTELjj{{l|x}}(i32 2, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrINTELjj{{l|x}}(i32 2, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrINTELjj{{l|x}}(i32 2, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrINTELjj{{l|x}}(i32 2, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorINTELjj{{l|x}}(i32 2, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorINTELjj{{l|x}}(i32 2, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorINTELjj{{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_GroupIMulINTELjj{{l|x}}(i32 2, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulINTELjj{{l|x}}(i32 2, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulINTELjj{{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_GroupBitwiseAndINTELjj{{m|y}}(i32 2, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndINTELjj{{m|y}}(i32 2, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndINTELjj{{m|y}}(i32 2, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrINTELjj{{m|y}}(i32 2, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrINTELjj{{m|y}}(i32 2, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrINTELjj{{m|y}}(i32 2, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorINTELjj{{m|y}}(i32 2, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorINTELjj{{m|y}}(i32 2, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorINTELjj{{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_GroupIMulINTELjj{{m|y}}(i32 2, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulINTELjj{{m|y}}(i32 2, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulINTELjj{{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_GroupFMulINTELjjDF16_(i32 2, i32 0, half 0xH3C00) +// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMulINTELjjDF16_(i32 2, i32 1, half 0xH3C00) +// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMulINTELjjDF16_(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_GroupFMulINTELjjf(i32 2, i32 0, float 1.000000e+00) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMulINTELjjf(i32 2, i32 1, float 1.000000e+00) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMulINTELjjf(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_GroupFMulINTELjjd(i32 2, i32 0, double 1.000000e+00) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMulINTELjjd(i32 2, i32 1, double 1.000000e+00) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMulINTELjjd(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_GroupBitwiseAndINTELjji(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjji(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjji(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(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_GroupIMulINTELjji(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjji(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjji(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_GroupBitwiseAndINTELjjj(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjjj(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjjj(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(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_GroupIMulINTELjjj(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjjj(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjjj(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_GroupBitwiseAndINTELjji(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjji(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjji(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(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_GroupIMulINTELjji(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjji(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjji(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_GroupBitwiseAndINTELjjj(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjjj(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjjj(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(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_GroupIMulINTELjjj(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjjj(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjjj(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_GroupBitwiseAndINTELjji(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjji(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjji(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(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_GroupIMulINTELjji(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjji(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjji(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_GroupBitwiseAndINTELjjj(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjjj(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjjj(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(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_GroupIMulINTELjjj(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjjj(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjjj(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_GroupBitwiseAndINTELjj{{l|x}}(i32 3, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndINTELjj{{l|x}}(i32 3, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndINTELjj{{l|x}}(i32 3, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrINTELjj{{l|x}}(i32 3, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrINTELjj{{l|x}}(i32 3, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrINTELjj{{l|x}}(i32 3, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorINTELjj{{l|x}}(i32 3, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorINTELjj{{l|x}}(i32 3, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorINTELjj{{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_GroupIMulINTELjj{{l|x}}(i32 3, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulINTELjj{{l|x}}(i32 3, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulINTELjj{{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_GroupBitwiseAndINTELjj{{m|y}}(i32 3, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndINTELjj{{m|y}}(i32 3, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndINTELjj{{m|y}}(i32 3, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrINTELjj{{m|y}}(i32 3, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrINTELjj{{m|y}}(i32 3, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrINTELjj{{m|y}}(i32 3, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorINTELjj{{m|y}}(i32 3, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorINTELjj{{m|y}}(i32 3, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorINTELjj{{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_GroupIMulINTELjj{{m|y}}(i32 3, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulINTELjj{{m|y}}(i32 3, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulINTELjj{{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_GroupFMulINTELjjDF16_(i32 3, i32 0, half 0xH3C00) +// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMulINTELjjDF16_(i32 3, i32 1, half 0xH3C00) +// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMulINTELjjDF16_(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_GroupFMulINTELjjf(i32 3, i32 0, float 1.000000e+00) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMulINTELjjf(i32 3, i32 1, float 1.000000e+00) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMulINTELjjf(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_GroupFMulINTELjjd(i32 3, i32 0, double 1.000000e+00) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMulINTELjjd(i32 3, i32 1, double 1.000000e+00) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMulINTELjjd(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) From f035628b2a4da098a2ac0e2f0c6fd42e30f0d790 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy Date: Thu, 3 Mar 2022 11:51:44 +0300 Subject: [PATCH 2/4] Update instruction names --- clang/lib/Driver/ToolChains/Clang.cpp | 6 +- clang/lib/Sema/SPIRVBuiltins.td | 6 +- sycl/include/sycl/ext/oneapi/functional.hpp | 20 +- .../check_device_code/group_operations.cpp | 420 +++++++++--------- 4 files changed, 226 insertions(+), 226 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 9f9cf581fa909..1b41519a24062 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -9158,7 +9158,8 @@ void SPIRVTranslator::ConstructJob(Compilation &C, const JobAction &JA, std::string DefaultExtArg = ",+SPV_EXT_shader_atomic_float_add,+SPV_EXT_shader_atomic_float_min_max" ",+SPV_KHR_no_integer_wrap_decoration,+SPV_KHR_float_controls" - ",+SPV_KHR_expect_assume,+SPV_KHR_linkonce_odr"; + ",+SPV_KHR_expect_assume,+SPV_KHR_linkonce_odr" + ",+SPV_KHR_uniform_group_instructions"; std::string INTELExtArg = ",+SPV_INTEL_subgroups,+SPV_INTEL_media_block_io" ",+SPV_INTEL_device_side_avc_motion_estimation" @@ -9188,8 +9189,7 @@ 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_uniform_group_instructions"; + ",+SPV_INTEL_hw_thread_queries"; 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 15bbc74796e8e..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", "GroupIMulINTEL", "GroupBitwiseOrINTEL", - "GroupBitwiseXorINTEL", "GroupBitwiseAndINTEL"] in { +foreach name = ["GroupIAdd", "GroupIMulKHR", "GroupBitwiseOrKHR", + "GroupBitwiseXorKHR", "GroupBitwiseAndKHR"] in { def : SPVBuiltin; } foreach name = ["GroupFAdd", "GroupFMin", "GroupFMax", - "GroupFMulINTEL"] in { + "GroupFMulKHR"] in { def : SPVBuiltin; } diff --git a/sycl/include/sycl/ext/oneapi/functional.hpp b/sycl/include/sycl/ext/oneapi/functional.hpp index 82b6e720b04d1..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) __SYCL_CALC_OVERLOAD(GroupOpIUnsigned, IAdd, sycl::plus) __SYCL_CALC_OVERLOAD(GroupOpFP, FAdd, sycl::plus) -__SYCL_CALC_OVERLOAD(GroupOpISigned, IMulINTEL, sycl::multiplies) -__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, IMulINTEL, sycl::multiplies) -__SYCL_CALC_OVERLOAD(GroupOpFP, FMulINTEL, sycl::multiplies) - -__SYCL_CALC_OVERLOAD(GroupOpISigned, BitwiseOrINTEL, sycl::bit_or) -__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, BitwiseOrINTEL, sycl::bit_or) -__SYCL_CALC_OVERLOAD(GroupOpISigned, BitwiseXorINTEL, sycl::bit_xor) -__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, BitwiseXorINTEL, sycl::bit_xor) -__SYCL_CALC_OVERLOAD(GroupOpISigned, BitwiseAndINTEL, sycl::bit_and) -__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, BitwiseAndINTEL, sycl::bit_and) +__SYCL_CALC_OVERLOAD(GroupOpISigned, IMulKHR, sycl::multiplies) +__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, IMulKHR, sycl::multiplies) +__SYCL_CALC_OVERLOAD(GroupOpFP, FMulKHR, sycl::multiplies) + +__SYCL_CALC_OVERLOAD(GroupOpISigned, BitwiseOrKHR, sycl::bit_or) +__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, BitwiseOrKHR, sycl::bit_or) +__SYCL_CALC_OVERLOAD(GroupOpISigned, BitwiseXorKHR, sycl::bit_xor) +__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, BitwiseXorKHR, sycl::bit_xor) +__SYCL_CALC_OVERLOAD(GroupOpISigned, BitwiseAndKHR, sycl::bit_and) +__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, BitwiseAndKHR, sycl::bit_and) #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 308607549c315..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 [[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_GroupBitwiseAndINTELjji(i32 2, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjji(i32 2, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjji(i32 2, i32 2, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 2, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 2, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 2, i32 2, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(i32 2, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(i32 2, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(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_GroupIMulINTELjji(i32 2, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjji(i32 2, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjji(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_GroupBitwiseAndINTELjjj(i32 2, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjjj(i32 2, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjjj(i32 2, i32 2, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 2, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 2, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 2, i32 2, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(i32 2, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(i32 2, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(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_GroupIMulINTELjjj(i32 2, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjjj(i32 2, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjjj(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_GroupBitwiseAndINTELjji(i32 2, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjji(i32 2, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjji(i32 2, i32 2, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 2, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 2, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 2, i32 2, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(i32 2, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(i32 2, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(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_GroupIMulINTELjji(i32 2, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjji(i32 2, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjji(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_GroupBitwiseAndINTELjjj(i32 2, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjjj(i32 2, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjjj(i32 2, i32 2, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 2, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 2, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 2, i32 2, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(i32 2, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(i32 2, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(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_GroupIMulINTELjjj(i32 2, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjjj(i32 2, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjjj(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_GroupBitwiseAndINTELjji(i32 2, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjji(i32 2, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjji(i32 2, i32 2, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 2, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 2, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 2, i32 2, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(i32 2, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(i32 2, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(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_GroupIMulINTELjji(i32 2, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjji(i32 2, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjji(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_GroupBitwiseAndINTELjjj(i32 2, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjjj(i32 2, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjjj(i32 2, i32 2, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 2, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 2, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 2, i32 2, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(i32 2, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(i32 2, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(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_GroupIMulINTELjjj(i32 2, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjjj(i32 2, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjjj(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_GroupBitwiseAndINTELjj{{l|x}}(i32 2, i32 0, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndINTELjj{{l|x}}(i32 2, i32 1, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndINTELjj{{l|x}}(i32 2, i32 2, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrINTELjj{{l|x}}(i32 2, i32 0, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrINTELjj{{l|x}}(i32 2, i32 1, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrINTELjj{{l|x}}(i32 2, i32 2, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorINTELjj{{l|x}}(i32 2, i32 0, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorINTELjj{{l|x}}(i32 2, i32 1, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorINTELjj{{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_GroupIMulINTELjj{{l|x}}(i32 2, i32 0, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulINTELjj{{l|x}}(i32 2, i32 1, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulINTELjj{{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_GroupBitwiseAndINTELjj{{m|y}}(i32 2, i32 0, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndINTELjj{{m|y}}(i32 2, i32 1, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndINTELjj{{m|y}}(i32 2, i32 2, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrINTELjj{{m|y}}(i32 2, i32 0, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrINTELjj{{m|y}}(i32 2, i32 1, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrINTELjj{{m|y}}(i32 2, i32 2, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorINTELjj{{m|y}}(i32 2, i32 0, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorINTELjj{{m|y}}(i32 2, i32 1, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorINTELjj{{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_GroupIMulINTELjj{{m|y}}(i32 2, i32 0, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulINTELjj{{m|y}}(i32 2, i32 1, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulINTELjj{{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_GroupFMulINTELjjDF16_(i32 2, i32 0, half 0xH3C00) -// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMulINTELjjDF16_(i32 2, i32 1, half 0xH3C00) -// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMulINTELjjDF16_(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_GroupFMulINTELjjf(i32 2, i32 0, float 1.000000e+00) -// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMulINTELjjf(i32 2, i32 1, float 1.000000e+00) -// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMulINTELjjf(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_GroupFMulINTELjjd(i32 2, i32 0, double 1.000000e+00) -// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMulINTELjjd(i32 2, i32 1, double 1.000000e+00) -// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMulINTELjjd(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_GroupBitwiseAndINTELjji(i32 3, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjji(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjji(i32 3, i32 2, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 3, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 3, i32 2, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(i32 3, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(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_GroupIMulINTELjji(i32 3, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjji(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjji(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_GroupBitwiseAndINTELjjj(i32 3, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjjj(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjjj(i32 3, i32 2, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 3, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 3, i32 2, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(i32 3, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(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_GroupIMulINTELjjj(i32 3, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjjj(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjjj(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_GroupBitwiseAndINTELjji(i32 3, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjji(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjji(i32 3, i32 2, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 3, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 3, i32 2, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(i32 3, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(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_GroupIMulINTELjji(i32 3, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjji(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjji(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_GroupBitwiseAndINTELjjj(i32 3, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjjj(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjjj(i32 3, i32 2, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 3, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 3, i32 2, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(i32 3, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(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_GroupIMulINTELjjj(i32 3, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjjj(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjjj(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_GroupBitwiseAndINTELjji(i32 3, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjji(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjji(i32 3, i32 2, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 3, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjji(i32 3, i32 2, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(i32 3, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjji(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_GroupIMulINTELjji(i32 3, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjji(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjji(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_GroupBitwiseAndINTELjjj(i32 3, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjjj(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndINTELjjj(i32 3, i32 2, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 3, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrINTELjjj(i32 3, i32 2, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(i32 3, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorINTELjjj(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_GroupIMulINTELjjj(i32 3, i32 0, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjjj(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulINTELjjj(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_GroupBitwiseAndINTELjj{{l|x}}(i32 3, i32 0, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndINTELjj{{l|x}}(i32 3, i32 1, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndINTELjj{{l|x}}(i32 3, i32 2, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrINTELjj{{l|x}}(i32 3, i32 0, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrINTELjj{{l|x}}(i32 3, i32 1, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrINTELjj{{l|x}}(i32 3, i32 2, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorINTELjj{{l|x}}(i32 3, i32 0, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorINTELjj{{l|x}}(i32 3, i32 1, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorINTELjj{{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_GroupIMulINTELjj{{l|x}}(i32 3, i32 0, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulINTELjj{{l|x}}(i32 3, i32 1, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulINTELjj{{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_GroupBitwiseAndINTELjj{{m|y}}(i32 3, i32 0, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndINTELjj{{m|y}}(i32 3, i32 1, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndINTELjj{{m|y}}(i32 3, i32 2, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrINTELjj{{m|y}}(i32 3, i32 0, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrINTELjj{{m|y}}(i32 3, i32 1, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrINTELjj{{m|y}}(i32 3, i32 2, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorINTELjj{{m|y}}(i32 3, i32 0, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorINTELjj{{m|y}}(i32 3, i32 1, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorINTELjj{{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_GroupIMulINTELjj{{m|y}}(i32 3, i32 0, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulINTELjj{{m|y}}(i32 3, i32 1, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulINTELjj{{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_GroupFMulINTELjjDF16_(i32 3, i32 0, half 0xH3C00) -// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMulINTELjjDF16_(i32 3, i32 1, half 0xH3C00) -// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMulINTELjjDF16_(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_GroupFMulINTELjjf(i32 3, i32 0, float 1.000000e+00) -// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMulINTELjjf(i32 3, i32 1, float 1.000000e+00) -// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMulINTELjjf(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_GroupFMulINTELjjd(i32 3, i32 0, double 1.000000e+00) -// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMulINTELjjd(i32 3, i32 1, double 1.000000e+00) -// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMulINTELjjd(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) From 121db8b3acf4132e7f8e02800582a89d41d1d657 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy Date: Thu, 3 Mar 2022 14:25:53 +0300 Subject: [PATCH 3/4] Address review comments --- clang/lib/Driver/ToolChains/Clang.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 1b41519a24062..820ce503385ae 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -9158,8 +9158,7 @@ void SPIRVTranslator::ConstructJob(Compilation &C, const JobAction &JA, std::string DefaultExtArg = ",+SPV_EXT_shader_atomic_float_add,+SPV_EXT_shader_atomic_float_min_max" ",+SPV_KHR_no_integer_wrap_decoration,+SPV_KHR_float_controls" - ",+SPV_KHR_expect_assume,+SPV_KHR_linkonce_odr" - ",+SPV_KHR_uniform_group_instructions"; + ",+SPV_KHR_expect_assume,+SPV_KHR_linkonce_odr"; std::string INTELExtArg = ",+SPV_INTEL_subgroups,+SPV_INTEL_media_block_io" ",+SPV_INTEL_device_side_avc_motion_estimation" @@ -9189,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) { From 8bd1e8c586357a1d215663ba643b955f34232f08 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy Date: Fri, 18 Mar 2022 20:02:02 +0300 Subject: [PATCH 4/4] Update check-clang test --- clang/test/Driver/sycl-spirv-ext.c | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) 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