From 8e18e42b4197f2b450a45431d45d90cd9a6d0352 Mon Sep 17 00:00:00 2001 From: Wenju He Date: Thu, 13 Mar 2025 02:00:51 -0700 Subject: [PATCH 1/6] [SYCL] Change SPIR-V Enum token type from unsigned int to int for groups builtins Motivation is the same as PR #17438, i.e. unifying SPIR-V builtin mangling to enhance SYCL AOT support for backend targets that bypass SPIR-V generation. --- clang/lib/Sema/SPIRVBuiltins.td | 54 +- .../include/libspirv/spirv_builtins.h | 280 +++--- .../amdgcn-amdhsa/async/wait_group_events.cl | 2 +- .../lib/amdgcn-amdhsa/group/collectives.cl | 24 +- .../lib/amdgcn-amdhsa/group/group_ballot.cl | 2 +- .../async/async_work_group_strided_copy.inc | 14 +- .../lib/generic/async/wait_group_events.cl | 2 +- libclc/libspirv/lib/generic/float16.cl | 24 +- .../async/async_work_group_strided_copy.cl | 8 +- .../ptx-nvidiacl/async/wait_group_events.cl | 2 +- .../lib/ptx-nvidiacl/group/collectives.cl | 44 +- .../ptx-nvidiacl/group/group_non_uniform.cl | 4 +- libdevice/nativecpu_utils.cpp | 10 +- sycl/include/sycl/__spirv/spirv_ops.hpp | 88 +- sycl/include/sycl/__spirv/spirv_types.hpp | 14 +- sycl/source/spirv_ops.cpp | 2 +- sycl/test/abi/sycl_symbols_linux.dump | 2 +- sycl/test/abi/sycl_symbols_windows.dump | 2 +- .../check_device_code/group_operations.cpp | 948 +++++++++--------- .../test/check_device_code/sub_group_mask.cpp | 2 +- 20 files changed, 765 insertions(+), 763 deletions(-) diff --git a/clang/lib/Sema/SPIRVBuiltins.td b/clang/lib/Sema/SPIRVBuiltins.td index 955c58c5532ab..2e005346bd1ab 100644 --- a/clang/lib/Sema/SPIRVBuiltins.td +++ b/clang/lib/Sema/SPIRVBuiltins.td @@ -911,52 +911,68 @@ foreach name = ["MemoryBarrier"] in { // 3.32.21. Group and Subgroup Instructions foreach name = ["GroupAsyncCopy"] in { - // TODO: Allow enum flags instead of UInt ? - // TODO: We should enforce that the UInt must be a literal. - def : SPVBuiltin, PointerType, GlobalAS>, Size, Size, Event], Attr.Convergent>; - def : SPVBuiltin, PointerType, LocalAS>, Size, Size, Event], Attr.Convergent>; + // TODO: Allow enum flags instead of Int ? + // TODO: We should enforce that the Int must be a literal. + def : SPVBuiltin, + PointerType, GlobalAS>, Size, Size, + Event], + Attr.Convergent>; + def : SPVBuiltin, + PointerType, LocalAS>, Size, Size, + Event], + Attr.Convergent>; } foreach name = ["GroupWaitEvents"] in { - def : SPVBuiltin], Attr.Convergent>; - def : SPVBuiltin], Attr.Convergent>; - def : SPVBuiltin], Attr.Convergent>; + def : SPVBuiltin], + Attr.Convergent>; + def : SPVBuiltin], + Attr.Convergent>; + def : SPVBuiltin], + Attr.Convergent>; } foreach name = ["GroupAll", "GroupAny"] in { - def : SPVBuiltin; + def : SPVBuiltin; } foreach name = ["GroupBroadcast"] in { foreach IDType = TLAllInts.List in { - def : SPVBuiltin; - def : SPVBuiltin], Attr.Convergent>; - def : SPVBuiltin], Attr.Convergent>; - def : SPVBuiltin; - def : SPVBuiltin], Attr.Convergent>; - def : SPVBuiltin], Attr.Convergent>; + def : SPVBuiltin; + def : SPVBuiltin], + Attr.Convergent>; + def : SPVBuiltin], + Attr.Convergent>; + def : SPVBuiltin; + def : SPVBuiltin], + Attr.Convergent>; + def : SPVBuiltin], + Attr.Convergent>; } } foreach name = ["GroupIAdd", "GroupIMulKHR", "GroupBitwiseOrKHR", "GroupBitwiseXorKHR", "GroupBitwiseAndKHR"] in { - def : SPVBuiltin; + def : SPVBuiltin; } foreach name = ["GroupFAdd", "GroupFMin", "GroupFMax", "GroupFMulKHR"] in { - def : SPVBuiltin; + def : SPVBuiltin; } foreach name = ["GroupUMin", "GroupUMax"] in { - def : SPVBuiltin; + def : SPVBuiltin; } foreach name = ["GroupSMin", "GroupSMax"] in { - def : SPVBuiltin; + def : SPVBuiltin; } // TODO: These builtins need to support vectors of bool. foreach name = ["GroupLogicalAndKHR", "GroupLogicalOrKHR"] in { - def : SPVBuiltin; + def : SPVBuiltin; } diff --git a/libclc/libspirv/include/libspirv/spirv_builtins.h b/libclc/libspirv/include/libspirv/spirv_builtins.h index 8a3f9070a1aee..0d10b275c0622 100644 --- a/libclc/libspirv/include/libspirv/spirv_builtins.h +++ b/libclc/libspirv/include/libspirv/spirv_builtins.h @@ -9852,49 +9852,49 @@ _CLC_OVERLOAD _CLC_DECL _CLC_CONSTFN __clc_vec16_int8_t #endif _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t, __clc_char_t __local *, __clc_char_t const __global *, + __clc_int32_t, __clc_char_t __local *, __clc_char_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t, __clc_char_t __global *, __clc_char_t const __local *, + __clc_int32_t, __clc_char_t __global *, __clc_char_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec2_char_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec2_char_t __local *, __clc_vec2_char_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec2_char_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec2_char_t __global *, __clc_vec2_char_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec3_char_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec3_char_t __local *, __clc_vec3_char_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec3_char_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec3_char_t __global *, __clc_vec3_char_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec4_char_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec4_char_t __local *, __clc_vec4_char_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec4_char_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec4_char_t __global *, __clc_vec4_char_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec8_char_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec8_char_t __local *, __clc_vec8_char_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec8_char_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec8_char_t __global *, __clc_vec8_char_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec16_char_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec16_char_t __local *, __clc_vec16_char_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec16_char_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec16_char_t __global *, __clc_vec16_char_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( @@ -9904,135 +9904,135 @@ _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( __clc_uint32_t, __clc_int8_t __global *, __clc_int8_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec2_int8_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec2_int8_t __local *, __clc_vec2_int8_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec2_int8_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec2_int8_t __global *, __clc_vec2_int8_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec3_int8_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec3_int8_t __local *, __clc_vec3_int8_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec3_int8_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec3_int8_t __global *, __clc_vec3_int8_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec4_int8_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec4_int8_t __local *, __clc_vec4_int8_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec4_int8_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec4_int8_t __global *, __clc_vec4_int8_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec8_int8_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec8_int8_t __local *, __clc_vec8_int8_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec8_int8_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec8_int8_t __global *, __clc_vec8_int8_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec16_int8_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec16_int8_t __local *, __clc_vec16_int8_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec16_int8_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec16_int8_t __global *, __clc_vec16_int8_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t, __clc_int16_t __local *, __clc_int16_t const __global *, + __clc_int32_t, __clc_int16_t __local *, __clc_int16_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t, __clc_int16_t __global *, __clc_int16_t const __local *, + __clc_int32_t, __clc_int16_t __global *, __clc_int16_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec2_int16_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec2_int16_t __local *, __clc_vec2_int16_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec2_int16_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec2_int16_t __global *, __clc_vec2_int16_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec3_int16_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec3_int16_t __local *, __clc_vec3_int16_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec3_int16_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec3_int16_t __global *, __clc_vec3_int16_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec4_int16_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec4_int16_t __local *, __clc_vec4_int16_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec4_int16_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec4_int16_t __global *, __clc_vec4_int16_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec8_int16_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec8_int16_t __local *, __clc_vec8_int16_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec8_int16_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec8_int16_t __global *, __clc_vec8_int16_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec16_int16_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec16_int16_t __local *, __clc_vec16_int16_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec16_int16_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec16_int16_t __global *, __clc_vec16_int16_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t, __clc_int32_t __local *, __clc_int32_t const __global *, + __clc_int32_t, __clc_int32_t __local *, __clc_int32_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( __clc_uint32_t, __clc_int32_t __global *, __clc_int32_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec2_int32_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec2_int32_t __local *, __clc_vec2_int32_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec2_int32_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec2_int32_t __global *, __clc_vec2_int32_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec3_int32_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec3_int32_t __local *, __clc_vec3_int32_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec3_int32_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec3_int32_t __global *, __clc_vec3_int32_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec4_int32_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec4_int32_t __local *, __clc_vec4_int32_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec4_int32_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec4_int32_t __global *, __clc_vec4_int32_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec8_int32_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec8_int32_t __local *, __clc_vec8_int32_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec8_int32_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec8_int32_t __global *, __clc_vec8_int32_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec16_int32_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec16_int32_t __local *, __clc_vec16_int32_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec16_int32_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec16_int32_t __global *, __clc_vec16_int32_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( @@ -10042,376 +10042,376 @@ _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( __clc_uint32_t, __clc_int64_t __global *, __clc_int64_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec2_int64_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec2_int64_t __local *, __clc_vec2_int64_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec2_int64_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec2_int64_t __global *, __clc_vec2_int64_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec3_int64_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec3_int64_t __local *, __clc_vec3_int64_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec3_int64_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec3_int64_t __global *, __clc_vec3_int64_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec4_int64_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec4_int64_t __local *, __clc_vec4_int64_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec4_int64_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec4_int64_t __global *, __clc_vec4_int64_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec8_int64_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec8_int64_t __local *, __clc_vec8_int64_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec8_int64_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec8_int64_t __global *, __clc_vec8_int64_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec16_int64_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec16_int64_t __local *, __clc_vec16_int64_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec16_int64_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec16_int64_t __global *, __clc_vec16_int64_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t, __clc_uint8_t __local *, __clc_uint8_t const __global *, + __clc_int32_t, __clc_uint8_t __local *, __clc_uint8_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t, __clc_uint8_t __global *, __clc_uint8_t const __local *, + __clc_int32_t, __clc_uint8_t __global *, __clc_uint8_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec2_uint8_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec2_uint8_t __local *, __clc_vec2_uint8_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec2_uint8_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec2_uint8_t __global *, __clc_vec2_uint8_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec3_uint8_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec3_uint8_t __local *, __clc_vec3_uint8_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec3_uint8_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec3_uint8_t __global *, __clc_vec3_uint8_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec4_uint8_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec4_uint8_t __local *, __clc_vec4_uint8_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec4_uint8_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec4_uint8_t __global *, __clc_vec4_uint8_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec8_uint8_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec8_uint8_t __local *, __clc_vec8_uint8_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec8_uint8_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec8_uint8_t __global *, __clc_vec8_uint8_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec16_uint8_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec16_uint8_t __local *, __clc_vec16_uint8_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec16_uint8_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec16_uint8_t __global *, __clc_vec16_uint8_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t, __clc_uint16_t __local *, __clc_uint16_t const __global *, + __clc_int32_t, __clc_uint16_t __local *, __clc_uint16_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t, __clc_uint16_t __global *, __clc_uint16_t const __local *, + __clc_int32_t, __clc_uint16_t __global *, __clc_uint16_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec2_uint16_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec2_uint16_t __local *, __clc_vec2_uint16_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec2_uint16_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec2_uint16_t __global *, __clc_vec2_uint16_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec3_uint16_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec3_uint16_t __local *, __clc_vec3_uint16_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec3_uint16_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec3_uint16_t __global *, __clc_vec3_uint16_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec4_uint16_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec4_uint16_t __local *, __clc_vec4_uint16_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec4_uint16_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec4_uint16_t __global *, __clc_vec4_uint16_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec8_uint16_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec8_uint16_t __local *, __clc_vec8_uint16_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec8_uint16_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec8_uint16_t __global *, __clc_vec8_uint16_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec16_uint16_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec16_uint16_t __local *, __clc_vec16_uint16_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec16_uint16_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec16_uint16_t __global *, __clc_vec16_uint16_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t, __clc_uint32_t __local *, __clc_uint32_t const __global *, + __clc_int32_t, __clc_uint32_t __local *, __clc_uint32_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t, __clc_uint32_t __global *, __clc_uint32_t const __local *, + __clc_int32_t, __clc_uint32_t __global *, __clc_uint32_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec2_uint32_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec2_uint32_t __local *, __clc_vec2_uint32_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec2_uint32_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec2_uint32_t __global *, __clc_vec2_uint32_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec3_uint32_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec3_uint32_t __local *, __clc_vec3_uint32_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec3_uint32_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec3_uint32_t __global *, __clc_vec3_uint32_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec4_uint32_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec4_uint32_t __local *, __clc_vec4_uint32_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec4_uint32_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec4_uint32_t __global *, __clc_vec4_uint32_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec8_uint32_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec8_uint32_t __local *, __clc_vec8_uint32_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec8_uint32_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec8_uint32_t __global *, __clc_vec8_uint32_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec16_uint32_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec16_uint32_t __local *, __clc_vec16_uint32_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec16_uint32_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec16_uint32_t __global *, __clc_vec16_uint32_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t, __clc_uint64_t __local *, __clc_uint64_t const __global *, + __clc_int32_t, __clc_uint64_t __local *, __clc_uint64_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t, __clc_uint64_t __global *, __clc_uint64_t const __local *, + __clc_int32_t, __clc_uint64_t __global *, __clc_uint64_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec2_uint64_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec2_uint64_t __local *, __clc_vec2_uint64_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec2_uint64_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec2_uint64_t __global *, __clc_vec2_uint64_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec3_uint64_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec3_uint64_t __local *, __clc_vec3_uint64_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec3_uint64_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec3_uint64_t __global *, __clc_vec3_uint64_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec4_uint64_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec4_uint64_t __local *, __clc_vec4_uint64_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec4_uint64_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec4_uint64_t __global *, __clc_vec4_uint64_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec8_uint64_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec8_uint64_t __local *, __clc_vec8_uint64_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec8_uint64_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec8_uint64_t __global *, __clc_vec8_uint64_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec16_uint64_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec16_uint64_t __local *, __clc_vec16_uint64_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec16_uint64_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec16_uint64_t __global *, __clc_vec16_uint64_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t, __clc_fp32_t __local *, __clc_fp32_t const __global *, + __clc_int32_t, __clc_fp32_t __local *, __clc_fp32_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t, __clc_fp32_t __global *, __clc_fp32_t const __local *, + __clc_int32_t, __clc_fp32_t __global *, __clc_fp32_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec2_fp32_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec2_fp32_t __local *, __clc_vec2_fp32_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec2_fp32_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec2_fp32_t __global *, __clc_vec2_fp32_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec3_fp32_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec3_fp32_t __local *, __clc_vec3_fp32_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec3_fp32_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec3_fp32_t __global *, __clc_vec3_fp32_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec4_fp32_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec4_fp32_t __local *, __clc_vec4_fp32_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec4_fp32_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec4_fp32_t __global *, __clc_vec4_fp32_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec8_fp32_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec8_fp32_t __local *, __clc_vec8_fp32_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec8_fp32_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec8_fp32_t __global *, __clc_vec8_fp32_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec16_fp32_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec16_fp32_t __local *, __clc_vec16_fp32_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec16_fp32_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec16_fp32_t __global *, __clc_vec16_fp32_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); #ifdef cl_khr_fp64 _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t, __clc_fp64_t __local *, __clc_fp64_t const __global *, + __clc_int32_t, __clc_fp64_t __local *, __clc_fp64_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t, __clc_fp64_t __global *, __clc_fp64_t const __local *, + __clc_int32_t, __clc_fp64_t __global *, __clc_fp64_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec2_fp64_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec2_fp64_t __local *, __clc_vec2_fp64_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec2_fp64_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec2_fp64_t __global *, __clc_vec2_fp64_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec3_fp64_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec3_fp64_t __local *, __clc_vec3_fp64_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec3_fp64_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec3_fp64_t __global *, __clc_vec3_fp64_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec4_fp64_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec4_fp64_t __local *, __clc_vec4_fp64_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec4_fp64_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec4_fp64_t __global *, __clc_vec4_fp64_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec8_fp64_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec8_fp64_t __local *, __clc_vec8_fp64_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec8_fp64_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec8_fp64_t __global *, __clc_vec8_fp64_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec16_fp64_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec16_fp64_t __local *, __clc_vec16_fp64_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec16_fp64_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec16_fp64_t __global *, __clc_vec16_fp64_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); #endif #ifdef cl_khr_fp16 _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t, __clc_fp16_t __local *, __clc_fp16_t const __global *, + __clc_int32_t, __clc_fp16_t __local *, __clc_fp16_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t, __clc_fp16_t __global *, __clc_fp16_t const __local *, + __clc_int32_t, __clc_fp16_t __global *, __clc_fp16_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec2_fp16_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec2_fp16_t __local *, __clc_vec2_fp16_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec2_fp16_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec2_fp16_t __global *, __clc_vec2_fp16_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec3_fp16_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec3_fp16_t __local *, __clc_vec3_fp16_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec3_fp16_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec3_fp16_t __global *, __clc_vec3_fp16_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec4_fp16_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec4_fp16_t __local *, __clc_vec4_fp16_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec4_fp16_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec4_fp16_t __global *, __clc_vec4_fp16_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec8_fp16_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec8_fp16_t __local *, __clc_vec8_fp16_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec8_fp16_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec8_fp16_t __global *, __clc_vec8_fp16_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec16_fp16_t __local *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec16_fp16_t __local *, __clc_vec16_fp16_t const __global *, __clc_size_t, __clc_size_t, __clc_event_t); _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT __clc_event_t -__spirv_GroupAsyncCopy(__clc_uint32_t, __clc_vec16_fp16_t __global *, +__spirv_GroupAsyncCopy(__clc_int32_t, __clc_vec16_fp16_t __global *, __clc_vec16_fp16_t const __local *, __clc_size_t, __clc_size_t, __clc_event_t); #endif _CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT void -__spirv_GroupWaitEvents(__clc_uint32_t, __clc_int32_t, __clc_event_t *); +__spirv_GroupWaitEvents(__clc_int32_t, __clc_int32_t, __clc_event_t *); _CLC_OVERLOAD _CLC_DECL _CLC_CONSTFN __clc_bool_t __spirv_IsFinite(__clc_fp32_t); diff --git a/libclc/libspirv/lib/amdgcn-amdhsa/async/wait_group_events.cl b/libclc/libspirv/lib/amdgcn-amdhsa/async/wait_group_events.cl index 701668c73e6b6..6437cc8eba4ac 100644 --- a/libclc/libspirv/lib/amdgcn-amdhsa/async/wait_group_events.cl +++ b/libclc/libspirv/lib/amdgcn-amdhsa/async/wait_group_events.cl @@ -8,7 +8,7 @@ #include -_CLC_OVERLOAD _CLC_DEF void __spirv_GroupWaitEvents(unsigned int scope, +_CLC_OVERLOAD _CLC_DEF void __spirv_GroupWaitEvents(int scope, int num_events, event_t *event_list) { __spirv_ControlBarrier(scope, Workgroup, SequentiallyConsistent); diff --git a/libclc/libspirv/lib/amdgcn-amdhsa/group/collectives.cl b/libclc/libspirv/lib/amdgcn-amdhsa/group/collectives.cl index 3aa99d8a9180c..eb1c1f615d52b 100644 --- a/libclc/libspirv/lib/amdgcn-amdhsa/group/collectives.cl +++ b/libclc/libspirv/lib/amdgcn-amdhsa/group/collectives.cl @@ -196,7 +196,7 @@ __CLC_SUBGROUP_COLLECTIVE(LogicalAndKHR, __CLC_LOGICAL_AND, bool, a, true) #define __CLC_GROUP_COLLECTIVE_INNER(SPIRV_NAME, CLC_NAME, OP, TYPE, IDENTITY) \ _CLC_DEF _CLC_OVERLOAD _CLC_CONVERGENT TYPE __CLC_APPEND( \ - __spirv_Group, SPIRV_NAME)(uint scope, uint op, TYPE x) { \ + __spirv_Group, SPIRV_NAME)(int scope, uint op, TYPE x) { \ TYPE carry = IDENTITY; \ /* Perform GroupOperation within sub-group */ \ TYPE sg_x = __CLC_APPEND(__clc__Subgroup, CLC_NAME)(op, x, &carry); \ @@ -253,11 +253,11 @@ __CLC_SUBGROUP_COLLECTIVE(LogicalAndKHR, __CLC_LOGICAL_AND, bool, a, true) __CLC_GROUP_COLLECTIVE(Any, __CLC_OR, bool, false); __CLC_GROUP_COLLECTIVE(All, __CLC_AND, bool, true); -_CLC_DEF _CLC_OVERLOAD _CLC_CONVERGENT bool __spirv_GroupAny(uint scope, +_CLC_DEF _CLC_OVERLOAD _CLC_CONVERGENT bool __spirv_GroupAny(int scope, bool predicate) { return __spirv_GroupAny(scope, Reduce, predicate); } -_CLC_DEF _CLC_OVERLOAD _CLC_CONVERGENT bool __spirv_GroupAll(uint scope, +_CLC_DEF _CLC_OVERLOAD _CLC_CONVERGENT bool __spirv_GroupAll(int scope, bool predicate) { return __spirv_GroupAll(scope, Reduce, predicate); } @@ -343,13 +343,13 @@ __CLC_GROUP_COLLECTIVE(LogicalAndKHR, __CLC_LOGICAL_AND, bool, true) // half requires additional mangled entry points #define __CLC_GROUP_COLLECTIVE__DF16(MANGLED_NAME, SPIRV_DISPATCH) \ - _CLC_DEF _CLC_CONVERGENT half MANGLED_NAME(uint scope, uint op, half x) { \ + _CLC_DEF _CLC_CONVERGENT half MANGLED_NAME(int scope, uint op, half x) { \ return SPIRV_DISPATCH(scope, op, x); \ } -__CLC_GROUP_COLLECTIVE__DF16(_Z17__spirv_GroupFAddjjDF16_, __spirv_GroupFAdd) -__CLC_GROUP_COLLECTIVE__DF16(_Z17__spirv_GroupFMinjjDF16_, __spirv_GroupFMin) -__CLC_GROUP_COLLECTIVE__DF16(_Z17__spirv_GroupFMaxjjDF16_, __spirv_GroupFMax) -__CLC_GROUP_COLLECTIVE__DF16(_Z20__spirv_GroupFMulKHRjjDF16_, +__CLC_GROUP_COLLECTIVE__DF16(_Z17__spirv_GroupFAddiiDF16_, __spirv_GroupFAdd) +__CLC_GROUP_COLLECTIVE__DF16(_Z17__spirv_GroupFMiniiDF16_, __spirv_GroupFMin) +__CLC_GROUP_COLLECTIVE__DF16(_Z17__spirv_GroupFMaxiiDF16_, __spirv_GroupFMax) +__CLC_GROUP_COLLECTIVE__DF16(_Z20__spirv_GroupFMulKHRiiDF16_, __spirv_GroupFMulKHR) #undef __CLC_GROUP_COLLECTIVE__DF16 @@ -380,7 +380,7 @@ long __clc__3d_to_linear_local_id(ulong3 id) { #define __CLC_GROUP_BROADCAST(TYPE, TYPE_MANGLED) \ _CLC_DEF _CLC_OVERLOAD _CLC_CONVERGENT TYPE __spirv_GroupBroadcast( \ - uint scope, TYPE x, ulong local_id) { \ + int scope, TYPE x, ulong local_id) { \ if (scope == Subgroup) { \ return _Z28__spirv_SubgroupShuffleINTELI##TYPE_MANGLED##ET_S0_j( \ x, local_id); \ @@ -396,17 +396,17 @@ long __clc__3d_to_linear_local_id(ulong3 id) { return result; \ } \ _CLC_DEF _CLC_OVERLOAD _CLC_CONVERGENT TYPE __spirv_GroupBroadcast( \ - uint scope, TYPE x, ulong2 local_id) { \ + int scope, TYPE x, ulong2 local_id) { \ ulong linear_local_id = __clc__2d_to_linear_local_id(local_id); \ return __spirv_GroupBroadcast(scope, x, linear_local_id); \ } \ _CLC_DEF _CLC_OVERLOAD _CLC_CONVERGENT TYPE __spirv_GroupBroadcast( \ - uint scope, TYPE x, ulong3 local_id) { \ + int scope, TYPE x, ulong3 local_id) { \ ulong linear_local_id = __clc__3d_to_linear_local_id(local_id); \ return __spirv_GroupBroadcast(scope, x, linear_local_id); \ } \ _CLC_DEF _CLC_OVERLOAD _CLC_CONVERGENT TYPE __spirv_GroupBroadcast( \ - uint scope, TYPE x, uint local_id) { \ + int scope, TYPE x, uint local_id) { \ return __spirv_GroupBroadcast(scope, x, (ulong)local_id); \ } __CLC_GROUP_BROADCAST(char, a); diff --git a/libclc/libspirv/lib/amdgcn-amdhsa/group/group_ballot.cl b/libclc/libspirv/lib/amdgcn-amdhsa/group/group_ballot.cl index 7208f788898da..49fa34a8cd7ab 100644 --- a/libclc/libspirv/lib/amdgcn-amdhsa/group/group_ballot.cl +++ b/libclc/libspirv/lib/amdgcn-amdhsa/group/group_ballot.cl @@ -13,7 +13,7 @@ #define ICMP_NE 33 _CLC_DEF _CLC_CONVERGENT __clc_vec4_uint32_t -_Z29__spirv_GroupNonUniformBallotjb(unsigned flag, bool predicate) { +_Z29__spirv_GroupNonUniformBallotib(int flag, bool predicate) { // only support subgroup for now if (flag != Subgroup) { __builtin_trap(); diff --git a/libclc/libspirv/lib/generic/async/async_work_group_strided_copy.inc b/libclc/libspirv/lib/generic/async/async_work_group_strided_copy.inc index 513c93b6f74e1..d85ea78d44a58 100644 --- a/libclc/libspirv/lib/generic/async/async_work_group_strided_copy.inc +++ b/libclc/libspirv/lib/generic/async/async_work_group_strided_copy.inc @@ -6,18 +6,16 @@ // //===----------------------------------------------------------------------===// -_CLC_OVERLOAD _CLC_DEF event_t -__spirv_GroupAsyncCopy(unsigned int scope, global __CLC_GENTYPE *dst, - const local __CLC_GENTYPE *src, size_t num_gentypes, - size_t stride, event_t event) { +_CLC_OVERLOAD _CLC_DEF event_t __spirv_GroupAsyncCopy( + int scope, global __CLC_GENTYPE *dst, const local __CLC_GENTYPE *src, + size_t num_gentypes, size_t stride, event_t event) { STRIDED_COPY(global, local, stride, 1); return event; } -_CLC_OVERLOAD _CLC_DEF event_t -__spirv_GroupAsyncCopy(unsigned int scope, local __CLC_GENTYPE *dst, - const global __CLC_GENTYPE *src, size_t num_gentypes, - size_t stride, event_t event) { +_CLC_OVERLOAD _CLC_DEF event_t __spirv_GroupAsyncCopy( + int scope, local __CLC_GENTYPE *dst, const global __CLC_GENTYPE *src, + size_t num_gentypes, size_t stride, event_t event) { STRIDED_COPY(local, global, 1, stride); return event; } diff --git a/libclc/libspirv/lib/generic/async/wait_group_events.cl b/libclc/libspirv/lib/generic/async/wait_group_events.cl index 9a6475a3f626c..10590e68d1e6f 100644 --- a/libclc/libspirv/lib/generic/async/wait_group_events.cl +++ b/libclc/libspirv/lib/generic/async/wait_group_events.cl @@ -8,7 +8,7 @@ #include -_CLC_OVERLOAD _CLC_DEF void __spirv_GroupWaitEvents(unsigned int scope, +_CLC_OVERLOAD _CLC_DEF void __spirv_GroupWaitEvents(int scope, int num_events, event_t *event_list) { __spirv_ControlBarrier(scope, Workgroup, SequentiallyConsistent); diff --git a/libclc/libspirv/lib/generic/float16.cl b/libclc/libspirv/lib/generic/float16.cl index 3315ba531d785..0799c8ab18a12 100644 --- a/libclc/libspirv/lib/generic/float16.cl +++ b/libclc/libspirv/lib/generic/float16.cl @@ -3210,7 +3210,7 @@ _CLC_OVERLOAD _CLC_DEF _CLC_CONSTFN __clc_vec16_int8_t __spirv_FUnordNotEqual( } _CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t args_0, __clc_float16_t __local *args_1, + __clc_int32_t args_0, __clc_float16_t __local *args_1, __clc_float16_t const __global *args_2, __clc_size_t args_3, __clc_size_t args_4, __clc_event_t args_5) { return __spirv_GroupAsyncCopy(args_0, (__clc_fp16_t __local *)(args_1), @@ -3219,7 +3219,7 @@ _CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( } _CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t args_0, __clc_float16_t __global *args_1, + __clc_int32_t args_0, __clc_float16_t __global *args_1, __clc_float16_t const __local *args_2, __clc_size_t args_3, __clc_size_t args_4, __clc_event_t args_5) { return __spirv_GroupAsyncCopy(args_0, (__clc_fp16_t __global *)(args_1), @@ -3228,7 +3228,7 @@ _CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( } _CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t args_0, __clc_vec2_float16_t __local *args_1, + __clc_int32_t args_0, __clc_vec2_float16_t __local *args_1, __clc_vec2_float16_t const __global *args_2, __clc_size_t args_3, __clc_size_t args_4, __clc_event_t args_5) { return __spirv_GroupAsyncCopy(args_0, (__clc_vec2_fp16_t __local *)(args_1), @@ -3237,7 +3237,7 @@ _CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( } _CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t args_0, __clc_vec2_float16_t __global *args_1, + __clc_int32_t args_0, __clc_vec2_float16_t __global *args_1, __clc_vec2_float16_t const __local *args_2, __clc_size_t args_3, __clc_size_t args_4, __clc_event_t args_5) { return __spirv_GroupAsyncCopy(args_0, (__clc_vec2_fp16_t __global *)(args_1), @@ -3246,7 +3246,7 @@ _CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( } _CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t args_0, __clc_vec3_float16_t __local *args_1, + __clc_int32_t args_0, __clc_vec3_float16_t __local *args_1, __clc_vec3_float16_t const __global *args_2, __clc_size_t args_3, __clc_size_t args_4, __clc_event_t args_5) { return __spirv_GroupAsyncCopy(args_0, (__clc_vec3_fp16_t __local *)(args_1), @@ -3255,7 +3255,7 @@ _CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( } _CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t args_0, __clc_vec3_float16_t __global *args_1, + __clc_int32_t args_0, __clc_vec3_float16_t __global *args_1, __clc_vec3_float16_t const __local *args_2, __clc_size_t args_3, __clc_size_t args_4, __clc_event_t args_5) { return __spirv_GroupAsyncCopy(args_0, (__clc_vec3_fp16_t __global *)(args_1), @@ -3264,7 +3264,7 @@ _CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( } _CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t args_0, __clc_vec4_float16_t __local *args_1, + __clc_int32_t args_0, __clc_vec4_float16_t __local *args_1, __clc_vec4_float16_t const __global *args_2, __clc_size_t args_3, __clc_size_t args_4, __clc_event_t args_5) { return __spirv_GroupAsyncCopy(args_0, (__clc_vec4_fp16_t __local *)(args_1), @@ -3273,7 +3273,7 @@ _CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( } _CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t args_0, __clc_vec4_float16_t __global *args_1, + __clc_int32_t args_0, __clc_vec4_float16_t __global *args_1, __clc_vec4_float16_t const __local *args_2, __clc_size_t args_3, __clc_size_t args_4, __clc_event_t args_5) { return __spirv_GroupAsyncCopy(args_0, (__clc_vec4_fp16_t __global *)(args_1), @@ -3282,7 +3282,7 @@ _CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( } _CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t args_0, __clc_vec8_float16_t __local *args_1, + __clc_int32_t args_0, __clc_vec8_float16_t __local *args_1, __clc_vec8_float16_t const __global *args_2, __clc_size_t args_3, __clc_size_t args_4, __clc_event_t args_5) { return __spirv_GroupAsyncCopy(args_0, (__clc_vec8_fp16_t __local *)(args_1), @@ -3291,7 +3291,7 @@ _CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( } _CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t args_0, __clc_vec8_float16_t __global *args_1, + __clc_int32_t args_0, __clc_vec8_float16_t __global *args_1, __clc_vec8_float16_t const __local *args_2, __clc_size_t args_3, __clc_size_t args_4, __clc_event_t args_5) { return __spirv_GroupAsyncCopy(args_0, (__clc_vec8_fp16_t __global *)(args_1), @@ -3300,7 +3300,7 @@ _CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( } _CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t args_0, __clc_vec16_float16_t __local *args_1, + __clc_int32_t args_0, __clc_vec16_float16_t __local *args_1, __clc_vec16_float16_t const __global *args_2, __clc_size_t args_3, __clc_size_t args_4, __clc_event_t args_5) { return __spirv_GroupAsyncCopy(args_0, (__clc_vec16_fp16_t __local *)(args_1), @@ -3309,7 +3309,7 @@ _CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( } _CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT __clc_event_t __spirv_GroupAsyncCopy( - __clc_uint32_t args_0, __clc_vec16_float16_t __global *args_1, + __clc_int32_t args_0, __clc_vec16_float16_t __global *args_1, __clc_vec16_float16_t const __local *args_2, __clc_size_t args_3, __clc_size_t args_4, __clc_event_t args_5) { return __spirv_GroupAsyncCopy(args_0, (__clc_vec16_fp16_t __global *)(args_1), diff --git a/libclc/libspirv/lib/ptx-nvidiacl/async/async_work_group_strided_copy.cl b/libclc/libspirv/lib/ptx-nvidiacl/async/async_work_group_strided_copy.cl index ad8041eec3e5e..438ad9ccf6ab7 100644 --- a/libclc/libspirv/lib/ptx-nvidiacl/async/async_work_group_strided_copy.cl +++ b/libclc/libspirv/lib/ptx-nvidiacl/async/async_work_group_strided_copy.cl @@ -18,7 +18,7 @@ int __clc_nvvm_reflect_arch(); #define __CLC_GROUP_CP_ASYNC_DST_GLOBAL(TYPE) \ _CLC_OVERLOAD _CLC_DEF event_t __spirv_GroupAsyncCopy( \ - unsigned int scope, __attribute__((address_space(1))) TYPE *dst, \ + int scope, __attribute__((address_space(1))) TYPE *dst, \ const __attribute__((address_space(3))) TYPE *src, size_t num_gentypes, \ size_t stride, event_t event) { \ STRIDED_COPY(__attribute__((address_space(1))), \ @@ -64,7 +64,7 @@ __CLC_GROUP_CP_ASYNC_DST_GLOBAL(uchar16); #define __CLC_GROUP_CP_ASYNC_4(TYPE) \ _CLC_DEF _CLC_OVERLOAD _CLC_CONVERGENT event_t __spirv_GroupAsyncCopy( \ - unsigned int scope, __attribute__((address_space(3))) TYPE *dst, \ + int scope, __attribute__((address_space(3))) TYPE *dst, \ const __attribute__((address_space(1))) TYPE *src, size_t num_gentypes, \ size_t stride, event_t event) { \ if (__clc_nvvm_reflect_arch() >= 800) { \ @@ -95,7 +95,7 @@ __CLC_GROUP_CP_ASYNC_4(uchar4); #define __CLC_GROUP_CP_ASYNC_8(TYPE) \ _CLC_DEF _CLC_OVERLOAD _CLC_CONVERGENT event_t __spirv_GroupAsyncCopy( \ - unsigned int scope, __attribute__((address_space(3))) TYPE *dst, \ + int scope, __attribute__((address_space(3))) TYPE *dst, \ const __attribute__((address_space(1))) TYPE *src, size_t num_gentypes, \ size_t stride, event_t event) { \ if (__clc_nvvm_reflect_arch() >= 800) { \ @@ -129,7 +129,7 @@ __CLC_GROUP_CP_ASYNC_8(uchar8); #define __CLC_GROUP_CP_ASYNC_16(TYPE) \ _CLC_DEF _CLC_OVERLOAD _CLC_CONVERGENT event_t __spirv_GroupAsyncCopy( \ - unsigned int scope, __attribute__((address_space(3))) TYPE *dst, \ + int scope, __attribute__((address_space(3))) TYPE *dst, \ const __attribute__((address_space(1))) TYPE *src, size_t num_gentypes, \ size_t stride, event_t event) { \ if (__clc_nvvm_reflect_arch() >= 800) { \ diff --git a/libclc/libspirv/lib/ptx-nvidiacl/async/wait_group_events.cl b/libclc/libspirv/lib/ptx-nvidiacl/async/wait_group_events.cl index 52dfd19a5e508..ae329324fd935 100644 --- a/libclc/libspirv/lib/ptx-nvidiacl/async/wait_group_events.cl +++ b/libclc/libspirv/lib/ptx-nvidiacl/async/wait_group_events.cl @@ -11,7 +11,7 @@ int __clc_nvvm_reflect_arch(); -_CLC_OVERLOAD _CLC_DEF void __spirv_GroupWaitEvents(unsigned int scope, +_CLC_OVERLOAD _CLC_DEF void __spirv_GroupWaitEvents(int scope, int num_events, event_t *event_list) { if (__clc_nvvm_reflect_arch() >= 800) { diff --git a/libclc/libspirv/lib/ptx-nvidiacl/group/collectives.cl b/libclc/libspirv/lib/ptx-nvidiacl/group/collectives.cl index 13b4c1f96354a..6e090d68178c3 100644 --- a/libclc/libspirv/lib/ptx-nvidiacl/group/collectives.cl +++ b/libclc/libspirv/lib/ptx-nvidiacl/group/collectives.cl @@ -204,13 +204,13 @@ __clc__SubgroupShuffleUp(complex_double x, uint delta) { // TODO: Implement InclusiveScan/ExclusiveScan // Currently only Reduce is required (for GroupAny and GroupAll) _CLC_DEF _CLC_OVERLOAD _CLC_CONVERGENT bool -__clc__SubgroupBitwiseOr(uint op, bool predicate, bool *carry) { +__clc__SubgroupBitwiseOr(int op, bool predicate, bool *carry) { bool result = __nvvm_vote_any_sync(__clc__membermask(), predicate); *carry = result; return result; } _CLC_DEF _CLC_OVERLOAD _CLC_CONVERGENT bool -__clc__SubgroupBitwiseAny(uint op, bool predicate, bool *carry) { +__clc__SubgroupBitwiseAny(int op, bool predicate, bool *carry) { bool result = __nvvm_vote_all_sync(__clc__membermask(), predicate); *carry = result; return result; @@ -477,7 +477,7 @@ __CLC_SUBGROUP_COLLECTIVE(LogicalAndKHR, __CLC_LOGICAL_AND, bool, true) #define __CLC_GROUP_COLLECTIVE_OUTER(SPIRV_NAME, CLC_NAME, OP, TYPE, IDENTITY) \ _CLC_DEF _CLC_OVERLOAD _CLC_CONVERGENT TYPE __CLC_APPEND( \ - __spirv_Group, SPIRV_NAME)(uint scope, uint op, TYPE x) { \ + __spirv_Group, SPIRV_NAME)(int scope, int op, TYPE x) { \ __CLC_GROUP_COLLECTIVE_INNER(CLC_NAME, OP, TYPE, IDENTITY) \ } @@ -495,18 +495,18 @@ __CLC_SUBGROUP_COLLECTIVE(LogicalAndKHR, __CLC_LOGICAL_AND, bool, true) #define __CLC_GROUP_COLLECTIVE_MANUAL_MANGLE(SPIRV_NAME_MANGLED, CLC_NAME, OP, \ TYPE, IDENTITY) \ - _CLC_DEF _CLC_CONVERGENT TYPE SPIRV_NAME_MANGLED(uint scope, uint op, \ + _CLC_DEF _CLC_CONVERGENT TYPE SPIRV_NAME_MANGLED(int scope, int op, \ TYPE x) { \ __CLC_GROUP_COLLECTIVE_INNER(CLC_NAME, OP, TYPE, IDENTITY) \ } __CLC_GROUP_COLLECTIVE(BitwiseOr, __CLC_OR, bool, false); __CLC_GROUP_COLLECTIVE(BitwiseAny, __CLC_AND, bool, true); -_CLC_DEF _CLC_OVERLOAD _CLC_CONVERGENT bool __spirv_GroupAny(uint scope, +_CLC_DEF _CLC_OVERLOAD _CLC_CONVERGENT bool __spirv_GroupAny(int scope, bool predicate) { return __spirv_GroupBitwiseOr(scope, Reduce, predicate); } -_CLC_DEF _CLC_OVERLOAD _CLC_CONVERGENT bool __spirv_GroupAll(uint scope, +_CLC_DEF _CLC_OVERLOAD _CLC_CONVERGENT bool __spirv_GroupAll(int scope, bool predicate) { return __spirv_GroupBitwiseAny(scope, Reduce, predicate); } @@ -536,13 +536,13 @@ __CLC_GROUP_COLLECTIVE(FMulKHR, __CLC_MUL, float, 1) __CLC_GROUP_COLLECTIVE(FMulKHR, __CLC_MUL, double, 1) __CLC_GROUP_COLLECTIVE_MANUAL_MANGLE( - _Z22__spirv_GroupCMulINTELjjN5__spv12complex_halfE, CMulINTEL, + _Z22__spirv_GroupCMulINTELiiN5__spv12complex_halfE, CMulINTEL, __CLC_COMPLEX_MUL, complex_half, ((complex_half){1, 0})) __CLC_GROUP_COLLECTIVE_MANUAL_MANGLE( - _Z22__spirv_GroupCMulINTELjjN5__spv13complex_floatE, CMulINTEL, + _Z22__spirv_GroupCMulINTELiiN5__spv13complex_floatE, CMulINTEL, __CLC_COMPLEX_MUL, complex_float, ((complex_float){1, 0})) __CLC_GROUP_COLLECTIVE_MANUAL_MANGLE( - _Z22__spirv_GroupCMulINTELjjN5__spv14complex_doubleE, CMulINTEL, + _Z22__spirv_GroupCMulINTELiiN5__spv14complex_doubleE, CMulINTEL, __CLC_COMPLEX_MUL, complex_double, ((complex_double){1, 0})) __CLC_GROUP_COLLECTIVE(SMin, __CLC_MIN, char, CHAR_MAX) @@ -602,13 +602,13 @@ __CLC_GROUP_COLLECTIVE(LogicalAndKHR, __CLC_LOGICAL_AND, bool, true) // half requires additional mangled entry points #define __CLC_GROUP_COLLECTIVE__DF16(MANGLED_NAME, SPIRV_DISPATCH) \ - _CLC_DEF _CLC_CONVERGENT half MANGLED_NAME(uint scope, uint op, half x) { \ + _CLC_DEF _CLC_CONVERGENT half MANGLED_NAME(int scope, uint op, half x) { \ return SPIRV_DISPATCH(scope, op, x); \ } -__CLC_GROUP_COLLECTIVE__DF16(_Z17__spirv_GroupFAddjjDF16_, __spirv_GroupFAdd) -__CLC_GROUP_COLLECTIVE__DF16(_Z17__spirv_GroupFMinjjDF16_, __spirv_GroupFMin) -__CLC_GROUP_COLLECTIVE__DF16(_Z17__spirv_GroupFMaxjjDF16_, __spirv_GroupFMax) -__CLC_GROUP_COLLECTIVE__DF16(_Z20__spirv_GroupFMulKHRjjDF16_, +__CLC_GROUP_COLLECTIVE__DF16(_Z17__spirv_GroupFAddiiDF16_, __spirv_GroupFAdd) +__CLC_GROUP_COLLECTIVE__DF16(_Z17__spirv_GroupFMiniiDF16_, __spirv_GroupFMin) +__CLC_GROUP_COLLECTIVE__DF16(_Z17__spirv_GroupFMaxiiDF16_, __spirv_GroupFMax) +__CLC_GROUP_COLLECTIVE__DF16(_Z20__spirv_GroupFMulKHRiiDF16_, __spirv_GroupFMulKHR) #undef __CLC_GROUP_COLLECTIVE__DF16 @@ -639,7 +639,7 @@ long __clc__3d_to_linear_local_id(ulong3 id) { #define __CLC_GROUP_BROADCAST(TYPE) \ _CLC_DEF _CLC_OVERLOAD _CLC_CONVERGENT TYPE __spirv_GroupBroadcast( \ - uint scope, TYPE x, ulong local_id) { \ + int scope, TYPE x, ulong local_id) { \ if (scope == Subgroup) { \ return __clc__SubgroupShuffle(x, local_id); \ } \ @@ -654,17 +654,17 @@ long __clc__3d_to_linear_local_id(ulong3 id) { return result; \ } \ _CLC_DEF _CLC_OVERLOAD _CLC_CONVERGENT TYPE __spirv_GroupBroadcast( \ - uint scope, TYPE x, ulong2 local_id) { \ + int scope, TYPE x, ulong2 local_id) { \ ulong linear_local_id = __clc__2d_to_linear_local_id(local_id); \ return __spirv_GroupBroadcast(scope, x, linear_local_id); \ } \ _CLC_DEF _CLC_OVERLOAD _CLC_CONVERGENT TYPE __spirv_GroupBroadcast( \ - uint scope, TYPE x, ulong3 local_id) { \ + int scope, TYPE x, ulong3 local_id) { \ ulong linear_local_id = __clc__3d_to_linear_local_id(local_id); \ return __spirv_GroupBroadcast(scope, x, linear_local_id); \ } \ _CLC_DEF _CLC_OVERLOAD _CLC_CONVERGENT TYPE __spirv_GroupBroadcast( \ - uint scope, TYPE x, uint local_id) { \ + int scope, TYPE x, uint local_id) { \ return __spirv_GroupBroadcast(scope, x, (ulong)local_id); \ } __CLC_GROUP_BROADCAST(char); @@ -681,18 +681,18 @@ __CLC_GROUP_BROADCAST(double) // half requires additional mangled entry points _CLC_DEF _CLC_CONVERGENT half -_Z17__spirv_GroupBroadcastjDF16_m(uint scope, half x, ulong local_id) { +_Z17__spirv_GroupBroadcastiDF16_m(int scope, half x, ulong local_id) { return __spirv_GroupBroadcast(scope, x, local_id); } _CLC_DEF _CLC_CONVERGENT half -_Z17__spirv_GroupBroadcastjDF16_Dv2_m(uint scope, half x, ulong2 local_id) { +_Z17__spirv_GroupBroadcastiDF16_Dv2_m(int scope, half x, ulong2 local_id) { return __spirv_GroupBroadcast(scope, x, local_id); } _CLC_DEF _CLC_CONVERGENT half -_Z17__spirv_GroupBroadcastjDF16_Dv3_m(uint scope, half x, ulong3 local_id) { +_Z17__spirv_GroupBroadcastiDF16_Dv3_m(int scope, half x, ulong3 local_id) { return __spirv_GroupBroadcast(scope, x, local_id); } -_CLC_DEF _CLC_CONVERGENT half _Z22__spirv_GroupBroadcastjDF16_j(uint scope, +_CLC_DEF _CLC_CONVERGENT half _Z22__spirv_GroupBroadcastiDF16_j(int scope, half x, uint local_id) { return __spirv_GroupBroadcast(scope, x, (ulong)local_id); diff --git a/libclc/libspirv/lib/ptx-nvidiacl/group/group_non_uniform.cl b/libclc/libspirv/lib/ptx-nvidiacl/group/group_non_uniform.cl index 671dbb4535d76..3474f17666080 100644 --- a/libclc/libspirv/lib/ptx-nvidiacl/group/group_non_uniform.cl +++ b/libclc/libspirv/lib/ptx-nvidiacl/group/group_non_uniform.cl @@ -13,7 +13,7 @@ #include _CLC_DEF _CLC_CONVERGENT __clc_vec4_uint32_t -_Z29__spirv_GroupNonUniformBallotjb(unsigned flag, bool predicate) { +_Z29__spirv_GroupNonUniformBallotib(int flag, bool predicate) { // only support subgroup for now if (flag != Subgroup) { __builtin_trap(); @@ -38,7 +38,7 @@ _Z29__spirv_GroupNonUniformBallotjb(unsigned flag, bool predicate) { _CLC_DEF _CLC_CONVERGENT uint _Z37__spirv_GroupNonUniformBallotBitCountN5__spv5Scope4FlagEiDv4_j( - uint scope, uint flag, __clc_vec4_uint32_t mask) { + int scope, int flag, __clc_vec4_uint32_t mask) { // here we assume scope == __spv::Scope::Subgroup // flag == InclusiveScan is not yet implemented if (flag == Reduce) { diff --git a/libdevice/nativecpu_utils.cpp b/libdevice/nativecpu_utils.cpp index eb5c3ff2ebdf4..b36c3f8344e0c 100644 --- a/libdevice/nativecpu_utils.cpp +++ b/libdevice/nativecpu_utils.cpp @@ -103,7 +103,7 @@ DefSubgroupBlockINTEL(uint8_t) DefSubgroupBlockINTEL(uint16_t) #define DefineGOp1(spir_sfx, name)\ DEVICE_EXTERN_C bool __mux_sub_group_##name##_i1(bool);\ DEVICE_EXTERN_C bool __mux_work_group_##name##_i1(uint32_t id, bool val);\ -DEVICE_EXTERNAL bool __spirv_Group ## spir_sfx(unsigned g, bool val) {\ +DEVICE_EXTERNAL bool __spirv_Group ## spir_sfx(int32_t g, bool val) {\ if (__spv::Scope::Flag::Subgroup == g)\ return __mux_sub_group_##name##_i1(val);\ else if (__spv::Scope::Flag::Workgroup == g)\ @@ -124,7 +124,7 @@ DefineGOp1(All, all) DEVICE_EXTERN_C MuxType __mux_work_group_scan_inclusive_##mux_sfx(uint32_t, \ MuxType); \ DEVICE_EXTERN_C MuxType __mux_work_group_reduce_##mux_sfx(uint32_t, MuxType);\ - DEVICE_EXTERNAL Type __spirv_Group##spir_sfx(uint32_t g, uint32_t id, \ + DEVICE_EXTERNAL Type __spirv_Group##spir_sfx(int32_t g, int32_t id, \ Type v) { \ if (__spv::Scope::Flag::Subgroup == g) { \ if (static_cast(__spv::GroupOperation::InclusiveScan) == id) \ @@ -202,7 +202,7 @@ DefineLogicalGroupOp(bool, bool, i1) int32_t sg_lid); #define DefineBroadCastImpl(Type, Sfx, MuxType, IDType) \ - DEVICE_EXTERNAL Type __spirv_GroupBroadcast(uint32_t g, Type v, \ + DEVICE_EXTERNAL Type __spirv_GroupBroadcast(int32_t g, Type v, \ IDType l) { \ if (__spv::Scope::Flag::Subgroup == g) \ return __mux_sub_group_broadcast_##Sfx(v, l); \ @@ -210,7 +210,7 @@ DefineLogicalGroupOp(bool, bool, i1) return __mux_work_group_broadcast_##Sfx(0, v, l, 0, 0); \ } \ \ - DEVICE_EXTERNAL Type __spirv_GroupBroadcast(uint32_t g, Type v, \ + DEVICE_EXTERNAL Type __spirv_GroupBroadcast(int32_t g, Type v, \ sycl::vec::vector_t l) { \ if (__spv::Scope::Flag::Subgroup == g) \ return __mux_sub_group_broadcast_##Sfx(v, l[0]); \ @@ -218,7 +218,7 @@ DefineLogicalGroupOp(bool, bool, i1) return __mux_work_group_broadcast_##Sfx(0, v, l[0], l[1], 0); \ } \ \ - DEVICE_EXTERNAL Type __spirv_GroupBroadcast(uint32_t g, Type v, \ + DEVICE_EXTERNAL Type __spirv_GroupBroadcast(int32_t g, Type v, \ sycl::vec::vector_t l) { \ if (__spv::Scope::Flag::Subgroup == g) \ return __mux_sub_group_broadcast_##Sfx(v, l[0]); \ diff --git a/sycl/include/sycl/__spirv/spirv_ops.hpp b/sycl/include/sycl/__spirv/spirv_ops.hpp index 5800190f539a0..79271a8d88a93 100644 --- a/sycl/include/sycl/__spirv/spirv_ops.hpp +++ b/sycl/include/sycl/__spirv/spirv_ops.hpp @@ -879,7 +879,7 @@ extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT __ocl_vec_t - __spirv_GroupNonUniformBallot(uint32_t Execution, bool Predicate) noexcept; + __spirv_GroupNonUniformBallot(int32_t Execution, bool Predicate) noexcept; // TODO: I'm not 100% sure that these NonUniform instructions should be // convergent Following precedent set for GroupNonUniformBallot above @@ -907,137 +907,125 @@ __spirv_GroupNonUniformAny(__spv::Scope::Flag, bool); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformSMin(__spv::Scope::Flag, unsigned int, ValueT); +__spirv_GroupNonUniformSMin(__spv::Scope::Flag, int, ValueT); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformUMin(__spv::Scope::Flag, unsigned int, ValueT); +__spirv_GroupNonUniformUMin(__spv::Scope::Flag, int, ValueT); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformFMin(__spv::Scope::Flag, unsigned int, ValueT); +__spirv_GroupNonUniformFMin(__spv::Scope::Flag, int, ValueT); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformSMax(__spv::Scope::Flag, unsigned int, ValueT); +__spirv_GroupNonUniformSMax(__spv::Scope::Flag, int, ValueT); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformUMax(__spv::Scope::Flag, unsigned int, ValueT); +__spirv_GroupNonUniformUMax(__spv::Scope::Flag, int, ValueT); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformFMax(__spv::Scope::Flag, unsigned int, ValueT); +__spirv_GroupNonUniformFMax(__spv::Scope::Flag, int, ValueT); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformIAdd(__spv::Scope::Flag, unsigned int, ValueT); +__spirv_GroupNonUniformIAdd(__spv::Scope::Flag, int, ValueT); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformFAdd(__spv::Scope::Flag, unsigned int, ValueT); +__spirv_GroupNonUniformFAdd(__spv::Scope::Flag, int, ValueT); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformIMul(__spv::Scope::Flag, unsigned int, ValueT); +__spirv_GroupNonUniformIMul(__spv::Scope::Flag, int, ValueT); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformFMul(__spv::Scope::Flag, unsigned int, ValueT); +__spirv_GroupNonUniformFMul(__spv::Scope::Flag, int, ValueT); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformBitwiseOr(__spv::Scope::Flag, unsigned int, ValueT); +__spirv_GroupNonUniformBitwiseOr(__spv::Scope::Flag, int, ValueT); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformBitwiseXor(__spv::Scope::Flag, unsigned int, ValueT); +__spirv_GroupNonUniformBitwiseXor(__spv::Scope::Flag, int, ValueT); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformBitwiseAnd(__spv::Scope::Flag, unsigned int, ValueT); +__spirv_GroupNonUniformBitwiseAnd(__spv::Scope::Flag, int, ValueT); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformLogicalOr(__spv::Scope::Flag, unsigned int, ValueT); +__spirv_GroupNonUniformLogicalOr(__spv::Scope::Flag, int, ValueT); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformLogicalAnd(__spv::Scope::Flag, unsigned int, ValueT); +__spirv_GroupNonUniformLogicalAnd(__spv::Scope::Flag, int, ValueT); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformSMin(__spv::Scope::Flag, unsigned int, ValueT, - unsigned int); +__spirv_GroupNonUniformSMin(__spv::Scope::Flag, int, ValueT, unsigned int); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformUMin(__spv::Scope::Flag, unsigned int, ValueT, - unsigned int); +__spirv_GroupNonUniformUMin(__spv::Scope::Flag, int, ValueT, unsigned int); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformFMin(__spv::Scope::Flag, unsigned int, ValueT, - unsigned int); +__spirv_GroupNonUniformFMin(__spv::Scope::Flag, int, ValueT, unsigned int); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformSMax(__spv::Scope::Flag, unsigned int, ValueT, - unsigned int); +__spirv_GroupNonUniformSMax(__spv::Scope::Flag, int, ValueT, unsigned int); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformUMax(__spv::Scope::Flag, unsigned int, ValueT, - unsigned int); +__spirv_GroupNonUniformUMax(__spv::Scope::Flag, int, ValueT, unsigned int); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformFMax(__spv::Scope::Flag, unsigned int, ValueT, - unsigned int); +__spirv_GroupNonUniformFMax(__spv::Scope::Flag, int, ValueT, unsigned int); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformIAdd(__spv::Scope::Flag, unsigned int, ValueT, - unsigned int); +__spirv_GroupNonUniformIAdd(__spv::Scope::Flag, int, ValueT, unsigned int); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformFAdd(__spv::Scope::Flag, unsigned int, ValueT, - unsigned int); +__spirv_GroupNonUniformFAdd(__spv::Scope::Flag, int, ValueT, unsigned int); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformIMul(__spv::Scope::Flag, unsigned int, ValueT, - unsigned int); +__spirv_GroupNonUniformIMul(__spv::Scope::Flag, int, ValueT, unsigned int); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformFMul(__spv::Scope::Flag, unsigned int, ValueT, - unsigned int); +__spirv_GroupNonUniformFMul(__spv::Scope::Flag, int, ValueT, unsigned int); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformBitwiseOr(__spv::Scope::Flag, unsigned int, ValueT, - unsigned int); +__spirv_GroupNonUniformBitwiseOr(__spv::Scope::Flag, int, ValueT, unsigned int); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformBitwiseXor(__spv::Scope::Flag, unsigned int, ValueT, +__spirv_GroupNonUniformBitwiseXor(__spv::Scope::Flag, int, ValueT, unsigned int); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformBitwiseAnd(__spv::Scope::Flag, unsigned int, ValueT, +__spirv_GroupNonUniformBitwiseAnd(__spv::Scope::Flag, int, ValueT, unsigned int); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformLogicalOr(__spv::Scope::Flag, unsigned int, ValueT, - unsigned int); +__spirv_GroupNonUniformLogicalOr(__spv::Scope::Flag, int, ValueT, unsigned int); template __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT ValueT -__spirv_GroupNonUniformLogicalAnd(__spv::Scope::Flag, unsigned int, ValueT, +__spirv_GroupNonUniformLogicalAnd(__spv::Scope::Flag, int, ValueT, unsigned int); extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT void @@ -1204,9 +1192,9 @@ extern __DPCPP_SYCL_EXTERNAL void __spirv_TaskSequenceReleaseINTEL( template __SYCL_CONVERGENT__ extern __ocl_event_t -__SYCL_OpGroupAsyncCopyGlobalToLocal(__spv::Scope::Flag, dataT *Dest, - const dataT *Src, size_t NumElements, - size_t Stride, __ocl_event_t) noexcept { +__SYCL_OpGroupAsyncCopyGlobalToLocal(int32_t, dataT *Dest, const dataT *Src, + size_t NumElements, size_t Stride, + __ocl_event_t) noexcept { for (size_t i = 0; i < NumElements; i++) { Dest[i] = Src[i * Stride]; } @@ -1216,9 +1204,9 @@ __SYCL_OpGroupAsyncCopyGlobalToLocal(__spv::Scope::Flag, dataT *Dest, template __SYCL_CONVERGENT__ extern __ocl_event_t -__SYCL_OpGroupAsyncCopyLocalToGlobal(__spv::Scope::Flag, dataT *Dest, - const dataT *Src, size_t NumElements, - size_t Stride, __ocl_event_t) noexcept { +__SYCL_OpGroupAsyncCopyLocalToGlobal(int32_t, dataT *Dest, const dataT *Src, + size_t NumElements, size_t Stride, + __ocl_event_t) noexcept { for (size_t i = 0; i < NumElements; i++) { Dest[i * Stride] = Src[i]; } @@ -1237,6 +1225,6 @@ __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT void __spirv_MemoryBarrier(__spv::Scope Memory, uint32_t Semantics) noexcept; __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT void -__spirv_GroupWaitEvents(__spv::Scope Execution, uint32_t NumEvents, +__spirv_GroupWaitEvents(int32_t Execution, uint32_t NumEvents, __ocl_event_t *WaitEvents) noexcept; #endif // !__SYCL_DEVICE_ONLY__ diff --git a/sycl/include/sycl/__spirv/spirv_types.hpp b/sycl/include/sycl/__spirv/spirv_types.hpp index 3a082b3fdf3aa..1aeb4cdd88c50 100644 --- a/sycl/include/sycl/__spirv/spirv_types.hpp +++ b/sycl/include/sycl/__spirv/spirv_types.hpp @@ -24,7 +24,7 @@ namespace __spv { struct Scope { - enum Flag : uint32_t { + enum Flag : int32_t { CrossDevice = 0, Device = 1, Workgroup = 2, @@ -34,13 +34,13 @@ struct Scope { constexpr Scope(Flag flag) : flag_value(flag) {} - constexpr operator uint32_t() const { return flag_value; } + constexpr operator int32_t() const { return flag_value; } Flag flag_value; }; struct StorageClass { - enum Flag : uint32_t { + enum Flag : int32_t { UniformConstant = 0, Input = 1, Uniform = 2, @@ -75,13 +75,13 @@ struct StorageClass { Max = 0x7fffffff, }; constexpr StorageClass(Flag flag) : flag_value(flag) {} - constexpr operator uint32_t() const { return flag_value; } + constexpr operator int32_t() const { return flag_value; } Flag flag_value; }; struct MemorySemanticsMask { - enum Flag : uint32_t { + enum Flag : int32_t { None = 0x0, Acquire = 0x2, Release = 0x4, @@ -97,12 +97,12 @@ struct MemorySemanticsMask { constexpr MemorySemanticsMask(Flag flag) : flag_value(flag) {} - constexpr operator uint32_t() const { return flag_value; } + constexpr operator int32_t() const { return flag_value; } Flag flag_value; }; -enum class GroupOperation : uint32_t { +enum class GroupOperation : int32_t { Reduce = 0, InclusiveScan = 1, ExclusiveScan = 2, diff --git a/sycl/source/spirv_ops.cpp b/sycl/source/spirv_ops.cpp index fdaa7e1834eae..23bde9f4a2ddd 100644 --- a/sycl/source/spirv_ops.cpp +++ b/sycl/source/spirv_ops.cpp @@ -16,7 +16,7 @@ // This operation is NOP on HOST as all operations there are blocking and // by the moment this function was called, the operations generating // the __ocl_event_t objects had already been finished. -__SYCL_EXPORT void __spirv_GroupWaitEvents(__spv::Scope Execution, +__SYCL_EXPORT void __spirv_GroupWaitEvents(int32_t Execution, uint32_t NumEvents, __ocl_event_t *WaitEvents) noexcept { (void)Execution; diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index cbe08c9c12d4f..f2394a2372be1 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -10,7 +10,7 @@ _Z20__spirv_ocl_prefetchPKcm _Z21__spirv_MemoryBarrierN5__spv5ScopeEj _Z22__spirv_ControlBarrierN5__spv5ScopeES0_j -_Z23__spirv_GroupWaitEventsN5__spv5ScopeEjPPv +_Z23__spirv_GroupWaitEventsijPPv _ZN4sycl3_V110__abs_implENS0_3vecIaLi16EEE _ZN4sycl3_V110__abs_implENS0_3vecIaLi1EEE _ZN4sycl3_V110__abs_implENS0_3vecIaLi2EEE diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 06907b875aac5..55426827a04f8 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -3473,7 +3473,7 @@ ?__smoothstep_impl@_V1@sycl@@YAMMMM@Z ?__smoothstep_impl@_V1@sycl@@YANNNN@Z ?__spirv_ControlBarrier@@YAXUScope@__spv@@0I@Z -?__spirv_GroupWaitEvents@@YAXUScope@__spv@@IPEAPEAX@Z +?__spirv_GroupWaitEvents@@YAXHIPEAPEAX@Z ?__spirv_MemoryBarrier@@YAXUScope@__spv@@I@Z ?__spirv_ocl_prefetch@@YAXPEBD_K@Z ?__sqrt_impl@_V1@sycl@@YA?AV?$vec@M$00@12@V312@@Z diff --git a/sycl/test/check_device_code/group_operations.cpp b/sycl/test/check_device_code/group_operations.cpp index 7145b38c4f254..83b26c118ae4d 100644 --- a/sycl/test/check_device_code/group_operations.cpp +++ b/sycl/test/check_device_code/group_operations.cpp @@ -54,521 +54,521 @@ 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_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_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) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastji{{m|y}}(i32 2, i32 1, i64 0) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastji{{m|y}}(i32 2, i32 1, i64 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastji{{m|y}}(i32 2, i32 1, i64 2) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiii(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiii(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiii(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiii(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiii(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiii(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiii(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiii(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiii(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxiii(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxiii(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxiii(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMiniii(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMiniii(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMiniii(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiii(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiii(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiii(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiii(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiii(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiii(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastii{{m|y}}(i32 2, i32 1, i64 0) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastii{{m|y}}(i32 2, i32 1, i64 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastii{{m|y}}(i32 2, i32 1, i64 2) // uint8_t -// 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_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) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjj{{m|y}}(i32 2, i32 1, i64 0) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjj{{m|y}}(i32 2, i32 1, i64 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjj{{m|y}}(i32 2, i32 1, i64 2) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiij(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiij(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiij(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiij(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiij(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiij(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiij(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiij(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiij(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxiij(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxiij(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxiij(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMiniij(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMiniij(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMiniij(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiij(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiij(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiij(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiij(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiij(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiij(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastij{{m|y}}(i32 2, i32 1, i64 0) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastij{{m|y}}(i32 2, i32 1, i64 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastij{{m|y}}(i32 2, i32 1, i64 2) // int16_t -// 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_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) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastji{{m|y}}(i32 2, i32 1, i64 0) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastji{{m|y}}(i32 2, i32 1, i64 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastji{{m|y}}(i32 2, i32 1, i64 2) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiii(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiii(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiii(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiii(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiii(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiii(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiii(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiii(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiii(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxiii(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxiii(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxiii(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMiniii(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMiniii(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMiniii(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiii(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiii(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiii(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiii(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiii(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiii(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastii{{m|y}}(i32 2, i32 1, i64 0) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastii{{m|y}}(i32 2, i32 1, i64 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastii{{m|y}}(i32 2, i32 1, i64 2) // uint16_t -// 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_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) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjj{{m|y}}(i32 2, i32 1, i64 0) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjj{{m|y}}(i32 2, i32 1, i64 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjj{{m|y}}(i32 2, i32 1, i64 2) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiij(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiij(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiij(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiij(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiij(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiij(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiij(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiij(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiij(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxiij(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxiij(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxiij(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMiniij(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMiniij(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMiniij(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiij(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiij(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiij(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiij(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiij(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiij(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastij{{m|y}}(i32 2, i32 1, i64 0) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastij{{m|y}}(i32 2, i32 1, i64 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastij{{m|y}}(i32 2, i32 1, i64 2) // int32_t -// 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_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) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastji{{m|y}}(i32 2, i32 1, i64 0) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastji{{m|y}}(i32 2, i32 1, i64 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastji{{m|y}}(i32 2, i32 1, i64 2) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiii(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiii(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiii(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiii(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiii(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiii(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiii(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiii(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiii(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxiii(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxiii(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxiii(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMiniii(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMiniii(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMiniii(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiii(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiii(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiii(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiii(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiii(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiii(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastii{{m|y}}(i32 2, i32 1, i64 0) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastii{{m|y}}(i32 2, i32 1, i64 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastii{{m|y}}(i32 2, i32 1, i64 2) // uint32_t -// 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_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) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjj{{m|y}}(i32 2, i32 1, i64 0) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjj{{m|y}}(i32 2, i32 1, i64 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjj{{m|y}}(i32 2, i32 1, i64 2) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiij(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiij(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiij(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiij(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiij(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiij(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiij(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiij(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiij(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxiij(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxiij(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxiij(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMiniij(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMiniij(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMiniij(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiij(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiij(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiij(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiij(i32 2, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiij(i32 2, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiij(i32 2, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastij{{m|y}}(i32 2, i32 1, i64 0) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastij{{m|y}}(i32 2, i32 1, i64 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastij{{m|y}}(i32 2, i32 1, i64 2) // int64_t (Linux: long, Windows: long long) -// 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_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) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBroadcastj{{l|x}}{{m|y}}(i32 2, i64 1, i64 0) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBroadcastj{{l|x}}{{m|y}}(i32 2, i64 1, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBroadcastj{{l|x}}{{m|y}}(i32 2, i64 1, i64 2) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndKHRii{{l|x}}(i32 2, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndKHRii{{l|x}}(i32 2, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndKHRii{{l|x}}(i32 2, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrKHRii{{l|x}}(i32 2, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrKHRii{{l|x}}(i32 2, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrKHRii{{l|x}}(i32 2, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorKHRii{{l|x}}(i32 2, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorKHRii{{l|x}}(i32 2, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorKHRii{{l|x}}(i32 2, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupSMaxii{{l|x}}(i32 2, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupSMaxii{{l|x}}(i32 2, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupSMaxii{{l|x}}(i32 2, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupSMinii{{l|x}}(i32 2, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupSMinii{{l|x}}(i32 2, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupSMinii{{l|x}}(i32 2, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulKHRii{{l|x}}(i32 2, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulKHRii{{l|x}}(i32 2, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulKHRii{{l|x}}(i32 2, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIAddii{{l|x}}(i32 2, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIAddii{{l|x}}(i32 2, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIAddii{{l|x}}(i32 2, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBroadcasti{{l|x}}{{m|y}}(i32 2, i64 1, i64 0) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBroadcasti{{l|x}}{{m|y}}(i32 2, i64 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBroadcasti{{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_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_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) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBroadcastj{{m|y}}{{m|y}}(i32 2, i64 1, i64 0) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBroadcastj{{m|y}}{{m|y}}(i32 2, i64 1, i64 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBroadcastj{{m|y}}{{m|y}}(i32 2, i64 1, i64 2) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndKHRii{{m|y}}(i32 2, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndKHRii{{m|y}}(i32 2, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndKHRii{{m|y}}(i32 2, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrKHRii{{m|y}}(i32 2, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrKHRii{{m|y}}(i32 2, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrKHRii{{m|y}}(i32 2, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorKHRii{{m|y}}(i32 2, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorKHRii{{m|y}}(i32 2, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorKHRii{{m|y}}(i32 2, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupUMaxii{{m|y}}(i32 2, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupUMaxii{{m|y}}(i32 2, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupUMaxii{{m|y}}(i32 2, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupUMinii{{m|y}}(i32 2, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupUMinii{{m|y}}(i32 2, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupUMinii{{m|y}}(i32 2, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulKHRii{{m|y}}(i32 2, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulKHRii{{m|y}}(i32 2, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulKHRii{{m|y}}(i32 2, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIAddii{{m|y}}(i32 2, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIAddii{{m|y}}(i32 2, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIAddii{{m|y}}(i32 2, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBroadcasti{{m|y}}{{m|y}}(i32 2, i64 1, i64 0) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBroadcasti{{m|y}}{{m|y}}(i32 2, i64 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBroadcasti{{m|y}}{{m|y}}(i32 2, i64 1, i64 2) // half (15360 = 0xH3C00 = 1) -// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMaxjjDF16_(i32 2, i32 0, half 0xH3C00) -// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMaxjjDF16_(i32 2, i32 1, half 0xH3C00) -// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMaxjjDF16_(i32 2, i32 2, half 0xH3C00) -// 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_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) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjj{{m|y}}(i32 2, i32 15360, i64 0) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjj{{m|y}}(i32 2, i32 15360, i64 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjj{{m|y}}(i32 2, i32 15360, i64 2) +// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMaxiiDF16_(i32 2, i32 0, half 0xH3C00) +// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMaxiiDF16_(i32 2, i32 1, half 0xH3C00) +// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMaxiiDF16_(i32 2, i32 2, half 0xH3C00) +// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMiniiDF16_(i32 2, i32 0, half 0xH3C00) +// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMiniiDF16_(i32 2, i32 1, half 0xH3C00) +// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMiniiDF16_(i32 2, i32 2, half 0xH3C00) +// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMulKHRiiDF16_(i32 2, i32 0, half 0xH3C00) +// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMulKHRiiDF16_(i32 2, i32 1, half 0xH3C00) +// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMulKHRiiDF16_(i32 2, i32 2, half 0xH3C00) +// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFAddiiDF16_(i32 2, i32 0, half 0xH3C00) +// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFAddiiDF16_(i32 2, i32 1, half 0xH3C00) +// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFAddiiDF16_(i32 2, i32 2, half 0xH3C00) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastij{{m|y}}(i32 2, i32 15360, i64 0) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastij{{m|y}}(i32 2, i32 15360, i64 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastij{{m|y}}(i32 2, i32 15360, i64 2) // float -// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMaxjjf(i32 2, i32 0, float 1.000000e+00) -// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMaxjjf(i32 2, i32 1, float 1.000000e+00) -// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMaxjjf(i32 2, i32 2, float 1.000000e+00) -// 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_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) -// CHECK: call spir_func float @_Z[[#]]__spirv_GroupBroadcastjf{{m|y}}(i32 2, float 1.000000e+00, i64 0) -// CHECK: call spir_func float @_Z[[#]]__spirv_GroupBroadcastjf{{m|y}}(i32 2, float 1.000000e+00, i64 1) -// CHECK: call spir_func float @_Z[[#]]__spirv_GroupBroadcastjf{{m|y}}(i32 2, float 1.000000e+00, i64 2) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMaxiif(i32 2, i32 0, float 1.000000e+00) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMaxiif(i32 2, i32 1, float 1.000000e+00) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMaxiif(i32 2, i32 2, float 1.000000e+00) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMiniif(i32 2, i32 0, float 1.000000e+00) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMiniif(i32 2, i32 1, float 1.000000e+00) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMiniif(i32 2, i32 2, float 1.000000e+00) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMulKHRiif(i32 2, i32 0, float 1.000000e+00) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMulKHRiif(i32 2, i32 1, float 1.000000e+00) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMulKHRiif(i32 2, i32 2, float 1.000000e+00) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFAddiif(i32 2, i32 0, float 1.000000e+00) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFAddiif(i32 2, i32 1, float 1.000000e+00) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFAddiif(i32 2, i32 2, float 1.000000e+00) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupBroadcastif{{m|y}}(i32 2, float 1.000000e+00, i64 0) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupBroadcastif{{m|y}}(i32 2, float 1.000000e+00, i64 1) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupBroadcastif{{m|y}}(i32 2, float 1.000000e+00, i64 2) // double -// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMaxjjd(i32 2, i32 0, double 1.000000e+00) -// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMaxjjd(i32 2, i32 1, double 1.000000e+00) -// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMaxjjd(i32 2, i32 2, double 1.000000e+00) -// 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_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) -// CHECK: call spir_func double @_Z[[#]]__spirv_GroupBroadcastjd{{m|y}}(i32 2, double 1.000000e+00, i64 0) -// CHECK: call spir_func double @_Z[[#]]__spirv_GroupBroadcastjd{{m|y}}(i32 2, double 1.000000e+00, i64 1) -// CHECK: call spir_func double @_Z[[#]]__spirv_GroupBroadcastjd{{m|y}}(i32 2, double 1.000000e+00, i64 2) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMaxiid(i32 2, i32 0, double 1.000000e+00) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMaxiid(i32 2, i32 1, double 1.000000e+00) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMaxiid(i32 2, i32 2, double 1.000000e+00) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMiniid(i32 2, i32 0, double 1.000000e+00) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMiniid(i32 2, i32 1, double 1.000000e+00) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMiniid(i32 2, i32 2, double 1.000000e+00) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMulKHRiid(i32 2, i32 0, double 1.000000e+00) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMulKHRiid(i32 2, i32 1, double 1.000000e+00) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMulKHRiid(i32 2, i32 2, double 1.000000e+00) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFAddiid(i32 2, i32 0, double 1.000000e+00) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFAddiid(i32 2, i32 1, double 1.000000e+00) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFAddiid(i32 2, i32 2, double 1.000000e+00) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupBroadcastid{{m|y}}(i32 2, double 1.000000e+00, i64 0) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupBroadcastid{{m|y}}(i32 2, double 1.000000e+00, i64 1) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupBroadcastid{{m|y}}(i32 2, double 1.000000e+00, i64 2) SYCL_EXTERNAL void test_sub_group(sub_group g) { test(g); } // int8_t -// 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_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) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjij(i32 3, i32 1, i32 0) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjij(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjij(i32 3, i32 1, i32 2) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiii(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiii(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiii(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiii(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiii(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiii(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiii(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiii(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiii(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxiii(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxiii(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxiii(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMiniii(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMiniii(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMiniii(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiii(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiii(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiii(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiii(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiii(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiii(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastiij(i32 3, i32 1, i32 0) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastiij(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastiij(i32 3, i32 1, i32 2) // uint8_t -// 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_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) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjjj(i32 3, i32 1, i32 0) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjjj(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjjj(i32 3, i32 1, i32 2) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiij(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiij(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiij(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiij(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiij(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiij(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiij(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiij(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiij(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxiij(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxiij(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxiij(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMiniij(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMiniij(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMiniij(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiij(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiij(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiij(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiij(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiij(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiij(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastijj(i32 3, i32 1, i32 0) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastijj(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastijj(i32 3, i32 1, i32 2) // int16_t -// 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_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) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjij(i32 3, i32 1, i32 0) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjij(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjij(i32 3, i32 1, i32 2) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiii(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiii(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiii(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiii(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiii(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiii(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiii(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiii(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiii(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxiii(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxiii(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxiii(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMiniii(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMiniii(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMiniii(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiii(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiii(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiii(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiii(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiii(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiii(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastiij(i32 3, i32 1, i32 0) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastiij(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastiij(i32 3, i32 1, i32 2) // uint16_t -// 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_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) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjjj(i32 3, i32 1, i32 0) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjjj(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjjj(i32 3, i32 1, i32 2) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiij(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiij(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiij(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiij(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiij(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiij(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiij(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiij(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiij(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxiij(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxiij(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxiij(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMiniij(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMiniij(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMiniij(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiij(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiij(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiij(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiij(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiij(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiij(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastijj(i32 3, i32 1, i32 0) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastijj(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastijj(i32 3, i32 1, i32 2) // int32_t -// 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_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) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjij(i32 3, i32 1, i32 0) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjij(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjij(i32 3, i32 1, i32 2) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiii(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiii(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiii(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiii(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiii(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiii(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiii(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiii(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiii(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxiii(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxiii(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMaxiii(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMiniii(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMiniii(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupSMiniii(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiii(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiii(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiii(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiii(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiii(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiii(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastiij(i32 3, i32 1, i32 0) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastiij(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastiij(i32 3, i32 1, i32 2) // uint32_t -// 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_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) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjjj(i32 3, i32 1, i32 0) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjjj(i32 3, i32 1, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjjj(i32 3, i32 1, i32 2) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiij(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiij(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseAndKHRiij(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiij(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiij(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseOrKHRiij(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiij(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiij(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBitwiseXorKHRiij(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxiij(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxiij(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMaxiij(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMiniij(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMiniij(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupUMiniij(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiij(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiij(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIMulKHRiij(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiij(i32 3, i32 0, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiij(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupIAddiij(i32 3, i32 2, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastijj(i32 3, i32 1, i32 0) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastijj(i32 3, i32 1, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastijj(i32 3, i32 1, i32 2) // int64_t (Linux: long, Windows: long long) -// 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_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) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBroadcastj{{l|x}}j(i32 3, i64 1, i32 0) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBroadcastj{{l|x}}j(i32 3, i64 1, i32 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBroadcastj{{l|x}}j(i32 3, i64 1, i32 2) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndKHRii{{l|x}}(i32 3, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndKHRii{{l|x}}(i32 3, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndKHRii{{l|x}}(i32 3, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrKHRii{{l|x}}(i32 3, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrKHRii{{l|x}}(i32 3, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrKHRii{{l|x}}(i32 3, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorKHRii{{l|x}}(i32 3, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorKHRii{{l|x}}(i32 3, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorKHRii{{l|x}}(i32 3, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupSMaxii{{l|x}}(i32 3, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupSMaxii{{l|x}}(i32 3, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupSMaxii{{l|x}}(i32 3, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupSMinii{{l|x}}(i32 3, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupSMinii{{l|x}}(i32 3, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupSMinii{{l|x}}(i32 3, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulKHRii{{l|x}}(i32 3, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulKHRii{{l|x}}(i32 3, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulKHRii{{l|x}}(i32 3, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIAddii{{l|x}}(i32 3, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIAddii{{l|x}}(i32 3, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIAddii{{l|x}}(i32 3, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBroadcasti{{l|x}}j(i32 3, i64 1, i32 0) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBroadcasti{{l|x}}j(i32 3, i64 1, i32 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBroadcasti{{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_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_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) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBroadcastj{{m|y}}j(i32 3, i64 1, i32 0) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBroadcastj{{m|y}}j(i32 3, i64 1, i32 1) -// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBroadcastj{{m|y}}j(i32 3, i64 1, i32 2) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndKHRii{{m|y}}(i32 3, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndKHRii{{m|y}}(i32 3, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseAndKHRii{{m|y}}(i32 3, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrKHRii{{m|y}}(i32 3, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrKHRii{{m|y}}(i32 3, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseOrKHRii{{m|y}}(i32 3, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorKHRii{{m|y}}(i32 3, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorKHRii{{m|y}}(i32 3, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBitwiseXorKHRii{{m|y}}(i32 3, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupUMaxii{{m|y}}(i32 3, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupUMaxii{{m|y}}(i32 3, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupUMaxii{{m|y}}(i32 3, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupUMinii{{m|y}}(i32 3, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupUMinii{{m|y}}(i32 3, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupUMinii{{m|y}}(i32 3, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulKHRii{{m|y}}(i32 3, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulKHRii{{m|y}}(i32 3, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIMulKHRii{{m|y}}(i32 3, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIAddii{{m|y}}(i32 3, i32 0, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIAddii{{m|y}}(i32 3, i32 1, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupIAddii{{m|y}}(i32 3, i32 2, i64 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBroadcasti{{m|y}}j(i32 3, i64 1, i32 0) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBroadcasti{{m|y}}j(i32 3, i64 1, i32 1) +// CHECK: call spir_func i64 @_Z[[#]]__spirv_GroupBroadcasti{{m|y}}j(i32 3, i64 1, i32 2) // half (15360 = 0xH3C00 = 1) -// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMaxjjDF16_(i32 3, i32 0, half 0xH3C00) -// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMaxjjDF16_(i32 3, i32 1, half 0xH3C00) -// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMaxjjDF16_(i32 3, i32 2, half 0xH3C00) -// 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_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) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjjj(i32 3, i32 15360, i32 0) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjjj(i32 3, i32 15360, i32 1) -// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastjjj(i32 3, i32 15360, i32 2) +// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMaxiiDF16_(i32 3, i32 0, half 0xH3C00) +// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMaxiiDF16_(i32 3, i32 1, half 0xH3C00) +// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMaxiiDF16_(i32 3, i32 2, half 0xH3C00) +// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMiniiDF16_(i32 3, i32 0, half 0xH3C00) +// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMiniiDF16_(i32 3, i32 1, half 0xH3C00) +// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMiniiDF16_(i32 3, i32 2, half 0xH3C00) +// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMulKHRiiDF16_(i32 3, i32 0, half 0xH3C00) +// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMulKHRiiDF16_(i32 3, i32 1, half 0xH3C00) +// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFMulKHRiiDF16_(i32 3, i32 2, half 0xH3C00) +// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFAddiiDF16_(i32 3, i32 0, half 0xH3C00) +// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFAddiiDF16_(i32 3, i32 1, half 0xH3C00) +// CHECK: call spir_func half @_Z[[#]]__spirv_GroupFAddiiDF16_(i32 3, i32 2, half 0xH3C00) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastijj(i32 3, i32 15360, i32 0) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastijj(i32 3, i32 15360, i32 1) +// CHECK: call spir_func i32 @_Z[[#]]__spirv_GroupBroadcastijj(i32 3, i32 15360, i32 2) // float -// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMaxjjf(i32 3, i32 0, float 1.000000e+00) -// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMaxjjf(i32 3, i32 1, float 1.000000e+00) -// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMaxjjf(i32 3, i32 2, float 1.000000e+00) -// 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_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) -// CHECK: call spir_func float @_Z[[#]]__spirv_GroupBroadcastjfj(i32 3, float 1.000000e+00, i32 0) -// CHECK: call spir_func float @_Z[[#]]__spirv_GroupBroadcastjfj(i32 3, float 1.000000e+00, i32 1) -// CHECK: call spir_func float @_Z[[#]]__spirv_GroupBroadcastjfj(i32 3, float 1.000000e+00, i32 2) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMaxiif(i32 3, i32 0, float 1.000000e+00) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMaxiif(i32 3, i32 1, float 1.000000e+00) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMaxiif(i32 3, i32 2, float 1.000000e+00) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMiniif(i32 3, i32 0, float 1.000000e+00) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMiniif(i32 3, i32 1, float 1.000000e+00) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMiniif(i32 3, i32 2, float 1.000000e+00) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMulKHRiif(i32 3, i32 0, float 1.000000e+00) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMulKHRiif(i32 3, i32 1, float 1.000000e+00) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFMulKHRiif(i32 3, i32 2, float 1.000000e+00) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFAddiif(i32 3, i32 0, float 1.000000e+00) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFAddiif(i32 3, i32 1, float 1.000000e+00) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupFAddiif(i32 3, i32 2, float 1.000000e+00) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupBroadcastifj(i32 3, float 1.000000e+00, i32 0) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupBroadcastifj(i32 3, float 1.000000e+00, i32 1) +// CHECK: call spir_func float @_Z[[#]]__spirv_GroupBroadcastifj(i32 3, float 1.000000e+00, i32 2) // double -// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMaxjjd(i32 3, i32 0, double 1.000000e+00) -// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMaxjjd(i32 3, i32 1, double 1.000000e+00) -// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMaxjjd(i32 3, i32 2, double 1.000000e+00) -// 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_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) -// CHECK: call spir_func double @_Z[[#]]__spirv_GroupBroadcastjdj(i32 3, double 1.000000e+00, i32 0) -// CHECK: call spir_func double @_Z[[#]]__spirv_GroupBroadcastjdj(i32 3, double 1.000000e+00, i32 1) -// CHECK: call spir_func double @_Z[[#]]__spirv_GroupBroadcastjdj(i32 3, double 1.000000e+00, i32 2) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMaxiid(i32 3, i32 0, double 1.000000e+00) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMaxiid(i32 3, i32 1, double 1.000000e+00) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMaxiid(i32 3, i32 2, double 1.000000e+00) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMiniid(i32 3, i32 0, double 1.000000e+00) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMiniid(i32 3, i32 1, double 1.000000e+00) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMiniid(i32 3, i32 2, double 1.000000e+00) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMulKHRiid(i32 3, i32 0, double 1.000000e+00) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMulKHRiid(i32 3, i32 1, double 1.000000e+00) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFMulKHRiid(i32 3, i32 2, double 1.000000e+00) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFAddiid(i32 3, i32 0, double 1.000000e+00) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFAddiid(i32 3, i32 1, double 1.000000e+00) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupFAddiid(i32 3, i32 2, double 1.000000e+00) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupBroadcastidj(i32 3, double 1.000000e+00, i32 0) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupBroadcastidj(i32 3, double 1.000000e+00, i32 1) +// CHECK: call spir_func double @_Z[[#]]__spirv_GroupBroadcastidj(i32 3, double 1.000000e+00, i32 2) diff --git a/sycl/test/check_device_code/sub_group_mask.cpp b/sycl/test/check_device_code/sub_group_mask.cpp index 2a7c0b14a914f..533cb584d1c97 100644 --- a/sycl/test/check_device_code/sub_group_mask.cpp +++ b/sycl/test/check_device_code/sub_group_mask.cpp @@ -7,4 +7,4 @@ using namespace sycl; SYCL_EXTERNAL void test_group_mask(sub_group g) { ext::oneapi::group_ballot(g, true); } -// CHECK: %{{.*}} = call spir_func <4 x i32> @_Z[[#]]__spirv_GroupNonUniformBallotjb(i32 {{.*}}, i1{{.*}}) +// CHECK: %{{.*}} = call spir_func <4 x i32> @_Z[[#]]__spirv_GroupNonUniformBallotib(i32 {{.*}}, i1{{.*}}) From 4681afb7b1983ff347ab42b4b36853d823f8b864 Mon Sep 17 00:00:00 2001 From: Wenju He Date: Thu, 13 Mar 2025 17:47:04 -0700 Subject: [PATCH 2/6] revert host code ABI change --- sycl/include/sycl/__spirv/spirv_ops.hpp | 2 +- sycl/source/spirv_ops.cpp | 2 +- sycl/test/abi/sycl_symbols_linux.dump | 2 +- sycl/test/abi/sycl_symbols_windows.dump | 2 +- 4 files changed, 4 insertions(+), 4 deletions(-) diff --git a/sycl/include/sycl/__spirv/spirv_ops.hpp b/sycl/include/sycl/__spirv/spirv_ops.hpp index 79271a8d88a93..e93b0f584e78e 100644 --- a/sycl/include/sycl/__spirv/spirv_ops.hpp +++ b/sycl/include/sycl/__spirv/spirv_ops.hpp @@ -1225,6 +1225,6 @@ __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT void __spirv_MemoryBarrier(__spv::Scope Memory, uint32_t Semantics) noexcept; __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT void -__spirv_GroupWaitEvents(int32_t Execution, uint32_t NumEvents, +__spirv_GroupWaitEvents(__spv::Scope Execution, uint32_t NumEvents, __ocl_event_t *WaitEvents) noexcept; #endif // !__SYCL_DEVICE_ONLY__ diff --git a/sycl/source/spirv_ops.cpp b/sycl/source/spirv_ops.cpp index 23bde9f4a2ddd..fdaa7e1834eae 100644 --- a/sycl/source/spirv_ops.cpp +++ b/sycl/source/spirv_ops.cpp @@ -16,7 +16,7 @@ // This operation is NOP on HOST as all operations there are blocking and // by the moment this function was called, the operations generating // the __ocl_event_t objects had already been finished. -__SYCL_EXPORT void __spirv_GroupWaitEvents(int32_t Execution, +__SYCL_EXPORT void __spirv_GroupWaitEvents(__spv::Scope Execution, uint32_t NumEvents, __ocl_event_t *WaitEvents) noexcept { (void)Execution; diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index f2394a2372be1..cbe08c9c12d4f 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -10,7 +10,7 @@ _Z20__spirv_ocl_prefetchPKcm _Z21__spirv_MemoryBarrierN5__spv5ScopeEj _Z22__spirv_ControlBarrierN5__spv5ScopeES0_j -_Z23__spirv_GroupWaitEventsijPPv +_Z23__spirv_GroupWaitEventsN5__spv5ScopeEjPPv _ZN4sycl3_V110__abs_implENS0_3vecIaLi16EEE _ZN4sycl3_V110__abs_implENS0_3vecIaLi1EEE _ZN4sycl3_V110__abs_implENS0_3vecIaLi2EEE diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 55426827a04f8..06907b875aac5 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -3473,7 +3473,7 @@ ?__smoothstep_impl@_V1@sycl@@YAMMMM@Z ?__smoothstep_impl@_V1@sycl@@YANNNN@Z ?__spirv_ControlBarrier@@YAXUScope@__spv@@0I@Z -?__spirv_GroupWaitEvents@@YAXHIPEAPEAX@Z +?__spirv_GroupWaitEvents@@YAXUScope@__spv@@IPEAPEAX@Z ?__spirv_MemoryBarrier@@YAXUScope@__spv@@I@Z ?__spirv_ocl_prefetch@@YAXPEBD_K@Z ?__sqrt_impl@_V1@sycl@@YA?AV?$vec@M$00@12@V312@@Z From 3c8287762e737c4e1a5aadd54332afefe9a28096 Mon Sep 17 00:00:00 2001 From: Wenju He Date: Thu, 13 Mar 2025 17:52:01 -0700 Subject: [PATCH 3/6] fix amdgcn-amdhsa group op type --- libclc/libspirv/lib/amdgcn-amdhsa/group/collectives.cl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/libclc/libspirv/lib/amdgcn-amdhsa/group/collectives.cl b/libclc/libspirv/lib/amdgcn-amdhsa/group/collectives.cl index eb1c1f615d52b..27218592b2487 100644 --- a/libclc/libspirv/lib/amdgcn-amdhsa/group/collectives.cl +++ b/libclc/libspirv/lib/amdgcn-amdhsa/group/collectives.cl @@ -196,7 +196,7 @@ __CLC_SUBGROUP_COLLECTIVE(LogicalAndKHR, __CLC_LOGICAL_AND, bool, a, true) #define __CLC_GROUP_COLLECTIVE_INNER(SPIRV_NAME, CLC_NAME, OP, TYPE, IDENTITY) \ _CLC_DEF _CLC_OVERLOAD _CLC_CONVERGENT TYPE __CLC_APPEND( \ - __spirv_Group, SPIRV_NAME)(int scope, uint op, TYPE x) { \ + __spirv_Group, SPIRV_NAME)(int scope, int op, TYPE x) { \ TYPE carry = IDENTITY; \ /* Perform GroupOperation within sub-group */ \ TYPE sg_x = __CLC_APPEND(__clc__Subgroup, CLC_NAME)(op, x, &carry); \ From bc42ed15d414eec77f232af783919754a0b69a82 Mon Sep 17 00:00:00 2001 From: Wenju He Date: Tue, 18 Mar 2025 17:30:06 -0700 Subject: [PATCH 4/6] add clang test --- .../spirv-builtin-lookup-group.cpp | 114 ++++++++++++++++++ 1 file changed, 114 insertions(+) create mode 100644 clang/test/CodeGenSPIRV/spirv-builtin-lookup-group.cpp diff --git a/clang/test/CodeGenSPIRV/spirv-builtin-lookup-group.cpp b/clang/test/CodeGenSPIRV/spirv-builtin-lookup-group.cpp new file mode 100644 index 0000000000000..dd41db065f57d --- /dev/null +++ b/clang/test/CodeGenSPIRV/spirv-builtin-lookup-group.cpp @@ -0,0 +1,114 @@ +// RUN: %clang_cc1 -cl-std=CL2.0 -fdeclare-spirv-builtins -emit-llvm %s -o - | FileCheck %s + +// Note -cl-std=CL2.0 is added to avoid error +// Assertion `!isNull() && "Cannot retrieve a NULL type pointer"' failed. +// Root cause is unknown. + +void group_async_copy(short __attribute__((opencl_local)) *dst , short const __attribute__((opencl_global)) *src, +event_t event) { + // CHECK-LABEL: @group_async_copy( + // CHECK: call ptr @_Z22__spirv_GroupAsyncCopyiPU7CLlocalsPU8CLglobalKsmm9ocl_event( + __spirv_GroupAsyncCopy(2, dst, src, 32, 16, event); +} + +void group_wait_event(event_t event) { + // CHECK-LABEL: @group_wait_event( + // CHECK: call void @_Z23__spirv_GroupWaitEventsiiPU9CLprivate9ocl_event( + __spirv_GroupWaitEvents(1, 0, &event); +} + +bool group_all(bool predicate) { + // CHECK-LABEL: @group_all( + // CHECK: call zeroext i1 @_Z16__spirv_GroupAllib( + return __spirv_GroupAll(2, predicate); +} + +bool group_any(bool predicate) { + // CHECK-LABEL: @group_any( + // CHECK: call zeroext i1 @_Z16__spirv_GroupAnyib( + return __spirv_GroupAny(2, predicate); +} + +char group_broad_cast(char a) { + // CHECK-LABEL: @group_broad_cast( + // CHECK: call i32 @_Z22__spirv_GroupBroadcastiij( + return __spirv_GroupBroadcast(2, a, 0u); +} + +int group_iadd(int a) { + // CHECK-LABEL: @group_iadd( + // CHECK: call i32 @_Z17__spirv_GroupIAddiii( + return __spirv_GroupIAdd(2, 2, a); +} + +int group_imul_khr(short a) { + // CHECK-LABEL: @group_imul_khr( + // CHECK: call signext i16 @_Z20__spirv_GroupIMulKHRiis( + return __spirv_GroupIMulKHR(2, 0, a); +} + +bool group_bitwise_or_khr(long a) { + // CHECK-LABEL: @group_bitwise_or_khr( + // CHECK: call i64 @_Z25__spirv_GroupBitwiseOrKHRiil( + return __spirv_GroupBitwiseOrKHR(2, 0, a); +} + +float group_fadd(float a) { + // CHECK-LABEL: @group_fadd( + // CHECK: call float @_Z17__spirv_GroupFAddiif( + return __spirv_GroupFAdd(2, 1, a); +} + +float group_fmin(float a) { + // CHECK-LABEL: @group_fmin( + // CHECK: call float @_Z17__spirv_GroupFMiniif( + return __spirv_GroupFMin(2, 0, a); +} + +float group_fmax(float a) { + // CHECK-LABEL: @group_fmax( + // CHECK: call float @_Z17__spirv_GroupFMaxiif( + return __spirv_GroupFMax(2, 2, a); +} + +float group_ful_khr(float a) { + // CHECK-LABEL: @group_ful_khr( + // CHECK: call float @_Z20__spirv_GroupFMulKHRiif( + return __spirv_GroupFMulKHR(2, 2, a); +} + +unsigned char group_umin(unsigned char a ) { + // CHECK-LABEL: @group_umin( + // CHECK: call zeroext i8 @_Z17__spirv_GroupUMiniih( + return __spirv_GroupUMin(2, 0, a); +} + +unsigned long group_umax(unsigned long a) { + // CHECK-LABEL: @group_umax( + // CHECK: call i64 @_Z17__spirv_GroupUMaxiim( + return __spirv_GroupUMax(2, 0, a); +} + +char group_smin(char a) { + // CHECK-LABEL: @group_smin( + // CHECK: call i32 @_Z17__spirv_GroupSMiniii( + return __spirv_GroupSMin(2, 0, a); +} + +short group_smax(short a) { + // CHECK-LABEL: @group_smax( + // CHECK: call signext i16 @_Z17__spirv_GroupSMaxiis( + return __spirv_GroupSMax(2, 0, a); +} + +bool group_logical_and_khr(bool a) { + // CHECK-LABEL: @group_logical_and_khr( + // CHECK: call zeroext i1 @_Z26__spirv_GroupLogicalAndKHRiib( + return __spirv_GroupLogicalAndKHR(2, 0, a); +} + +bool group_logical_or_khr(bool a) { + // CHECK-LABEL: @group_logical_or_khr( + // CHECK: call zeroext i1 @_Z25__spirv_GroupLogicalOrKHRiib( + return __spirv_GroupLogicalOrKHR(2, 0, a); +} From 7b78147b7513f3d0a5e40cc6d96d7cc7d27a6648 Mon Sep 17 00:00:00 2001 From: Wenju He Date: Wed, 19 Mar 2025 03:36:34 -0700 Subject: [PATCH 5/6] change clang test file to .cl --- ...roup.cpp => spirv-builtin-lookup-group.cl} | 42 +++++++++---------- 1 file changed, 19 insertions(+), 23 deletions(-) rename clang/test/CodeGenSPIRV/{spirv-builtin-lookup-group.cpp => spirv-builtin-lookup-group.cl} (61%) diff --git a/clang/test/CodeGenSPIRV/spirv-builtin-lookup-group.cpp b/clang/test/CodeGenSPIRV/spirv-builtin-lookup-group.cl similarity index 61% rename from clang/test/CodeGenSPIRV/spirv-builtin-lookup-group.cpp rename to clang/test/CodeGenSPIRV/spirv-builtin-lookup-group.cl index dd41db065f57d..0a31736f04122 100644 --- a/clang/test/CodeGenSPIRV/spirv-builtin-lookup-group.cpp +++ b/clang/test/CodeGenSPIRV/spirv-builtin-lookup-group.cl @@ -1,114 +1,110 @@ -// RUN: %clang_cc1 -cl-std=CL2.0 -fdeclare-spirv-builtins -emit-llvm %s -o - | FileCheck %s - -// Note -cl-std=CL2.0 is added to avoid error -// Assertion `!isNull() && "Cannot retrieve a NULL type pointer"' failed. -// Root cause is unknown. +// RUN: %clang_cc1 -triple=spir64 -cl-std=CL2.0 -fdeclare-spirv-builtins -emit-llvm %s -o - | FileCheck %s void group_async_copy(short __attribute__((opencl_local)) *dst , short const __attribute__((opencl_global)) *src, event_t event) { // CHECK-LABEL: @group_async_copy( - // CHECK: call ptr @_Z22__spirv_GroupAsyncCopyiPU7CLlocalsPU8CLglobalKsmm9ocl_event( + // CHECK: tail call spir_func target("spirv.Event") @_Z22__spirv_GroupAsyncCopyiPU3AS3sPU3AS1Ksmm9ocl_event( __spirv_GroupAsyncCopy(2, dst, src, 32, 16, event); } void group_wait_event(event_t event) { // CHECK-LABEL: @group_wait_event( - // CHECK: call void @_Z23__spirv_GroupWaitEventsiiPU9CLprivate9ocl_event( + // CHECK: call spir_func void @_Z23__spirv_GroupWaitEventsiiP9ocl_event( __spirv_GroupWaitEvents(1, 0, &event); } bool group_all(bool predicate) { // CHECK-LABEL: @group_all( - // CHECK: call zeroext i1 @_Z16__spirv_GroupAllib( + // CHECK: call spir_func zeroext i1 @_Z16__spirv_GroupAllib( return __spirv_GroupAll(2, predicate); } bool group_any(bool predicate) { // CHECK-LABEL: @group_any( - // CHECK: call zeroext i1 @_Z16__spirv_GroupAnyib( + // CHECK: call spir_func zeroext i1 @_Z16__spirv_GroupAnyib( return __spirv_GroupAny(2, predicate); } char group_broad_cast(char a) { // CHECK-LABEL: @group_broad_cast( - // CHECK: call i32 @_Z22__spirv_GroupBroadcastiij( + // CHECK: call spir_func i32 @_Z22__spirv_GroupBroadcastiij( return __spirv_GroupBroadcast(2, a, 0u); } int group_iadd(int a) { // CHECK-LABEL: @group_iadd( - // CHECK: call i32 @_Z17__spirv_GroupIAddiii( + // CHECK: call spir_func i32 @_Z17__spirv_GroupIAddiii( return __spirv_GroupIAdd(2, 2, a); } int group_imul_khr(short a) { // CHECK-LABEL: @group_imul_khr( - // CHECK: call signext i16 @_Z20__spirv_GroupIMulKHRiis( + // CHECK: call spir_func signext i16 @_Z20__spirv_GroupIMulKHRiis( return __spirv_GroupIMulKHR(2, 0, a); } bool group_bitwise_or_khr(long a) { // CHECK-LABEL: @group_bitwise_or_khr( - // CHECK: call i64 @_Z25__spirv_GroupBitwiseOrKHRiil( + // CHECK: call spir_func i64 @_Z25__spirv_GroupBitwiseOrKHRiil( return __spirv_GroupBitwiseOrKHR(2, 0, a); } float group_fadd(float a) { // CHECK-LABEL: @group_fadd( - // CHECK: call float @_Z17__spirv_GroupFAddiif( + // CHECK: call spir_func float @_Z17__spirv_GroupFAddiif( return __spirv_GroupFAdd(2, 1, a); } float group_fmin(float a) { // CHECK-LABEL: @group_fmin( - // CHECK: call float @_Z17__spirv_GroupFMiniif( + // CHECK: call spir_func float @_Z17__spirv_GroupFMiniif( return __spirv_GroupFMin(2, 0, a); } float group_fmax(float a) { // CHECK-LABEL: @group_fmax( - // CHECK: call float @_Z17__spirv_GroupFMaxiif( + // CHECK: call spir_func float @_Z17__spirv_GroupFMaxiif( return __spirv_GroupFMax(2, 2, a); } float group_ful_khr(float a) { // CHECK-LABEL: @group_ful_khr( - // CHECK: call float @_Z20__spirv_GroupFMulKHRiif( + // CHECK: call spir_func float @_Z20__spirv_GroupFMulKHRiif( return __spirv_GroupFMulKHR(2, 2, a); } unsigned char group_umin(unsigned char a ) { // CHECK-LABEL: @group_umin( - // CHECK: call zeroext i8 @_Z17__spirv_GroupUMiniih( + // CHECK: call spir_func zeroext i8 @_Z17__spirv_GroupUMiniih( return __spirv_GroupUMin(2, 0, a); } unsigned long group_umax(unsigned long a) { // CHECK-LABEL: @group_umax( - // CHECK: call i64 @_Z17__spirv_GroupUMaxiim( + // CHECK: call spir_func i64 @_Z17__spirv_GroupUMaxiim( return __spirv_GroupUMax(2, 0, a); } char group_smin(char a) { // CHECK-LABEL: @group_smin( - // CHECK: call i32 @_Z17__spirv_GroupSMiniii( + // CHECK: call spir_func i32 @_Z17__spirv_GroupSMiniii( return __spirv_GroupSMin(2, 0, a); } short group_smax(short a) { // CHECK-LABEL: @group_smax( - // CHECK: call signext i16 @_Z17__spirv_GroupSMaxiis( + // CHECK: call spir_func signext i16 @_Z17__spirv_GroupSMaxiis( return __spirv_GroupSMax(2, 0, a); } bool group_logical_and_khr(bool a) { // CHECK-LABEL: @group_logical_and_khr( - // CHECK: call zeroext i1 @_Z26__spirv_GroupLogicalAndKHRiib( + // CHECK: call spir_func zeroext i1 @_Z26__spirv_GroupLogicalAndKHRiib( return __spirv_GroupLogicalAndKHR(2, 0, a); } bool group_logical_or_khr(bool a) { // CHECK-LABEL: @group_logical_or_khr( - // CHECK: call zeroext i1 @_Z25__spirv_GroupLogicalOrKHRiib( + // CHECK: call spir_func zeroext i1 @_Z25__spirv_GroupLogicalOrKHRiib( return __spirv_GroupLogicalOrKHR(2, 0, a); } From a4b01fbf3311a82b9a2bbc04d027657092c5502d Mon Sep 17 00:00:00 2001 From: Wenju He Date: Fri, 21 Mar 2025 01:36:09 -0700 Subject: [PATCH 6/6] fix wrong merge conflict resolve --- libdevice/nativecpu_utils.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/libdevice/nativecpu_utils.cpp b/libdevice/nativecpu_utils.cpp index 03f9dcd953a0b..1b84903018c51 100644 --- a/libdevice/nativecpu_utils.cpp +++ b/libdevice/nativecpu_utils.cpp @@ -124,7 +124,7 @@ DefineGOp1(All, all) MuxType) noexcept; \ DEVICE_EXTERN_C MuxType __mux_work_group_scan_inclusive_##mux_sfx(uint32_t, \ MuxType) noexcept; \ - DEVICE_EXTERN_C MuxType __mux_work_group_reduce_##mux_sfx(uint32_t, MuxType); noexcept\ + DEVICE_EXTERN_C MuxType __mux_work_group_reduce_##mux_sfx(uint32_t, MuxType) noexcept;\ DEVICE_EXTERNAL Type __spirv_Group##spir_sfx(int32_t g, int32_t id, \ Type v) noexcept { \ if (__spv::Scope::Flag::Subgroup == g) { \ @@ -212,7 +212,7 @@ DefineLogicalGroupOp(bool, bool, i1) } \ \ DEVICE_EXTERNAL Type __spirv_GroupBroadcast(int32_t g, Type v, \ - sycl::vec::vector_t l) noexcept { \ + sycl::vec::vector_t l) noexcept { \ if (__spv::Scope::Flag::Subgroup == g) \ return __mux_sub_group_broadcast_##Sfx(v, l[0]); \ else \ @@ -220,7 +220,7 @@ DefineLogicalGroupOp(bool, bool, i1) } \ \ DEVICE_EXTERNAL Type __spirv_GroupBroadcast(int32_t g, Type v, \ - sycl::vec::vector_t l) noexcept { \ + sycl::vec::vector_t l) noexcept { \ if (__spv::Scope::Flag::Subgroup == g) \ return __mux_sub_group_broadcast_##Sfx(v, l[0]); \ else \