Skip to content

Commit 7f2971a

Browse files
wenju-hewenju-he
authored
[SYCL] Change SPIR-V Enum token type from unsigned int to int for barrier builtins (#17438)
Motivation is unifying SPIR-V builtin mangling to enhance SYCL AOT support for backend targets that bypass SPIR-V generation. Changing to signed int type aligns with * SPV-IR output of llvm-spirv translator. * Default underlying type of enum being int, e.g. enum defined in tablegened LLVM SPIR-V backend header and standard SPIR-V header. Changes are mainly made to following places: * clang/lib/Sema/SPIRVBuiltins.td * libclc/libspirv
1 parent a6a1a39 commit 7f2971a

35 files changed

+282
-264
lines changed

‎clang/lib/Sema/SPIRVBuiltins.td

+6-6
Original file line numberDiff line numberDiff line change
@@ -897,15 +897,15 @@ foreach name = ["BitCount"] in {
897897
// 3.32.20. Barrier Instructions
898898

899899
foreach name = ["ControlBarrier"] in {
900-
// TODO: Allow enum flags instead of UInt ?
901-
// TODO: We should enforce that the UInt must be a literal.
902-
def : SPVBuiltin<name, [Void, UInt, UInt, UInt], Attr.Convergent>;
900+
// TODO: Allow enum flags instead of Int ?
901+
// TODO: We should enforce that the Int must be a literal.
902+
def : SPVBuiltin<name, [Void, Int, Int, Int], Attr.Convergent>;
903903
}
904904

905905
foreach name = ["MemoryBarrier"] in {
906-
// TODO: Allow enum flags instead of UInt ?
907-
// TODO: We should enforce that the UInt must be a literal.
908-
def : SPVBuiltin<name, [Void, UInt, UInt]>;
906+
// TODO: Allow enum flags instead of Int ?
907+
// TODO: We should enforce that the Int must be a literal.
908+
def : SPVBuiltin<name, [Void, Int, Int]>;
909909
}
910910

911911
// 3.32.21. Group and Subgroup Instructions

‎clang/test/CodeGenSPIRV/spirv-builtin-lookup.cpp

+16
Original file line numberDiff line numberDiff line change
@@ -15,3 +15,19 @@ double acos(double val) {
1515
}
1616

1717
// CHECK: declare noundef double @_Z16__spirv_ocl_acosd(double noundef)
18+
19+
void control_barrier() {
20+
// CHECK-LABEL: @_Z15control_barrierv
21+
// CHECK: call void @_Z22__spirv_ControlBarrieriii
22+
__spirv_ControlBarrier(2, 2, 912);
23+
}
24+
25+
// CHECK: declare void @_Z22__spirv_ControlBarrieriii(i32 noundef, i32 noundef, i32 noundef)
26+
27+
void memory_barrier() {
28+
// CHECK-LABEL: @_Z14memory_barrierv
29+
// CHECK: call void @_Z21__spirv_MemoryBarrierii(
30+
__spirv_MemoryBarrier(0, 2);
31+
}
32+
33+
// CHECK: declare void @_Z21__spirv_MemoryBarrierii(i32 noundef, i32 noundef)

‎clang/test/CodeGenSYCL/Inputs/sycl.hpp

+1-1
Original file line numberDiff line numberDiff line change
@@ -19,7 +19,7 @@ extern "C" int printf(const char* fmt, ...);
1919

2020
#ifdef __SYCL_DEVICE_ONLY__
2121
__attribute__((convergent)) extern __attribute__((sycl_device)) void
22-
__spirv_ControlBarrier(int, int, int) noexcept;
22+
__spirv_ControlBarrier(int, int, int);
2323
#endif
2424

2525
// Dummy runtime classes to model SYCL API.

‎libclc/libspirv/include/libspirv/spirv_builtins.h

+3-3
Original file line numberDiff line numberDiff line change
@@ -176,7 +176,7 @@ _CLC_OVERLOAD _CLC_DECL _CLC_CONSTFN __clc_vec16_uint64_t
176176
__spirv_BitCount(__clc_vec16_uint64_t);
177177

178178
_CLC_OVERLOAD _CLC_DECL _CLC_CONVERGENT void
179-
__spirv_ControlBarrier(__clc_uint32_t, __clc_uint32_t, __clc_uint32_t);
179+
__spirv_ControlBarrier(__clc_int32_t, __clc_int32_t, __clc_int32_t);
180180

181181
_CLC_OVERLOAD
182182
_CLC_DECL _CLC_CONSTFN __clc_int8_t __spirv_ConvertFToS_Rchar(__clc_fp32_t);
@@ -10622,8 +10622,8 @@ _CLC_OVERLOAD _CLC_DECL _CLC_CONSTFN __clc_vec16_int8_t
1062210622
__spirv_LessOrGreater(__clc_vec16_fp16_t, __clc_vec16_fp16_t);
1062310623
#endif
1062410624

10625-
_CLC_OVERLOAD _CLC_DECL void __spirv_MemoryBarrier(__clc_uint32_t,
10626-
__clc_uint32_t);
10625+
_CLC_OVERLOAD _CLC_DECL void __spirv_MemoryBarrier(__clc_int32_t,
10626+
__clc_int32_t);
1062710627

1062810628
_CLC_OVERLOAD _CLC_DECL _CLC_CONSTFN __clc_bool_t __spirv_Ordered(__clc_fp32_t,
1062910629
__clc_fp32_t);

‎libclc/libspirv/lib/amdgcn-amdhsa/synchronization/barrier.cl

+8-8
Original file line numberDiff line numberDiff line change
@@ -31,8 +31,8 @@
3131
} \
3232
}
3333

34-
_CLC_INLINE void builtin_fence_order(unsigned int scope_memory,
35-
unsigned int order) {
34+
_CLC_INLINE void builtin_fence_order(int scope_memory,
35+
int order) {
3636
switch ((enum Scope)scope_memory) {
3737
case CrossDevice:
3838
BUILTIN_FENCE(order, "")
@@ -48,19 +48,19 @@ _CLC_INLINE void builtin_fence_order(unsigned int scope_memory,
4848
}
4949
#undef BUILTIN_FENCE
5050

51-
_CLC_DEF _CLC_OVERLOAD void __mem_fence(unsigned int scope_memory,
52-
unsigned int semantics) {
51+
_CLC_DEF _CLC_OVERLOAD void __mem_fence(int scope_memory,
52+
int semantics) {
5353
builtin_fence_order(scope_memory, semantics & 0x1F);
5454
}
5555

56-
_CLC_OVERLOAD _CLC_DEF void __spirv_MemoryBarrier(unsigned int scope_memory,
57-
unsigned int semantics) {
56+
_CLC_OVERLOAD _CLC_DEF void __spirv_MemoryBarrier(int scope_memory,
57+
int semantics) {
5858
__mem_fence(scope_memory, semantics);
5959
}
6060

6161
_CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT void
62-
__spirv_ControlBarrier(unsigned int scope_execution, unsigned int scope_memory,
63-
unsigned int semantics) {
62+
__spirv_ControlBarrier(int scope_execution, int scope_memory,
63+
int semantics) {
6464
if (semantics) {
6565
__mem_fence(scope_memory, semantics);
6666
}

‎libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_cmpxchg.cl

+1-1
Original file line numberDiff line numberDiff line change
@@ -11,7 +11,7 @@
1111
#include <libspirv/spirv_types.h>
1212

1313
int __clc_nvvm_reflect_arch();
14-
_CLC_OVERLOAD _CLC_DECL void __spirv_MemoryBarrier(unsigned int, unsigned int);
14+
_CLC_OVERLOAD _CLC_DECL void __spirv_MemoryBarrier(int, int);
1515

1616
#define __CLC_NVVM_ATOMIC_CAS_IMPL_ORDER(TYPE, TYPE_NV, TYPE_MANGLED_NV, OP, \
1717
ADDR_SPACE, ADDR_SPACE_NV, ORDER) \

‎libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_helpers.h

+1-1
Original file line numberDiff line numberDiff line change
@@ -13,7 +13,7 @@
1313
#include <libspirv/spirv_types.h>
1414

1515
extern int __clc_nvvm_reflect_arch();
16-
_CLC_OVERLOAD _CLC_DECL void __spirv_MemoryBarrier(unsigned int, unsigned int);
16+
_CLC_OVERLOAD _CLC_DECL void __spirv_MemoryBarrier(int, int);
1717

1818
#define __CLC_NVVM_ATOMIC_IMPL_ORDER(TYPE, TYPE_NV, TYPE_MANGLED_NV, OP, \
1919
ADDR_SPACE, ADDR_SPACE_NV, ORDER) \

‎libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_load.cl

+1-1
Original file line numberDiff line numberDiff line change
@@ -11,7 +11,7 @@
1111
#include <libspirv/spirv_types.h>
1212

1313
extern int __clc_nvvm_reflect_arch();
14-
_CLC_OVERLOAD _CLC_DECL void __spirv_MemoryBarrier(unsigned int, unsigned int);
14+
_CLC_OVERLOAD _CLC_DECL void __spirv_MemoryBarrier(int, int);
1515

1616
#define __CLC_NVVM_ATOMIC_LOAD_IMPL_ORDER(TYPE, TYPE_NV, TYPE_MANGLED_NV, \
1717
ADDR_SPACE, ADDR_SPACE_NV, ORDER) \

‎libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_store.cl

+1-1
Original file line numberDiff line numberDiff line change
@@ -11,7 +11,7 @@
1111
#include <libspirv/spirv_types.h>
1212

1313
extern int __clc_nvvm_reflect_arch();
14-
_CLC_OVERLOAD _CLC_DECL void __spirv_MemoryBarrier(unsigned int, unsigned int);
14+
_CLC_OVERLOAD _CLC_DECL void __spirv_MemoryBarrier(int, int);
1515

1616
#define __CLC_NVVM_ATOMIC_STORE_IMPL_ORDER(TYPE, TYPE_NV, TYPE_MANGLED_NV, \
1717
ADDR_SPACE, ADDR_SPACE_NV, ORDER) \

‎libclc/libspirv/lib/ptx-nvidiacl/synchronization/barrier.cl

+6-6
Original file line numberDiff line numberDiff line change
@@ -11,8 +11,8 @@
1111

1212
int __clc_nvvm_reflect_arch();
1313

14-
_CLC_OVERLOAD _CLC_DEF void __spirv_MemoryBarrier(unsigned int memory,
15-
unsigned int semantics) {
14+
_CLC_OVERLOAD _CLC_DEF void __spirv_MemoryBarrier(int memory,
15+
int semantics) {
1616

1717
// for sm_70 and above membar becomes semantically identical to fence.sc.
1818
// However sm_70 and above also introduces a lightweight fence.acq_rel that
@@ -21,7 +21,7 @@ _CLC_OVERLOAD _CLC_DEF void __spirv_MemoryBarrier(unsigned int memory,
2121
// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-membar-fence
2222
// for details.
2323

24-
unsigned int order = semantics & 0x1F;
24+
int order = semantics & 0x1F;
2525
if (__clc_nvvm_reflect_arch() < 700 ||
2626
order == SequentiallyConsistent) {
2727
if (memory == CrossDevice) {
@@ -43,9 +43,9 @@ _CLC_OVERLOAD _CLC_DEF void __spirv_MemoryBarrier(unsigned int memory,
4343
}
4444

4545
_CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT void
46-
__spirv_ControlBarrier(unsigned int scope, unsigned int memory,
47-
unsigned int semantics) {
48-
unsigned int order = semantics & 0x1F;
46+
__spirv_ControlBarrier(int scope, int memory,
47+
int semantics) {
48+
int order = semantics & 0x1F;
4949
if (scope == Subgroup) {
5050
// use a full mask as barriers are required to be convergent and exited
5151
// threads can safely be in the mask

‎libclc/test/binding/core/ControlBarrier.cl

+2-2
Original file line numberDiff line numberDiff line change
@@ -16,7 +16,7 @@
1616
// CHECK-NOT: declare {{.*}} @_Z
1717
// CHECK-NOT: call {{[^ ]*}} bitcast
1818
__attribute__((overloadable)) void
19-
test___spirv_ControlBarrier(__clc_uint32_t args_0, __clc_uint32_t args_1,
20-
__clc_uint32_t args_2) {
19+
test___spirv_ControlBarrier(__clc_int32_t args_0, __clc_int32_t args_1,
20+
__clc_int32_t args_2) {
2121
__spirv_ControlBarrier(args_0, args_1, args_2);
2222
}

‎libclc/test/binding/core/MemoryBarrier.cl

+1-1
Original file line numberDiff line numberDiff line change
@@ -16,6 +16,6 @@
1616
// CHECK-NOT: declare {{.*}} @_Z
1717
// CHECK-NOT: call {{[^ ]*}} bitcast
1818
__attribute__((overloadable)) void
19-
test___spirv_MemoryBarrier(__clc_uint32_t args_0, __clc_uint32_t args_1) {
19+
test___spirv_MemoryBarrier(__clc_int32_t args_0, __clc_int32_t args_1) {
2020
__spirv_MemoryBarrier(args_0, args_1);
2121
}

‎libdevice/include/sanitizer_defs.hpp

+1-1
Original file line numberDiff line numberDiff line change
@@ -49,7 +49,7 @@ extern SYCL_EXTERNAL __SYCL_PRIVATE__ void *
4949
__spirv_GenericCastToPtrExplicit_ToPrivate(void *, int);
5050

5151
extern SYCL_EXTERNAL __attribute__((convergent)) void
52-
__spirv_ControlBarrier(uint32_t Execution, uint32_t Memory, uint32_t Semantics);
52+
__spirv_ControlBarrier(int32_t Execution, int32_t Memory, int32_t Semantics);
5353

5454
extern "C" SYCL_EXTERNAL void __devicelib_exit();
5555

‎libdevice/nativecpu_utils.cpp

+5-6
Original file line numberDiff line numberDiff line change
@@ -33,19 +33,18 @@ using __nativecpu_state = native_cpu::state;
3333
#define OCL_GLOBAL __attribute__((opencl_global))
3434
#define OCL_PRIVATE __attribute__((opencl_private))
3535

36-
DEVICE_EXTERN_C void __mux_work_group_barrier(uint32_t id, uint32_t scope,
37-
uint32_t semantics);
36+
DEVICE_EXTERN_C void __mux_work_group_barrier(int32_t id, int32_t scope,
37+
int32_t semantics);
3838
__SYCL_CONVERGENT__ DEVICE_EXTERNAL void
39-
__spirv_ControlBarrier(uint32_t Execution, uint32_t Memory,
40-
uint32_t Semantics) {
39+
__spirv_ControlBarrier(int32_t Execution, int32_t Memory, int32_t Semantics) {
4140
if (__spv::Scope::Flag::Workgroup == Execution)
4241
// todo: check id and args; use mux constants
4342
__mux_work_group_barrier(0, Execution, Semantics);
4443
}
4544

46-
DEVICE_EXTERN_C void __mux_mem_barrier(uint32_t scope, uint32_t semantics);
45+
DEVICE_EXTERN_C void __mux_mem_barrier(int32_t scope, int32_t semantics);
4746
__SYCL_CONVERGENT__ DEVICE_EXTERNAL void
48-
__spirv_MemoryBarrier(uint32_t Memory, uint32_t Semantics) {
47+
__spirv_MemoryBarrier(int32_t Memory, int32_t Semantics) {
4948
__mux_mem_barrier(Memory, Semantics);
5049
}
5150

‎llvm/lib/SYCLLowerIR/LowerWGScope.cpp

+11-8
Original file line numberDiff line numberDiff line change
@@ -162,15 +162,15 @@ enum class AddrSpace : unsigned {
162162
Output = 6
163163
};
164164

165-
enum class Scope : unsigned {
165+
enum class Scope : int {
166166
CrossDevice = 0,
167167
Device = 1,
168168
Workgroup = 2,
169169
Subgroup = 3,
170170
Invocation = 4,
171171
};
172172

173-
enum class MemorySemantics : unsigned {
173+
enum class MemorySemantics : int {
174174
None = 0x0,
175175
Acquire = 0x2,
176176
Release = 0x4,
@@ -990,7 +990,7 @@ Value *spirv::genPseudoLocalID(Instruction &Before, const Triple &TT) {
990990
// uint32_t Semantics) noexcept;
991991
Instruction *spirv::genWGBarrier(Instruction &Before, const Triple &TT) {
992992
Module &M = *Before.getModule();
993-
StringRef Name = "_Z22__spirv_ControlBarrierjjj";
993+
StringRef Name = "_Z22__spirv_ControlBarrieriii";
994994
LLVMContext &Ctx = Before.getContext();
995995
Type *ScopeTy = Type::getInt32Ty(Ctx);
996996
Type *SemanticsTy = Type::getInt32Ty(Ctx);
@@ -1006,11 +1006,14 @@ Instruction *spirv::genWGBarrier(Instruction &Before, const Triple &TT) {
10061006

10071007
IRBuilder<> Bld(Ctx);
10081008
Bld.SetInsertPoint(&Before);
1009-
auto ArgExec = ConstantInt::get(ScopeTy, asUInt(spirv::Scope::Workgroup));
1010-
auto ArgMem = ConstantInt::get(ScopeTy, asUInt(spirv::Scope::Workgroup));
1011-
auto ArgSema = ConstantInt::get(
1012-
ScopeTy, asUInt(spirv::MemorySemantics::SequentiallyConsistent) |
1013-
asUInt(spirv::MemorySemantics::WorkgroupMemory));
1009+
auto ArgExec = ConstantInt::getSigned(
1010+
ScopeTy, static_cast<int>(spirv::Scope::Workgroup));
1011+
auto ArgMem = ConstantInt::getSigned(
1012+
ScopeTy, static_cast<int>(spirv::Scope::Workgroup));
1013+
auto ArgSema = ConstantInt::getSigned(
1014+
ScopeTy,
1015+
static_cast<int>(spirv::MemorySemantics::SequentiallyConsistent) |
1016+
static_cast<int>(spirv::MemorySemantics::WorkgroupMemory));
10141017
auto BarrierCall = Bld.CreateCall(FC, {ArgExec, ArgMem, ArgSema});
10151018
BarrierCall->addFnAttr(llvm::Attribute::Convergent);
10161019
if (TT.isSPIROrSPIRV())

‎llvm/test/SYCLLowerIR/LowerWGScope/barrier-calling-conv.ll

+2-2
Original file line numberDiff line numberDiff line change
@@ -12,7 +12,7 @@ target triple = "spir64-unknown-unknown"
1212

1313
define internal spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_ENKUlNS0_5groupILi1EEEE_clES5_(ptr addrspace(4) noundef align 1 dereferenceable_or_null(1) %this, ptr noundef byval(%"class.sycl::_V1::group") align 8 %group) !work_group_scope !0 {
1414
entry:
15-
; CHECK: call spir_func void @_Z22__spirv_ControlBarrierjjj(
15+
; CHECK: call spir_func void @_Z22__spirv_ControlBarrieriii(
1616

1717
%this.addr = alloca ptr addrspace(4), align 8
1818
%this.addr.ascast = addrspacecast ptr %this.addr to ptr addrspace(4)
@@ -22,6 +22,6 @@ entry:
2222
ret void
2323
}
2424

25-
; CHECK: declare spir_func void @_Z22__spirv_ControlBarrierjjj(
25+
; CHECK: declare spir_func void @_Z22__spirv_ControlBarrieriii(
2626

2727
!0 = !{}

‎llvm/test/SYCLLowerIR/addrspacecast_handling.ll

+6-6
Original file line numberDiff line numberDiff line change
@@ -14,28 +14,28 @@ define linkonce_odr dso_local spir_func void @foo(ptr addrspace(4) dereferenceab
1414
; CHECK-NEXT: bb:
1515
; CHECK-NEXT: [[TMP0:%.*]] = alloca ptr addrspace(4), align 8
1616
; CHECK-NEXT: [[TMP1:%.*]] = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationIndex, align 4
17-
; CHECK-NEXT: call void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272) #[[ATTR0:[0-9]+]]
17+
; CHECK-NEXT: call void @_Z22__spirv_ControlBarrieriii(i32 2, i32 2, i32 272) #[[ATTR0:[0-9]+]]
1818
; CHECK-NEXT: [[CMPZ3:%.*]] = icmp eq i64 [[TMP1]], 0
1919
; CHECK-NEXT: br i1 [[CMPZ3]], label [[LEADER:%.*]], label [[MERGE:%.*]]
2020
; CHECK: leader:
2121
; CHECK-NEXT: call void @llvm.memcpy.p3.p0.i64(ptr addrspace(3) align 8 @ArgShadow, ptr align 8 [[ARG1:%.*]], i64 8, i1 false)
2222
; CHECK-NEXT: br label [[MERGE]]
2323
; CHECK: merge:
24-
; CHECK-NEXT: call void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272) #[[ATTR0]]
24+
; CHECK-NEXT: call void @_Z22__spirv_ControlBarrieriii(i32 2, i32 2, i32 272) #[[ATTR0]]
2525
; CHECK-NEXT: call void @llvm.memcpy.p0.p3.i64(ptr align 8 [[ARG1]], ptr addrspace(3) align 8 @ArgShadow, i64 8, i1 false)
2626
; CHECK-NEXT: [[TMP4:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(4)
2727
; CHECK-NEXT: [[TMP5:%.*]] = alloca [[STRUCT_SPAM:%.*]], align 8
2828
; CHECK-NEXT: [[TMP6:%.*]] = addrspacecast ptr [[TMP5]] to ptr addrspace(4)
2929
; CHECK-NEXT: [[TMP7:%.*]] = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationIndex, align 4
30-
; CHECK-NEXT: call void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272) #[[ATTR0]]
30+
; CHECK-NEXT: call void @_Z22__spirv_ControlBarrieriii(i32 2, i32 2, i32 272) #[[ATTR0]]
3131
; CHECK-NEXT: [[CMPZ:%.*]] = icmp eq i64 [[TMP7]], 0
3232
; CHECK-NEXT: br i1 [[CMPZ]], label [[WG_LEADER:%.*]], label [[WG_CF:%.*]]
3333
; CHECK: wg_leader:
3434
; CHECK-NEXT: store ptr addrspace(4) [[ARG:%.*]], ptr addrspace(4) [[TMP4]], align 8
3535
; CHECK-NEXT: br label [[WG_CF]]
3636
; CHECK: wg_cf:
3737
; CHECK-NEXT: [[TMP8:%.*]] = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationIndex, align 4
38-
; CHECK-NEXT: call void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272) #[[ATTR0]]
38+
; CHECK-NEXT: call void @_Z22__spirv_ControlBarrieriii(i32 2, i32 2, i32 272) #[[ATTR0]]
3939
; CHECK-NEXT: [[CMPZ2:%.*]] = icmp eq i64 [[TMP8]], 0
4040
; CHECK-NEXT: br i1 [[CMPZ2]], label [[TESTMAT:%.*]], label [[LEADERMAT:%.*]]
4141
; CHECK: TestMat:
@@ -44,11 +44,11 @@ define linkonce_odr dso_local spir_func void @foo(ptr addrspace(4) dereferenceab
4444
; CHECK-NEXT: store ptr addrspace(4) [[MAT_LD]], ptr addrspace(3) @WGCopy, align 8
4545
; CHECK-NEXT: br label [[LEADERMAT]]
4646
; CHECK: LeaderMat:
47-
; CHECK-NEXT: call void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272) #[[ATTR0]]
47+
; CHECK-NEXT: call void @_Z22__spirv_ControlBarrieriii(i32 2, i32 2, i32 272) #[[ATTR0]]
4848
; CHECK-NEXT: [[MAT_LD1:%.*]] = load ptr addrspace(4), ptr addrspace(3) @WGCopy, align 8
4949
; CHECK-NEXT: store ptr addrspace(4) [[MAT_LD1]], ptr [[TMP0]], align 8
5050
; CHECK-NEXT: call void @llvm.memcpy.p0.p3.i64(ptr align 8 [[TMP5]], ptr addrspace(3) align 16 @WGCopy.1, i64 36, i1 false)
51-
; CHECK-NEXT: call void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272) #[[ATTR0]]
51+
; CHECK-NEXT: call void @_Z22__spirv_ControlBarrieriii(i32 2, i32 2, i32 272) #[[ATTR0]]
5252
; CHECK-NEXT: [[TMP11:%.*]] = addrspacecast ptr [[ARG1]] to ptr addrspace(4)
5353
; CHECK-NEXT: [[TMP12:%.*]] = addrspacecast ptr addrspace(4) [[TMP6]] to ptr
5454
; CHECK-NEXT: call spir_func void @widget(ptr addrspace(4) dereferenceable_or_null(32) [[TMP11]], ptr byval([[STRUCT_SPAM]]) align 8 [[TMP12]])

‎llvm/test/SYCLLowerIR/byval_arg.ll

+2-2
Original file line numberDiff line numberDiff line change
@@ -12,14 +12,14 @@
1212
define internal spir_func void @wibble(ptr byval(%struct.baz) %arg1) !work_group_scope !0 {
1313
; CHECK-LABEL: @wibble(
1414
; CHECK-NEXT: [[TMP1:%.*]] = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationIndex
15-
; CHECK-NEXT: call void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
15+
; CHECK-NEXT: call void @_Z22__spirv_ControlBarrieriii(i32 2, i32 2, i32 272)
1616
; CHECK-NEXT: [[CMPZ:%.*]] = icmp eq i64 [[TMP1]], 0
1717
; CHECK-NEXT: br i1 [[CMPZ]], label [[LEADER:%.*]], label [[MERGE:%.*]]
1818
; CHECK: leader:
1919
; CHECK-NEXT: call void @llvm.memcpy.p3.p0.i64(ptr addrspace(3) align 8 @[[SHADOW]], ptr [[ARG1:%.*]], i64 8, i1 false)
2020
; CHECK-NEXT: br label [[MERGE]]
2121
; CHECK: merge:
22-
; CHECK-NEXT: call void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
22+
; CHECK-NEXT: call void @_Z22__spirv_ControlBarrieriii(i32 2, i32 2, i32 272)
2323
; CHECK-NEXT: call void @llvm.memcpy.p0.p3.i64(ptr [[ARG1]], ptr addrspace(3) align 8 @[[SHADOW]], i64 8, i1 false)
2424
; CHECK-NEXT: ret void
2525
;

‎llvm/test/SYCLLowerIR/byval_arg_cast.ll

+4-4
Original file line numberDiff line numberDiff line change
@@ -20,24 +20,24 @@ define dso_local spir_func void @wombat(ptr byval(%struct.widget) align 8 %arg)
2020
; CHECK-LABEL: @wombat(
2121
; CHECK-NEXT: bb:
2222
; CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationIndex, align 4
23-
; CHECK-NEXT: call void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
23+
; CHECK-NEXT: call void @_Z22__spirv_ControlBarrieriii(i32 2, i32 2, i32 272)
2424
; CHECK-NEXT: [[CMPZ1:%.*]] = icmp eq i64 [[TMP0]], 0
2525
; CHECK-NEXT: br i1 [[CMPZ1]], label [[LEADER:%.*]], label [[MERGE:%.*]]
2626
; CHECK: leader:
2727
; CHECK-NEXT: call void @llvm.memcpy.p3.p0.i64(ptr addrspace(3) align 16 @[[SHADOW]], ptr align 8 [[ARG:%.*]], i64 32, i1 false)
2828
; CHECK-NEXT: br label [[MERGE]]
2929
; CHECK: merge:
30-
; CHECK-NEXT: call void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272) #0
30+
; CHECK-NEXT: call void @_Z22__spirv_ControlBarrieriii(i32 2, i32 2, i32 272) #0
3131
; CHECK-NEXT: call void @llvm.memcpy.p0.p3.i64(ptr align 8 [[ARG]], ptr addrspace(3) align 16 @[[SHADOW]], i64 32, i1 false)
3232
; CHECK-NEXT: [[TMP3:%.*]] = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationIndex, align 4
33-
; CHECK-NEXT: call void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
33+
; CHECK-NEXT: call void @_Z22__spirv_ControlBarrieriii(i32 2, i32 2, i32 272)
3434
; CHECK-NEXT: [[CMPZ:%.*]] = icmp eq i64 [[TMP3]], 0
3535
; CHECK-NEXT: br i1 [[CMPZ]], label [[WG_LEADER:%.*]], label [[WG_CF:%.*]]
3636
; CHECK: wg_leader:
3737
; CHECK-NEXT: call void @zot(ptr [[ARG]])
3838
; CHECK-NEXT: br label [[WG_CF]]
3939
; CHECK: wg_cf:
40-
; CHECK-NEXT: call void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272) #0
40+
; CHECK-NEXT: call void @_Z22__spirv_ControlBarrieriii(i32 2, i32 2, i32 272) #0
4141
; CHECK-NEXT: ret void
4242
;
4343
bb:

‎llvm/test/SYCLLowerIR/cast_shadow.ll

+1-1
Original file line numberDiff line numberDiff line change
@@ -14,7 +14,7 @@ target triple = "nvptx64-nvidia-cuda"
1414

1515
define internal void @wobble(ptr %arg, ptr byval(%struct.spam) %arg1) !work_group_scope !0 {
1616
; CHECK: call void @llvm.memcpy.p3.p0.i64(ptr addrspace(3) align 16 @[[SHADOW]], ptr [[ARG1:%.*]], i64 32, i1 false)
17-
; CHECK: call void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272) #0
17+
; CHECK: call void @_Z22__spirv_ControlBarrieriii(i32 2, i32 2, i32 272) #0
1818
; CHECK: call void @llvm.memcpy.p0.p3.i64(ptr [[TMP11:%.*]], ptr addrspace(3) align 16 @[[SHADOW]]
1919
; CHECK: call void @widget(ptr %arg1, ptr byval(%struct.quux) [[TMP2:%.*]])
2020
;

0 commit comments

Comments
 (0)