Skip to content

Commit 69305e3

Browse files
Victor Lomullergithub-actions[bot]
Victor Lomuller
authored and
github-actions[bot]
committed
Automerge: [SPIRV] Add more id and range builtIns (#143909)
The patch adds intrinsics and lowering logic for GlobalSize, GlobalOffset, SubgroupMaxSize, NumWorkgroups, WorkgroupSize, WorkgroupId, LocalInvocationId, GlobalInvocationId, SubgroupSize, NumSubgroups, SubgroupId and SubgroupLocalInvocationId SPIR-V builtins. The patch also extend spv_thread_id, spv_group_id and spv_thread_id_in_group to return anyint rather than i32. This allows the intrinsics to support the opencl environment. For each of the intrinsics, new clang builtins were added as well as a binding for the SPIR-V "friendly" format. The original format doesn't define such binding (uses global variables) but it is not possible to express the Input SC which is normally required by the environement specs, and using builtin functions is the most usual approach for other backend and programming models.
2 parents ef41bb3 + 27c9b55 commit 69305e3

File tree

19 files changed

+761
-37
lines changed

19 files changed

+761
-37
lines changed

‎clang/include/clang/Basic/BuiltinsSPIRVCL.td‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10,3 +10,6 @@ include "clang/Basic/BuiltinsSPIRVBase.td"
1010

1111
def generic_cast_to_ptr_explicit
1212
: SPIRVBuiltin<"void*(void*, int)", [NoThrow, Const, CustomTypeChecking]>;
13+
def global_size : SPIRVBuiltin<"size_t(int)", [NoThrow, Const]>;
14+
def global_offset : SPIRVBuiltin<"size_t(int)", [NoThrow, Const]>;
15+
def subgroup_max_size : SPIRVBuiltin<"uint32_t()", [NoThrow, Const]>;

‎clang/include/clang/Basic/BuiltinsSPIRVCommon.td‎

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -8,6 +8,16 @@
88

99
include "clang/Basic/BuiltinsSPIRVBase.td"
1010

11+
def num_workgroups : SPIRVBuiltin<"size_t(int)", [NoThrow, Const]>;
12+
def workgroup_size : SPIRVBuiltin<"size_t(int)", [NoThrow, Const]>;
13+
def workgroup_id : SPIRVBuiltin<"size_t(int)", [NoThrow, Const]>;
14+
def local_invocation_id : SPIRVBuiltin<"size_t(int)", [NoThrow, Const]>;
15+
def global_invocation_id : SPIRVBuiltin<"size_t(int)", [NoThrow, Const]>;
16+
def subgroup_size : SPIRVBuiltin<"uint32_t()", [NoThrow, Const]>;
17+
def num_subgroups : SPIRVBuiltin<"uint32_t()", [NoThrow, Const]>;
18+
def subgroup_id : SPIRVBuiltin<"uint32_t()", [NoThrow, Const]>;
19+
def subgroup_local_invocation_id : SPIRVBuiltin<"uint32_t()", [NoThrow, Const]>;
20+
1121
def distance : SPIRVBuiltin<"void(...)", [NoThrow, Const]>;
1222
def length : SPIRVBuiltin<"void(...)", [NoThrow, Const]>;
1323
def smoothstep : SPIRVBuiltin<"void(...)", [NoThrow, Const, CustomTypeChecking]>;

‎clang/lib/CodeGen/CGHLSLRuntime.cpp‎

Lines changed: 13 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -393,17 +393,27 @@ llvm::Value *CGHLSLRuntime::emitInputSemantic(IRBuilder<> &B,
393393
return B.CreateCall(FunctionCallee(GroupIndex));
394394
}
395395
if (D.hasAttr<HLSLSV_DispatchThreadIDAttr>()) {
396+
llvm::Intrinsic::ID IntrinID = getThreadIdIntrinsic();
396397
llvm::Function *ThreadIDIntrinsic =
397-
CGM.getIntrinsic(getThreadIdIntrinsic());
398+
llvm::Intrinsic::isOverloaded(IntrinID)
399+
? CGM.getIntrinsic(IntrinID, {CGM.Int32Ty})
400+
: CGM.getIntrinsic(IntrinID);
398401
return buildVectorInput(B, ThreadIDIntrinsic, Ty);
399402
}
400403
if (D.hasAttr<HLSLSV_GroupThreadIDAttr>()) {
404+
llvm::Intrinsic::ID IntrinID = getGroupThreadIdIntrinsic();
401405
llvm::Function *GroupThreadIDIntrinsic =
402-
CGM.getIntrinsic(getGroupThreadIdIntrinsic());
406+
llvm::Intrinsic::isOverloaded(IntrinID)
407+
? CGM.getIntrinsic(IntrinID, {CGM.Int32Ty})
408+
: CGM.getIntrinsic(IntrinID);
403409
return buildVectorInput(B, GroupThreadIDIntrinsic, Ty);
404410
}
405411
if (D.hasAttr<HLSLSV_GroupIDAttr>()) {
406-
llvm::Function *GroupIDIntrinsic = CGM.getIntrinsic(getGroupIdIntrinsic());
412+
llvm::Intrinsic::ID IntrinID = getGroupIdIntrinsic();
413+
llvm::Function *GroupIDIntrinsic =
414+
llvm::Intrinsic::isOverloaded(IntrinID)
415+
? CGM.getIntrinsic(IntrinID, {CGM.Int32Ty})
416+
: CGM.getIntrinsic(IntrinID);
407417
return buildVectorInput(B, GroupIDIntrinsic, Ty);
408418
}
409419
if (D.hasAttr<HLSLSV_PositionAttr>()) {

‎clang/lib/CodeGen/TargetBuiltins/SPIR.cpp‎

Lines changed: 42 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -97,6 +97,48 @@ Value *CodeGenFunction::EmitSPIRVBuiltinExpr(unsigned BuiltinID,
9797
Call->addRetAttr(llvm::Attribute::AttrKind::NoUndef);
9898
return Call;
9999
}
100+
case SPIRV::BI__builtin_spirv_num_workgroups:
101+
return Builder.CreateIntrinsic(
102+
/*ReturnType=*/getTypes().ConvertType(E->getType()),
103+
Intrinsic::spv_num_workgroups,
104+
ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr,
105+
"spv.num.workgroups");
106+
case SPIRV::BI__builtin_spirv_workgroup_size:
107+
return Builder.CreateIntrinsic(
108+
/*ReturnType=*/getTypes().ConvertType(E->getType()),
109+
Intrinsic::spv_workgroup_size,
110+
ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr,
111+
"spv.workgroup.size");
112+
case SPIRV::BI__builtin_spirv_workgroup_id:
113+
return Builder.CreateIntrinsic(
114+
/*ReturnType=*/getTypes().ConvertType(E->getType()),
115+
Intrinsic::spv_group_id,
116+
ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr,
117+
"spv.group.id");
118+
case SPIRV::BI__builtin_spirv_local_invocation_id:
119+
return Builder.CreateIntrinsic(
120+
/*ReturnType=*/getTypes().ConvertType(E->getType()),
121+
Intrinsic::spv_thread_id_in_group,
122+
ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr,
123+
"spv.thread.id.in.group");
124+
case SPIRV::BI__builtin_spirv_global_invocation_id:
125+
return Builder.CreateIntrinsic(
126+
/*ReturnType=*/getTypes().ConvertType(E->getType()),
127+
Intrinsic::spv_thread_id,
128+
ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr,
129+
"spv.thread.id");
130+
case SPIRV::BI__builtin_spirv_global_size:
131+
return Builder.CreateIntrinsic(
132+
/*ReturnType=*/getTypes().ConvertType(E->getType()),
133+
Intrinsic::spv_global_size,
134+
ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr,
135+
"spv.num.workgroups");
136+
case SPIRV::BI__builtin_spirv_global_offset:
137+
return Builder.CreateIntrinsic(
138+
/*ReturnType=*/getTypes().ConvertType(E->getType()),
139+
Intrinsic::spv_global_offset,
140+
ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr,
141+
"spv.global.offset");
100142
}
101143
return nullptr;
102144
}

‎clang/lib/Headers/__clang_spirv_builtins.h‎

Lines changed: 39 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -16,6 +16,12 @@
1616
#define __SPIRV_NOEXCEPT
1717
#endif
1818

19+
#pragma push_macro("__size_t")
20+
#pragma push_macro("__uint32_t")
21+
#pragma push_macro("__uint64_t")
22+
#define __size_t __SIZE_TYPE__
23+
#define __uint32_t __UINT32_TYPE__
24+
1925
#define __SPIRV_overloadable __attribute__((overloadable))
2026
#define __SPIRV_convergent __attribute__((convergent))
2127
#define __SPIRV_inline __attribute__((always_inline))
@@ -36,13 +42,41 @@
3642
// to establish if we can use the builtin alias. We disable builtin altogether
3743
// if we do not intent to use the backend. So instead of use target macros, rely
3844
// on a __has_builtin test.
39-
#if (__has_builtin(__builtin_spirv_generic_cast_to_ptr_explicit))
45+
#if (__has_builtin(__builtin_spirv_num_workgroups))
4046
#define __SPIRV_BUILTIN_ALIAS(builtin) \
4147
__attribute__((clang_builtin_alias(builtin)))
4248
#else
4349
#define __SPIRV_BUILTIN_ALIAS(builtin)
4450
#endif
4551

52+
// Builtin IDs and sizes
53+
54+
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_num_workgroups) __size_t
55+
__spirv_NumWorkgroups(int);
56+
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_workgroup_size) __size_t
57+
__spirv_WorkgroupSize(int);
58+
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_workgroup_id) __size_t
59+
__spirv_WorkgroupId(int);
60+
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_local_invocation_id) __size_t
61+
__spirv_LocalInvocationId(int);
62+
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_global_invocation_id) __size_t
63+
__spirv_GlobalInvocationId(int);
64+
65+
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_global_size) __size_t
66+
__spirv_GlobalSize(int);
67+
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_global_offset) __size_t
68+
__spirv_GlobalOffset(int);
69+
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_size) __uint32_t
70+
__spirv_SubgroupSize();
71+
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_max_size) __uint32_t
72+
__spirv_SubgroupMaxSize();
73+
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_num_subgroups) __uint32_t
74+
__spirv_NumSubgroups();
75+
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_id) __uint32_t
76+
__spirv_SubgroupId();
77+
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_local_invocation_id)
78+
__uint32_t __spirv_SubgroupLocalInvocationId();
79+
4680
// OpGenericCastToPtrExplicit
4781

4882
extern __SPIRV_overloadable
@@ -164,6 +198,10 @@ __spirv_GenericCastToPtr_ToPrivate(__generic const volatile void *p,
164198
return (__private const volatile void *)p;
165199
}
166200

201+
#pragma pop_macro("__size_t")
202+
#pragma pop_macro("__uint32_t")
203+
#pragma pop_macro("__uint64_t")
204+
167205
#undef __SPIRV_overloadable
168206
#undef __SPIRV_convergent
169207
#undef __SPIRV_inline

‎clang/test/CodeGenHLSL/semantics/DispatchThreadID.hlsl‎

Lines changed: 5 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -5,17 +5,19 @@
55

66
// CHECK: define void @foo()
77
// CHECK-DXIL: %[[#ID:]] = call i32 @llvm.[[TARGET]].thread.id(i32 0)
8-
// CHECK-SPIRV: %[[#ID:]] = call i32 @llvm.[[TARGET]].thread.id(i32 0)
8+
// CHECK-SPIRV: %[[#ID:]] = call i32 @llvm.[[TARGET]].thread.id.i32(i32 0)
99
// CHECK-DXIL: call void @{{.*}}foo{{.*}}(i32 %[[#ID]])
1010
// CHECK-SPIRV: call spir_func void @{{.*}}foo{{.*}}(i32 %[[#ID]])
1111
[shader("compute")]
1212
[numthreads(8,8,1)]
1313
void foo(uint Idx : SV_DispatchThreadID) {}
1414

1515
// CHECK: define void @bar()
16-
// CHECK: %[[#ID_X:]] = call i32 @llvm.[[TARGET]].thread.id(i32 0)
16+
// CHECK-DXIL: %[[#ID_X:]] = call i32 @llvm.[[TARGET]].thread.id(i32 0)
17+
// CHECK-SPIRV: %[[#ID_X:]] = call i32 @llvm.[[TARGET]].thread.id.i32(i32 0)
1718
// CHECK: %[[#ID_X_:]] = insertelement <2 x i32> poison, i32 %[[#ID_X]], i64 0
18-
// CHECK: %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].thread.id(i32 1)
19+
// CHECK-DXIL: %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].thread.id(i32 1)
20+
// CHECK-SPIRV: %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].thread.id.i32(i32 1)
1921
// CHECK: %[[#ID_XY:]] = insertelement <2 x i32> %[[#ID_X_]], i32 %[[#ID_Y]], i64 1
2022
// CHECK-DXIL: call void @{{.*}}bar{{.*}}(<2 x i32> %[[#ID_XY]])
2123
// CHECK-SPIRV: call spir_func void @{{.*}}bar{{.*}}(<2 x i32> %[[#ID_XY]])

‎clang/test/CodeGenHLSL/semantics/SV_GroupID.hlsl‎

Lines changed: 12 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -4,17 +4,20 @@
44
// Make sure SV_GroupID translated into dx.group.id for directx target and spv.group.id for spirv target.
55

66
// CHECK: define void @foo()
7-
// CHECK: %[[#ID:]] = call i32 @llvm.[[TARGET]].group.id(i32 0)
7+
// CHECK-DXIL: %[[#ID:]] = call i32 @llvm.[[TARGET]].group.id(i32 0)
8+
// CHECK-SPIRV: %[[#ID:]] = call i32 @llvm.[[TARGET]].group.id.i32(i32 0)
89
// CHECK-DXIL: call void @{{.*}}foo{{.*}}(i32 %[[#ID]])
910
// CHECK-SPIRV: call spir_func void @{{.*}}foo{{.*}}(i32 %[[#ID]])
1011
[shader("compute")]
1112
[numthreads(8,8,1)]
1213
void foo(uint Idx : SV_GroupID) {}
1314

1415
// CHECK: define void @bar()
15-
// CHECK: %[[#ID_X:]] = call i32 @llvm.[[TARGET]].group.id(i32 0)
16+
// CHECK-DXIL: %[[#ID_X:]] = call i32 @llvm.[[TARGET]].group.id(i32 0)
17+
// CHECK-SPIRV: %[[#ID_X:]] = call i32 @llvm.[[TARGET]].group.id.i32(i32 0)
1618
// CHECK: %[[#ID_X_:]] = insertelement <2 x i32> poison, i32 %[[#ID_X]], i64 0
17-
// CHECK: %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].group.id(i32 1)
19+
// CHECK-DXIL: %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].group.id(i32 1)
20+
// CHECK-SPIRV: %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].group.id.i32(i32 1)
1821
// CHECK: %[[#ID_XY:]] = insertelement <2 x i32> %[[#ID_X_]], i32 %[[#ID_Y]], i64 1
1922
// CHECK-DXIL: call void @{{.*}}bar{{.*}}(<2 x i32> %[[#ID_XY]])
2023
// CHECK-SPIRV: call spir_func void @{{.*}}bar{{.*}}(<2 x i32> %[[#ID_XY]])
@@ -23,11 +26,14 @@ void foo(uint Idx : SV_GroupID) {}
2326
void bar(uint2 Idx : SV_GroupID) {}
2427

2528
// CHECK: define void @test()
26-
// CHECK: %[[#ID_X:]] = call i32 @llvm.[[TARGET]].group.id(i32 0)
29+
// CHECK-DXIL: %[[#ID_X:]] = call i32 @llvm.[[TARGET]].group.id(i32 0)
30+
// CHECK-SPIRV: %[[#ID_X:]] = call i32 @llvm.[[TARGET]].group.id.i32(i32 0)
2731
// CHECK: %[[#ID_X_:]] = insertelement <3 x i32> poison, i32 %[[#ID_X]], i64 0
28-
// CHECK: %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].group.id(i32 1)
32+
// CHECK-DXIL: %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].group.id(i32 1)
33+
// CHECK-SPIRV: %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].group.id.i32(i32 1)
2934
// CHECK: %[[#ID_XY:]] = insertelement <3 x i32> %[[#ID_X_]], i32 %[[#ID_Y]], i64 1
30-
// CHECK: %[[#ID_Z:]] = call i32 @llvm.[[TARGET]].group.id(i32 2)
35+
// CHECK-DXIL: %[[#ID_Z:]] = call i32 @llvm.[[TARGET]].group.id(i32 2)
36+
// CHECK-SPIRV: %[[#ID_Z:]] = call i32 @llvm.[[TARGET]].group.id.i32(i32 2)
3137
// CHECK: %[[#ID_XYZ:]] = insertelement <3 x i32> %[[#ID_XY]], i32 %[[#ID_Z]], i64 2
3238
// CHECK-DXIL: call void @{{.*}}test{{.*}}(<3 x i32> %[[#ID_XYZ]])
3339
// CHECK-SPIRV: call spir_func void @{{.*}}test{{.*}}(<3 x i32> %[[#ID_XYZ]])

‎clang/test/CodeGenHLSL/semantics/SV_GroupThreadID.hlsl‎

Lines changed: 12 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -4,17 +4,20 @@
44
// Make sure SV_GroupThreadID translated into dx.thread.id.in.group for directx target and spv.thread.id.in.group for spirv target.
55

66
// CHECK: define void @foo()
7-
// CHECK: %[[#ID:]] = call i32 @llvm.[[TARGET]].thread.id.in.group(i32 0)
7+
// CHECK-DXIL: %[[#ID:]] = call i32 @llvm.[[TARGET]].thread.id.in.group(i32 0)
8+
// CHECK-SPIRV: %[[#ID:]] = call i32 @llvm.[[TARGET]].thread.id.in.group.i32(i32 0)
89
// CHECK-DXIL: call void @{{.*}}foo{{.*}}(i32 %[[#ID]])
910
// CHECK-SPIRV: call spir_func void @{{.*}}foo{{.*}}(i32 %[[#ID]])
1011
[shader("compute")]
1112
[numthreads(8,8,1)]
1213
void foo(uint Idx : SV_GroupThreadID) {}
1314

1415
// CHECK: define void @bar()
15-
// CHECK: %[[#ID_X:]] = call i32 @llvm.[[TARGET]].thread.id.in.group(i32 0)
16+
// CHECK-DXIL: %[[#ID_X:]] = call i32 @llvm.[[TARGET]].thread.id.in.group(i32 0)
17+
// CHECK-SPIRV: %[[#ID_X:]] = call i32 @llvm.[[TARGET]].thread.id.in.group.i32(i32 0)
1618
// CHECK: %[[#ID_X_:]] = insertelement <2 x i32> poison, i32 %[[#ID_X]], i64 0
17-
// CHECK: %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].thread.id.in.group(i32 1)
19+
// CHECK-DXIL: %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].thread.id.in.group(i32 1)
20+
// CHECK-SPIRV: %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].thread.id.in.group.i32(i32 1)
1821
// CHECK: %[[#ID_XY:]] = insertelement <2 x i32> %[[#ID_X_]], i32 %[[#ID_Y]], i64 1
1922
// CHECK-DXIL: call void @{{.*}}bar{{.*}}(<2 x i32> %[[#ID_XY]])
2023
// CHECK-SPIRV: call spir_func void @{{.*}}bar{{.*}}(<2 x i32> %[[#ID_XY]])
@@ -23,11 +26,14 @@ void foo(uint Idx : SV_GroupThreadID) {}
2326
void bar(uint2 Idx : SV_GroupThreadID) {}
2427

2528
// CHECK: define void @test()
26-
// CHECK: %[[#ID_X:]] = call i32 @llvm.[[TARGET]].thread.id.in.group(i32 0)
29+
// CHECK-DXIL: %[[#ID_X:]] = call i32 @llvm.[[TARGET]].thread.id.in.group(i32 0)
30+
// CHECK-SPIRV: %[[#ID_X:]] = call i32 @llvm.[[TARGET]].thread.id.in.group.i32(i32 0)
2731
// CHECK: %[[#ID_X_:]] = insertelement <3 x i32> poison, i32 %[[#ID_X]], i64 0
28-
// CHECK: %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].thread.id.in.group(i32 1)
32+
// CHECK-DXIL: %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].thread.id.in.group(i32 1)
33+
// CHECK-SPIRV: %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].thread.id.in.group.i32(i32 1)
2934
// CHECK: %[[#ID_XY:]] = insertelement <3 x i32> %[[#ID_X_]], i32 %[[#ID_Y]], i64 1
30-
// CHECK: %[[#ID_Z:]] = call i32 @llvm.[[TARGET]].thread.id.in.group(i32 2)
35+
// CHECK-DXIL: %[[#ID_Z:]] = call i32 @llvm.[[TARGET]].thread.id.in.group(i32 2)
36+
// CHECK-SPIRV: %[[#ID_Z:]] = call i32 @llvm.[[TARGET]].thread.id.in.group.i32(i32 2)
3137
// CHECK: %[[#ID_XYZ:]] = insertelement <3 x i32> %[[#ID_XY]], i32 %[[#ID_Z]], i64 2
3238
// CHECK-DXIL: call void @{{.*}}test{{.*}}(<3 x i32> %[[#ID_XYZ]])
3339
// CHECK-SPIRV: call spir_func void @{{.*}}test{{.*}}(<3 x i32> %[[#ID_XYZ]])

‎clang/test/CodeGenSPIRV/Builtins/ids_and_ranges.c‎

Lines changed: 106 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,106 @@
1+
// RUN: %clang_cc1 -O1 -triple spirv64 -fsycl-is-device %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,CHECK64
2+
// RUN: %clang_cc1 -O1 -triple spirv64 -cl-std=CL3.0 -x cl %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,CHECK64
3+
// RUN: %clang_cc1 -O1 -triple spirv32 -cl-std=CL3.0 -x cl %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,CHECK32
4+
5+
// CHECK: @test_num_workgroups(
6+
// CHECK-NEXT: [[ENTRY:.*:]]
7+
// CHECK64-NEXT: tail call i64 @llvm.spv.num.workgroups.i64(i32 0)
8+
// CHECK32-NEXT: tail call i32 @llvm.spv.num.workgroups.i32(i32 0)
9+
//
10+
unsigned int test_num_workgroups() {
11+
return __builtin_spirv_num_workgroups(0);
12+
}
13+
14+
// CHECK: @test_workgroup_size(
15+
// CHECK-NEXT: [[ENTRY:.*:]]
16+
// CHECK64-NEXT: tail call i64 @llvm.spv.workgroup.size.i64(i32 0)
17+
// CHECK32-NEXT: tail call i32 @llvm.spv.workgroup.size.i32(i32 0)
18+
//
19+
unsigned int test_workgroup_size() {
20+
return __builtin_spirv_workgroup_size(0);
21+
}
22+
23+
// CHECK: @test_workgroup_id(
24+
// CHECK-NEXT: [[ENTRY:.*:]]
25+
// CHECK64-NEXT: tail call i64 @llvm.spv.group.id.i64(i32 0)
26+
// CHECK32-NEXT: tail call i32 @llvm.spv.group.id.i32(i32 0)
27+
//
28+
unsigned int test_workgroup_id() {
29+
return __builtin_spirv_workgroup_id(0);
30+
}
31+
32+
// CHECK: @test_local_invocation_id(
33+
// CHECK-NEXT: [[ENTRY:.*:]]
34+
// CHECK64-NEXT: tail call i64 @llvm.spv.thread.id.in.group.i64(i32 0)
35+
// CHECK32-NEXT: tail call i32 @llvm.spv.thread.id.in.group.i32(i32 0)
36+
//
37+
unsigned int test_local_invocation_id() {
38+
return __builtin_spirv_local_invocation_id(0);
39+
}
40+
41+
// CHECK: @test_global_invocation_id(
42+
// CHECK-NEXT: [[ENTRY:.*:]]
43+
// CHECK64-NEXT: tail call i64 @llvm.spv.thread.id.i64(i32 0)
44+
// CHECK32-NEXT: tail call i32 @llvm.spv.thread.id.i32(i32 0)
45+
//
46+
unsigned int test_global_invocation_id() {
47+
return __builtin_spirv_global_invocation_id(0);
48+
}
49+
50+
// CHECK: @test_global_size(
51+
// CHECK-NEXT: [[ENTRY:.*:]]
52+
// CHECK64-NEXT: tail call i64 @llvm.spv.global.size.i64(i32 0)
53+
// CHECK32-NEXT: tail call i32 @llvm.spv.global.size.i32(i32 0)
54+
//
55+
unsigned int test_global_size() {
56+
return __builtin_spirv_global_size(0);
57+
}
58+
59+
// CHECK: @test_global_offset(
60+
// CHECK-NEXT: [[ENTRY:.*:]]
61+
// CHECK64-NEXT: tail call i64 @llvm.spv.global.offset.i64(i32 0)
62+
// CHECK32-NEXT: tail call i32 @llvm.spv.global.offset.i32(i32 0)
63+
//
64+
unsigned int test_global_offset() {
65+
return __builtin_spirv_global_offset(0);
66+
}
67+
68+
// CHECK: @test_subgroup_size(
69+
// CHECK-NEXT: [[ENTRY:.*:]]
70+
// CHECK-NEXT: tail call i32 @llvm.spv.subgroup.size()
71+
//
72+
unsigned int test_subgroup_size() {
73+
return __builtin_spirv_subgroup_size();
74+
}
75+
76+
// CHECK: @test_subgroup_max_size(
77+
// CHECK-NEXT: [[ENTRY:.*:]]
78+
// CHECK-NEXT: tail call i32 @llvm.spv.subgroup.max.size()
79+
//
80+
unsigned int test_subgroup_max_size() {
81+
return __builtin_spirv_subgroup_max_size();
82+
}
83+
84+
// CHECK: @test_num_subgroups(
85+
// CHECK-NEXT: [[ENTRY:.*:]]
86+
// CHECK-NEXT: tail call i32 @llvm.spv.num.subgroups()
87+
//
88+
unsigned int test_num_subgroups() {
89+
return __builtin_spirv_num_subgroups();
90+
}
91+
92+
// CHECK: @test_subgroup_id(
93+
// CHECK-NEXT: [[ENTRY:.*:]]
94+
// CHECK-NEXT: tail call i32 @llvm.spv.subgroup.id()
95+
//
96+
unsigned int test_subgroup_id() {
97+
return __builtin_spirv_subgroup_id();
98+
}
99+
100+
// CHECK: @test_subgroup_local_invocation_id(
101+
// CHECK-NEXT: [[ENTRY:.*:]]
102+
// CHECK-NEXT: tail call i32 @llvm.spv.subgroup.local.invocation.id()
103+
//
104+
unsigned int test_subgroup_local_invocation_id() {
105+
return __builtin_spirv_subgroup_local_invocation_id();
106+
}

0 commit comments

Comments
 (0)