Skip to content

Commit 735ce84

Browse files
NaghasanNaghasan
committed
[SPIRV] Add more id and range builtIns
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.
1 parent 3374263 commit 735ce84

File tree

16 files changed

+718
-22
lines changed

16 files changed

+718
-22
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: 4 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -394,16 +394,17 @@ llvm::Value *CGHLSLRuntime::emitInputSemantic(IRBuilder<> &B,
394394
}
395395
if (D.hasAttr<HLSLSV_DispatchThreadIDAttr>()) {
396396
llvm::Function *ThreadIDIntrinsic =
397-
CGM.getIntrinsic(getThreadIdIntrinsic());
397+
CGM.getIntrinsic(getThreadIdIntrinsic(), CGM.Int32Ty);
398398
return buildVectorInput(B, ThreadIDIntrinsic, Ty);
399399
}
400400
if (D.hasAttr<HLSLSV_GroupThreadIDAttr>()) {
401401
llvm::Function *GroupThreadIDIntrinsic =
402-
CGM.getIntrinsic(getGroupThreadIdIntrinsic());
402+
CGM.getIntrinsic(getGroupThreadIdIntrinsic(), CGM.Int32Ty);
403403
return buildVectorInput(B, GroupThreadIDIntrinsic, Ty);
404404
}
405405
if (D.hasAttr<HLSLSV_GroupIDAttr>()) {
406-
llvm::Function *GroupIDIntrinsic = CGM.getIntrinsic(getGroupIdIntrinsic());
406+
llvm::Function *GroupIDIntrinsic =
407+
CGM.getIntrinsic(getGroupIdIntrinsic(), CGM.Int32Ty);
407408
return buildVectorInput(B, GroupIDIntrinsic, Ty);
408409
}
409410
assert(false && "Unhandled parameter attribute");

‎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: 34 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -16,6 +16,11 @@
1616
#define __SPIRV_NOEXCEPT
1717
#endif
1818

19+
#if (!defined(__OPENCL_CPP_VERSION__) && !defined(__OPENCL_C_VERSION__))
20+
#include <stddef.h>
21+
#include <stdint.h>
22+
#endif
23+
1924
#define __SPIRV_overloadable __attribute__((overloadable))
2025
#define __SPIRV_convergent __attribute__((convergent))
2126
#define __SPIRV_inline __attribute__((always_inline))
@@ -36,13 +41,41 @@
3641
// to establish if we can use the builtin alias. We disable builtin altogether
3742
// if we do not intent to use the backend. So instead of use target macros, rely
3843
// on a __has_builtin test.
39-
#if (__has_builtin(__builtin_spirv_generic_cast_to_ptr_explicit))
44+
#if (__has_builtin(__builtin_spirv_num_workgroups))
4045
#define __SPIRV_BUILTIN_ALIAS(builtin) \
4146
__attribute__((clang_builtin_alias(builtin)))
4247
#else
4348
#define __SPIRV_BUILTIN_ALIAS(builtin)
4449
#endif
4550

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

4881
extern __SPIRV_overloadable

‎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+
}

‎clang/test/Headers/spirv_ids.cpp

Lines changed: 110 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,110 @@
1+
// RUN: %clang_cc1 -Wno-unused-value -O0 -internal-isystem %S/../../lib/Headers -include __clang_spirv_builtins.h -triple spirv64 -emit-llvm %s -fsycl-is-device -o - | FileCheck %s -check-prefixes=CHECK64
2+
// RUN: %clang_cc1 -Wno-unused-value -O0 -internal-isystem %S/../../lib/Headers -include __clang_spirv_builtins.h -triple spirv64 -emit-llvm %s -x cl -o - | FileCheck %s -check-prefixes=CHECK64
3+
// RUN: %clang_cc1 -Wno-unused-value -O0 -internal-isystem %S/../../lib/Headers -include __clang_spirv_builtins.h -triple spirv32 -emit-llvm %s -fsycl-is-device -o - | FileCheck %s -check-prefixes=CHECK32
4+
// RUN: %clang_cc1 -Wno-unused-value -O0 -internal-isystem %S/../../lib/Headers -include __clang_spirv_builtins.h -triple spirv32 -emit-llvm %s -x cl -o - | FileCheck %s -check-prefixes=CHECK32
5+
// RUN: %clang_cc1 -Wno-unused-value -O0 -internal-isystem %S/../../lib/Headers -include __clang_spirv_builtins.h -triple nvptx64 -emit-llvm %s -fsycl-is-device -o - | FileCheck %s -check-prefixes=NV
6+
7+
8+
// CHECK64: call i64 @llvm.spv.num.workgroups.i64(i32 0)
9+
// CHECK64: call i64 @llvm.spv.num.workgroups.i64(i32 1)
10+
// CHECK64: call i64 @llvm.spv.num.workgroups.i64(i32 2)
11+
// CHECK64: call i64 @llvm.spv.workgroup.size.i64(i32 0)
12+
// CHECK64: call i64 @llvm.spv.workgroup.size.i64(i32 1)
13+
// CHECK64: call i64 @llvm.spv.workgroup.size.i64(i32 2)
14+
// CHECK64: call i64 @llvm.spv.group.id.i64(i32 0)
15+
// CHECK64: call i64 @llvm.spv.group.id.i64(i32 1)
16+
// CHECK64: call i64 @llvm.spv.group.id.i64(i32 2)
17+
// CHECK64: call i64 @llvm.spv.thread.id.in.group.i64(i32 0)
18+
// CHECK64: call i64 @llvm.spv.thread.id.in.group.i64(i32 1)
19+
// CHECK64: call i64 @llvm.spv.thread.id.in.group.i64(i32 2)
20+
// CHECK64: call i64 @llvm.spv.thread.id.i64(i32 0)
21+
// CHECK64: call i64 @llvm.spv.thread.id.i64(i32 1)
22+
// CHECK64: call i64 @llvm.spv.thread.id.i64(i32 2)
23+
// CHECK64: call i64 @llvm.spv.global.size.i64(i32 0)
24+
// CHECK64: call i64 @llvm.spv.global.size.i64(i32 1)
25+
// CHECK64: call i64 @llvm.spv.global.size.i64(i32 2)
26+
// CHECK64: call i64 @llvm.spv.global.offset.i64(i32 0)
27+
// CHECK64: call i64 @llvm.spv.global.offset.i64(i32 1)
28+
// CHECK64: call i64 @llvm.spv.global.offset.i64(i32 2)
29+
// CHECK32: call i32 @llvm.spv.num.workgroups.i32(i32 0)
30+
// CHECK32: call i32 @llvm.spv.num.workgroups.i32(i32 1)
31+
// CHECK32: call i32 @llvm.spv.num.workgroups.i32(i32 2)
32+
// CHECK32: call i32 @llvm.spv.workgroup.size.i32(i32 0)
33+
// CHECK32: call i32 @llvm.spv.workgroup.size.i32(i32 1)
34+
// CHECK32: call i32 @llvm.spv.workgroup.size.i32(i32 2)
35+
// CHECK32: call i32 @llvm.spv.group.id.i32(i32 0)
36+
// CHECK32: call i32 @llvm.spv.group.id.i32(i32 1)
37+
// CHECK32: call i32 @llvm.spv.group.id.i32(i32 2)
38+
// CHECK32: call i32 @llvm.spv.thread.id.in.group.i32(i32 0)
39+
// CHECK32: call i32 @llvm.spv.thread.id.in.group.i32(i32 1)
40+
// CHECK32: call i32 @llvm.spv.thread.id.in.group.i32(i32 2)
41+
// CHECK32: call i32 @llvm.spv.thread.id.i32(i32 0)
42+
// CHECK32: call i32 @llvm.spv.thread.id.i32(i32 1)
43+
// CHECK32: call i32 @llvm.spv.thread.id.i32(i32 2)
44+
// CHECK32: call i32 @llvm.spv.global.size.i32(i32 0)
45+
// CHECK32: call i32 @llvm.spv.global.size.i32(i32 1)
46+
// CHECK32: call i32 @llvm.spv.global.size.i32(i32 2)
47+
// CHECK32: call i32 @llvm.spv.global.offset.i32(i32 0)
48+
// CHECK32: call i32 @llvm.spv.global.offset.i32(i32 1)
49+
// CHECK32: call i32 @llvm.spv.global.offset.i32(i32 2)
50+
// CHECK: call i32 @llvm.spv.subgroup.size()
51+
// CHECK: call i32 @llvm.spv.subgroup.max.size()
52+
// CHECK: call i32 @llvm.spv.num.subgroups()
53+
// CHECK: call i32 @llvm.spv.subgroup.id()
54+
// CHECK: call i32 @llvm.spv.subgroup.local.invocation.id()
55+
56+
// NV: call noundef i64 @_Z21__spirv_NumWorkgroupsi(i32 noundef 0) #2
57+
// NV: call noundef i64 @_Z21__spirv_NumWorkgroupsi(i32 noundef 1) #2
58+
// NV: call noundef i64 @_Z21__spirv_NumWorkgroupsi(i32 noundef 2) #2
59+
// NV: call noundef i64 @_Z21__spirv_WorkgroupSizei(i32 noundef 0) #2
60+
// NV: call noundef i64 @_Z21__spirv_WorkgroupSizei(i32 noundef 1) #2
61+
// NV: call noundef i64 @_Z21__spirv_WorkgroupSizei(i32 noundef 2) #2
62+
// NV: call noundef i64 @_Z19__spirv_WorkgroupIdi(i32 noundef 0) #2
63+
// NV: call noundef i64 @_Z19__spirv_WorkgroupIdi(i32 noundef 1) #2
64+
// NV: call noundef i64 @_Z19__spirv_WorkgroupIdi(i32 noundef 2) #2
65+
// NV: call noundef i64 @_Z25__spirv_LocalInvocationIdi(i32 noundef 0) #2
66+
// NV: call noundef i64 @_Z25__spirv_LocalInvocationIdi(i32 noundef 1) #2
67+
// NV: call noundef i64 @_Z25__spirv_LocalInvocationIdi(i32 noundef 2) #2
68+
// NV: call noundef i64 @_Z26__spirv_GlobalInvocationIdi(i32 noundef 0) #2
69+
// NV: call noundef i64 @_Z26__spirv_GlobalInvocationIdi(i32 noundef 1) #2
70+
// NV: call noundef i64 @_Z26__spirv_GlobalInvocationIdi(i32 noundef 2) #2
71+
// NV: call noundef i64 @_Z18__spirv_GlobalSizei(i32 noundef 0) #2
72+
// NV: call noundef i64 @_Z18__spirv_GlobalSizei(i32 noundef 1) #2
73+
// NV: call noundef i64 @_Z18__spirv_GlobalSizei(i32 noundef 2) #2
74+
// NV: call noundef i64 @_Z20__spirv_GlobalOffseti(i32 noundef 0) #2
75+
// NV: call noundef i64 @_Z20__spirv_GlobalOffseti(i32 noundef 1) #2
76+
// NV: call noundef i64 @_Z20__spirv_GlobalOffseti(i32 noundef 2) #2
77+
// NV: call noundef i32 @_Z20__spirv_SubgroupSizev() #2
78+
// NV: call noundef i32 @_Z23__spirv_SubgroupMaxSizev() #2
79+
// NV: call noundef i32 @_Z20__spirv_NumSubgroupsv() #2
80+
// NV: call noundef i32 @_Z18__spirv_SubgroupIdv() #2
81+
// NV: call noundef i32 @_Z33__spirv_SubgroupLocalInvocationIdv() #2
82+
83+
void test_id_and_range() {
84+
__spirv_NumWorkgroups(0);
85+
__spirv_NumWorkgroups(1);
86+
__spirv_NumWorkgroups(2);
87+
__spirv_WorkgroupSize(0);
88+
__spirv_WorkgroupSize(1);
89+
__spirv_WorkgroupSize(2);
90+
__spirv_WorkgroupId(0);
91+
__spirv_WorkgroupId(1);
92+
__spirv_WorkgroupId(2);
93+
__spirv_LocalInvocationId(0);
94+
__spirv_LocalInvocationId(1);
95+
__spirv_LocalInvocationId(2);
96+
__spirv_GlobalInvocationId(0);
97+
__spirv_GlobalInvocationId(1);
98+
__spirv_GlobalInvocationId(2);
99+
__spirv_GlobalSize(0);
100+
__spirv_GlobalSize(1);
101+
__spirv_GlobalSize(2);
102+
__spirv_GlobalOffset(0);
103+
__spirv_GlobalOffset(1);
104+
__spirv_GlobalOffset(2);
105+
unsigned int ssize = __spirv_SubgroupSize();
106+
unsigned int smax = __spirv_SubgroupMaxSize();
107+
unsigned int snum = __spirv_NumSubgroups();
108+
unsigned int sid = __spirv_SubgroupId();
109+
unsigned int sinvocid = __spirv_SubgroupLocalInvocationId();
110+
}

0 commit comments

Comments
 (0)