diff --git a/clang/lib/Sema/SPIRVBuiltins.td b/clang/lib/Sema/SPIRVBuiltins.td
index 4113e5e23aca1..2e46b37cf6974 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<name, [Event, UInt, PointerType<AGenTypeN, LocalAS>, PointerType<ConstType<AGenTypeN>, GlobalAS>, Size, Size, Event], Attr.Convergent>;
- def : SPVBuiltin<name, [Event, UInt, PointerType<AGenTypeN, GlobalAS>, PointerType<ConstType<AGenTypeN>, 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<name,
+ [Event, Int, PointerType<AGenTypeN, LocalAS>,
+ PointerType<ConstType<AGenTypeN>, GlobalAS>, Size, Size,
+ Event],
+ Attr.Convergent>;
+ def : SPVBuiltin<name,
+ [Event, Int, PointerType<AGenTypeN, GlobalAS>,
+ PointerType<ConstType<AGenTypeN>, LocalAS>, Size, Size,
+ Event],
+ Attr.Convergent>;
}
foreach name = ["GroupWaitEvents"] in {
- def : SPVBuiltin<name, [Void, UInt, Int, PointerType<Event, DefaultAS>], Attr.Convergent>;
- def : SPVBuiltin<name, [Void, UInt, Int, PointerType<Event, PrivateAS>], Attr.Convergent>;
- def : SPVBuiltin<name, [Void, UInt, Int, PointerType<Event, GenericAS>], Attr.Convergent>;
+ def : SPVBuiltin<name, [Void, Int, Int, PointerType<Event, DefaultAS>],
+ Attr.Convergent>;
+ def : SPVBuiltin<name, [Void, Int, Int, PointerType<Event, PrivateAS>],
+ Attr.Convergent>;
+ def : SPVBuiltin<name, [Void, Int, Int, PointerType<Event, GenericAS>],
+ Attr.Convergent>;
}
foreach name = ["GroupAll", "GroupAny"] in {
- def : SPVBuiltin<name, [Bool, UInt, Bool], Attr.Convergent>;
+ def : SPVBuiltin<name, [Bool, Int, Bool], Attr.Convergent>;
}
foreach name = ["GroupBroadcast"] in {
foreach IDType = TLAllInts.List in {
- def : SPVBuiltin<name, [AGenTypeN, UInt, AGenTypeN, IDType], Attr.Convergent>;
- def : SPVBuiltin<name, [AGenTypeN, UInt, AGenTypeN, VectorType<IDType, 2>], Attr.Convergent>;
- def : SPVBuiltin<name, [AGenTypeN, UInt, AGenTypeN, VectorType<IDType, 3>], Attr.Convergent>;
- def : SPVBuiltin<name, [Bool, UInt, Bool, IDType], Attr.Convergent>;
- def : SPVBuiltin<name, [Bool, UInt, Bool, VectorType<IDType, 2>], Attr.Convergent>;
- def : SPVBuiltin<name, [Bool, UInt, Bool, VectorType<IDType, 3>], Attr.Convergent>;
+ def : SPVBuiltin<name, [AGenTypeN, Int, AGenTypeN, IDType],
+ Attr.Convergent>;
+ def : SPVBuiltin<name, [AGenTypeN, Int, AGenTypeN, VectorType<IDType, 2>],
+ Attr.Convergent>;
+ def : SPVBuiltin<name, [AGenTypeN, Int, AGenTypeN, VectorType<IDType, 3>],
+ Attr.Convergent>;
+ def : SPVBuiltin<name, [Bool, Int, Bool, IDType], Attr.Convergent>;
+ def : SPVBuiltin<name, [Bool, Int, Bool, VectorType<IDType, 2>],
+ Attr.Convergent>;
+ def : SPVBuiltin<name, [Bool, Int, Bool, VectorType<IDType, 3>],
+ Attr.Convergent>;
}
}
foreach name = ["GroupIAdd", "GroupIMulKHR", "GroupBitwiseOrKHR",
"GroupBitwiseXorKHR", "GroupBitwiseAndKHR"] in {
- def : SPVBuiltin<name, [AIGenTypeN, UInt, UInt, AIGenTypeN], Attr.Convergent>;
+ def : SPVBuiltin<name, [AIGenTypeN, Int, Int, AIGenTypeN], Attr.Convergent>;
}
foreach name = ["GroupFAdd", "GroupFMin", "GroupFMax",
"GroupFMulKHR"] in {
- def : SPVBuiltin<name, [FGenTypeN, UInt, UInt, FGenTypeN], Attr.Convergent>;
+ def : SPVBuiltin<name, [FGenTypeN, Int, Int, FGenTypeN], Attr.Convergent>;
}
foreach name = ["GroupUMin", "GroupUMax"] in {
- def : SPVBuiltin<name, [AUIGenTypeN, UInt, UInt, AUIGenTypeN], Attr.Convergent>;
+ def : SPVBuiltin<name, [AUIGenTypeN, Int, Int, AUIGenTypeN], Attr.Convergent>;
}
foreach name = ["GroupSMin", "GroupSMax"] in {
- def : SPVBuiltin<name, [ASIGenTypeN, UInt, UInt, ASIGenTypeN], Attr.Convergent>;
+ def : SPVBuiltin<name, [ASIGenTypeN, Int, Int, ASIGenTypeN], Attr.Convergent>;
}
// TODO: These builtins need to support vectors of bool.
foreach name = ["GroupLogicalAndKHR", "GroupLogicalOrKHR"] in {
- def : SPVBuiltin<name, [Bool, UInt, UInt, Bool], Attr.Convergent>;
+ def : SPVBuiltin<name, [Bool, Int, Int, Bool], Attr.Convergent>;
}
diff --git a/clang/test/CodeGenSPIRV/spirv-builtin-lookup-group.cl b/clang/test/CodeGenSPIRV/spirv-builtin-lookup-group.cl
new file mode 100644
index 0000000000000..0a31736f04122
--- /dev/null
+++ b/clang/test/CodeGenSPIRV/spirv-builtin-lookup-group.cl
@@ -0,0 +1,110 @@
+// 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: 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 spir_func void @_Z23__spirv_GroupWaitEventsiiP9ocl_event(
+ __spirv_GroupWaitEvents(1, 0, &event);
+}
+
+bool group_all(bool predicate) {
+ // CHECK-LABEL: @group_all(
+ // 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 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 spir_func i32 @_Z22__spirv_GroupBroadcastiij(
+ return __spirv_GroupBroadcast(2, a, 0u);
+}
+
+int group_iadd(int a) {
+ // CHECK-LABEL: @group_iadd(
+ // 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 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 spir_func i64 @_Z25__spirv_GroupBitwiseOrKHRiil(
+ return __spirv_GroupBitwiseOrKHR(2, 0, a);
+}
+
+float group_fadd(float a) {
+ // CHECK-LABEL: @group_fadd(
+ // 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 spir_func float @_Z17__spirv_GroupFMiniif(
+ return __spirv_GroupFMin(2, 0, a);
+}
+
+float group_fmax(float a) {
+ // CHECK-LABEL: @group_fmax(
+ // 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 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 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 spir_func i64 @_Z17__spirv_GroupUMaxiim(
+ return __spirv_GroupUMax(2, 0, a);
+}
+
+char group_smin(char a) {
+ // CHECK-LABEL: @group_smin(
+ // 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 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 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 spir_func zeroext i1 @_Z25__spirv_GroupLogicalOrKHRiib(
+ return __spirv_GroupLogicalOrKHR(2, 0, a);
+}
diff --git a/libclc/libspirv/include/libspirv/spirv_builtins.h b/libclc/libspirv/include/libspirv/spirv_builtins.h
index 18483265d1f6a..ac3e871b3b0be 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 <libspirv/spirv.h>
-_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..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)(uint 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); \
@@ -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 <libspirv/spirv.h>
-_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 <libspirv/spirv_types.h>
_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 6c6ae6d021d02..1b84903018c51 100644
--- a/libdevice/nativecpu_utils.cpp
+++ b/libdevice/nativecpu_utils.cpp
@@ -104,7 +104,7 @@ DefSubgroupBlockINTEL(uint8_t) DefSubgroupBlockINTEL(uint16_t)
#define DefineGOp1(spir_sfx, name)\
DEVICE_EXTERN_C bool __mux_sub_group_##name##_i1(bool) noexcept;\
DEVICE_EXTERN_C bool __mux_work_group_##name##_i1(uint32_t id, bool val) noexcept;\
-DEVICE_EXTERNAL bool __spirv_Group ## spir_sfx(unsigned g, bool val) noexcept {\
+DEVICE_EXTERNAL bool __spirv_Group ## spir_sfx(int32_t g, bool val) noexcept {\
if (__spv::Scope::Flag::Subgroup == g)\
return __mux_sub_group_##name##_i1(val);\
else if (__spv::Scope::Flag::Workgroup == g)\
@@ -119,13 +119,13 @@ DefineGOp1(All, all)
#define DefineGOp(Type, MuxType, spir_sfx, mux_sfx) \
DEVICE_EXTERN_C MuxType __mux_sub_group_scan_inclusive_##mux_sfx(MuxType) noexcept; \
DEVICE_EXTERN_C MuxType __mux_sub_group_scan_exclusive_##mux_sfx(MuxType) noexcept; \
- DEVICE_EXTERN_C MuxType __mux_sub_group_reduce_##mux_sfx(MuxType) noexcept; \
+ DEVICE_EXTERN_C MuxType __mux_sub_group_reduce_##mux_sfx(MuxType) noexcept; \
DEVICE_EXTERN_C MuxType __mux_work_group_scan_exclusive_##mux_sfx(uint32_t, \
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_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) noexcept { \
if (__spv::Scope::Flag::Subgroup == g) { \
if (static_cast<unsigned>(__spv::GroupOperation::InclusiveScan) == id) \
@@ -203,7 +203,7 @@ DefineLogicalGroupOp(bool, bool, i1)
int32_t sg_lid) noexcept;
#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) noexcept { \
if (__spv::Scope::Flag::Subgroup == g) \
return __mux_sub_group_broadcast_##Sfx(v, l); \
@@ -211,16 +211,16 @@ 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, \
- sycl::vec<IDType, 2>::vector_t l) noexcept{ \
+ DEVICE_EXTERNAL Type __spirv_GroupBroadcast(int32_t g, Type v, \
+ sycl::vec<IDType, 2>::vector_t l) noexcept { \
if (__spv::Scope::Flag::Subgroup == g) \
return __mux_sub_group_broadcast_##Sfx(v, l[0]); \
else \
return __mux_work_group_broadcast_##Sfx(0, v, l[0], l[1], 0); \
} \
\
- DEVICE_EXTERNAL Type __spirv_GroupBroadcast(uint32_t g, Type v, \
- sycl::vec<IDType, 3>::vector_t l) noexcept{ \
+ DEVICE_EXTERNAL Type __spirv_GroupBroadcast(int32_t g, Type v, \
+ sycl::vec<IDType, 3>::vector_t l) noexcept { \
if (__spv::Scope::Flag::Subgroup == g) \
return __mux_sub_group_broadcast_##Sfx(v, l[0]); \
else \
diff --git a/sycl/include/sycl/__spirv/spirv_ops.hpp b/sycl/include/sycl/__spirv/spirv_ops.hpp
index f83067c8b2f4f..04cfec68e7679 100644
--- a/sycl/include/sycl/__spirv/spirv_ops.hpp
+++ b/sycl/include/sycl/__spirv/spirv_ops.hpp
@@ -883,7 +883,7 @@ extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<uint16_t, N>
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL
__SYCL_EXPORT __ocl_vec_t<uint32_t, 4>
- __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
@@ -911,137 +911,125 @@ __spirv_GroupNonUniformAny(__spv::Scope::Flag, bool);
template <typename ValueT>
__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 <typename ValueT>
__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 <typename ValueT>
__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 <typename ValueT>
__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 <typename ValueT>
__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 <typename ValueT>
__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 <typename ValueT>
__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 <typename ValueT>
__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 <typename ValueT>
__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 <typename ValueT>
__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 <typename ValueT>
__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 <typename ValueT>
__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 <typename ValueT>
__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 <typename ValueT>
__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 <typename ValueT>
__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 <typename ValueT>
__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 <typename ValueT>
__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 <typename ValueT>
__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 <typename ValueT>
__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 <typename ValueT>
__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 <typename ValueT>
__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 <typename ValueT>
__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 <typename ValueT>
__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 <typename ValueT>
__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 <typename ValueT>
__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 <typename ValueT>
__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 <typename ValueT>
__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 <typename ValueT>
__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 <typename ValueT>
__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 <typename ValueT>
__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
@@ -1208,9 +1196,9 @@ extern __DPCPP_SYCL_EXTERNAL void __spirv_TaskSequenceReleaseINTEL(
template <typename dataT>
__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];
}
@@ -1220,9 +1208,9 @@ __SYCL_OpGroupAsyncCopyGlobalToLocal(__spv::Scope::Flag, dataT *Dest,
template <typename dataT>
__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];
}
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/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 <typename G> [[gnu::always_inline]] void test(G g) {
SYCL_EXTERNAL void test_group(group<> g) { test(g); }
// int8_t
-// CHECK: call spir_func i32 @_Z[[#]]__spirv_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{{.*}})