Skip to content

Commit f4d81f6

Browse files
maarquitos14jsji
maarquitos14
authored and
jsji
committed
[SPIRV] id and range builtins integration for SYCL (#19639)
27c9b55 was added upstream, but because it has limited SYCL support, tests can't run properly.
1 parent 2f58eb0 commit f4d81f6

File tree

3 files changed

+37
-25
lines changed

3 files changed

+37
-25
lines changed

‎clang/lib/Headers/__clang_spirv_builtins.h‎

Lines changed: 12 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -53,30 +53,30 @@
5353

5454
// Builtin IDs and sizes
5555

56-
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_num_workgroups) __size_t
56+
extern __SPIRV_SYCL_EXTERNAL __SPIRV_BUILTIN_ALIAS(__builtin_spirv_num_workgroups) __size_t
5757
__spirv_NumWorkgroups(int);
58-
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_workgroup_size) __size_t
58+
extern __SPIRV_SYCL_EXTERNAL __SPIRV_BUILTIN_ALIAS(__builtin_spirv_workgroup_size) __size_t
5959
__spirv_WorkgroupSize(int);
60-
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_workgroup_id) __size_t
60+
extern __SPIRV_SYCL_EXTERNAL __SPIRV_BUILTIN_ALIAS(__builtin_spirv_workgroup_id) __size_t
6161
__spirv_WorkgroupId(int);
62-
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_local_invocation_id) __size_t
62+
extern __SPIRV_SYCL_EXTERNAL __SPIRV_BUILTIN_ALIAS(__builtin_spirv_local_invocation_id) __size_t
6363
__spirv_LocalInvocationId(int);
64-
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_global_invocation_id) __size_t
64+
extern __SPIRV_SYCL_EXTERNAL __SPIRV_BUILTIN_ALIAS(__builtin_spirv_global_invocation_id) __size_t
6565
__spirv_GlobalInvocationId(int);
6666

67-
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_global_size) __size_t
67+
extern __SPIRV_SYCL_EXTERNAL __SPIRV_BUILTIN_ALIAS(__builtin_spirv_global_size) __size_t
6868
__spirv_GlobalSize(int);
69-
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_global_offset) __size_t
69+
extern __SPIRV_SYCL_EXTERNAL __SPIRV_BUILTIN_ALIAS(__builtin_spirv_global_offset) __size_t
7070
__spirv_GlobalOffset(int);
71-
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_size) __uint32_t
71+
extern __SPIRV_SYCL_EXTERNAL __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_size) __uint32_t
7272
__spirv_SubgroupSize();
73-
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_max_size) __uint32_t
73+
extern __SPIRV_SYCL_EXTERNAL __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_max_size) __uint32_t
7474
__spirv_SubgroupMaxSize();
75-
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_num_subgroups) __uint32_t
75+
extern __SPIRV_SYCL_EXTERNAL __SPIRV_BUILTIN_ALIAS(__builtin_spirv_num_subgroups) __uint32_t
7676
__spirv_NumSubgroups();
77-
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_id) __uint32_t
77+
extern __SPIRV_SYCL_EXTERNAL __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_id) __uint32_t
7878
__spirv_SubgroupId();
79-
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_local_invocation_id)
79+
extern __SPIRV_SYCL_EXTERNAL __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_local_invocation_id)
8080
__uint32_t __spirv_SubgroupLocalInvocationId();
8181

8282
// OpGenericCastToPtrExplicit

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

Lines changed: 19 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -2,12 +2,19 @@
22
// RUN: %clang_cc1 -O1 -triple spirv64 -cl-std=CL3.0 -x cl %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,CHECK64
33
// RUN: %clang_cc1 -O1 -triple spirv32 -cl-std=CL3.0 -x cl %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,CHECK32
44

5+
#ifdef __SYCL_DEVICE_ONLY__
6+
#define SYCL_EXTERNAL __attribute__((sycl_device))
7+
#else
8+
#define SYCL_EXTERNAL
9+
#endif
10+
11+
512
// CHECK: @test_num_workgroups(
613
// CHECK-NEXT: [[ENTRY:.*:]]
714
// CHECK64-NEXT: tail call i64 @llvm.spv.num.workgroups.i64(i32 0)
815
// CHECK32-NEXT: tail call i32 @llvm.spv.num.workgroups.i32(i32 0)
916
//
10-
unsigned int test_num_workgroups() {
17+
SYCL_EXTERNAL unsigned int test_num_workgroups() {
1118
return __builtin_spirv_num_workgroups(0);
1219
}
1320

@@ -16,7 +23,7 @@ unsigned int test_num_workgroups() {
1623
// CHECK64-NEXT: tail call i64 @llvm.spv.workgroup.size.i64(i32 0)
1724
// CHECK32-NEXT: tail call i32 @llvm.spv.workgroup.size.i32(i32 0)
1825
//
19-
unsigned int test_workgroup_size() {
26+
SYCL_EXTERNAL unsigned int test_workgroup_size() {
2027
return __builtin_spirv_workgroup_size(0);
2128
}
2229

@@ -25,7 +32,7 @@ unsigned int test_workgroup_size() {
2532
// CHECK64-NEXT: tail call i64 @llvm.spv.group.id.i64(i32 0)
2633
// CHECK32-NEXT: tail call i32 @llvm.spv.group.id.i32(i32 0)
2734
//
28-
unsigned int test_workgroup_id() {
35+
SYCL_EXTERNAL unsigned int test_workgroup_id() {
2936
return __builtin_spirv_workgroup_id(0);
3037
}
3138

@@ -34,7 +41,7 @@ unsigned int test_workgroup_id() {
3441
// CHECK64-NEXT: tail call i64 @llvm.spv.thread.id.in.group.i64(i32 0)
3542
// CHECK32-NEXT: tail call i32 @llvm.spv.thread.id.in.group.i32(i32 0)
3643
//
37-
unsigned int test_local_invocation_id() {
44+
SYCL_EXTERNAL unsigned int test_local_invocation_id() {
3845
return __builtin_spirv_local_invocation_id(0);
3946
}
4047

@@ -43,7 +50,7 @@ unsigned int test_local_invocation_id() {
4350
// CHECK64-NEXT: tail call i64 @llvm.spv.thread.id.i64(i32 0)
4451
// CHECK32-NEXT: tail call i32 @llvm.spv.thread.id.i32(i32 0)
4552
//
46-
unsigned int test_global_invocation_id() {
53+
SYCL_EXTERNAL unsigned int test_global_invocation_id() {
4754
return __builtin_spirv_global_invocation_id(0);
4855
}
4956

@@ -52,7 +59,7 @@ unsigned int test_global_invocation_id() {
5259
// CHECK64-NEXT: tail call i64 @llvm.spv.global.size.i64(i32 0)
5360
// CHECK32-NEXT: tail call i32 @llvm.spv.global.size.i32(i32 0)
5461
//
55-
unsigned int test_global_size() {
62+
SYCL_EXTERNAL unsigned int test_global_size() {
5663
return __builtin_spirv_global_size(0);
5764
}
5865

@@ -61,46 +68,46 @@ unsigned int test_global_size() {
6168
// CHECK64-NEXT: tail call i64 @llvm.spv.global.offset.i64(i32 0)
6269
// CHECK32-NEXT: tail call i32 @llvm.spv.global.offset.i32(i32 0)
6370
//
64-
unsigned int test_global_offset() {
71+
SYCL_EXTERNAL unsigned int test_global_offset() {
6572
return __builtin_spirv_global_offset(0);
6673
}
6774

6875
// CHECK: @test_subgroup_size(
6976
// CHECK-NEXT: [[ENTRY:.*:]]
7077
// CHECK-NEXT: tail call i32 @llvm.spv.subgroup.size()
7178
//
72-
unsigned int test_subgroup_size() {
79+
SYCL_EXTERNAL unsigned int test_subgroup_size() {
7380
return __builtin_spirv_subgroup_size();
7481
}
7582

7683
// CHECK: @test_subgroup_max_size(
7784
// CHECK-NEXT: [[ENTRY:.*:]]
7885
// CHECK-NEXT: tail call i32 @llvm.spv.subgroup.max.size()
7986
//
80-
unsigned int test_subgroup_max_size() {
87+
SYCL_EXTERNAL unsigned int test_subgroup_max_size() {
8188
return __builtin_spirv_subgroup_max_size();
8289
}
8390

8491
// CHECK: @test_num_subgroups(
8592
// CHECK-NEXT: [[ENTRY:.*:]]
8693
// CHECK-NEXT: tail call i32 @llvm.spv.num.subgroups()
8794
//
88-
unsigned int test_num_subgroups() {
95+
SYCL_EXTERNAL unsigned int test_num_subgroups() {
8996
return __builtin_spirv_num_subgroups();
9097
}
9198

9299
// CHECK: @test_subgroup_id(
93100
// CHECK-NEXT: [[ENTRY:.*:]]
94101
// CHECK-NEXT: tail call i32 @llvm.spv.subgroup.id()
95102
//
96-
unsigned int test_subgroup_id() {
103+
SYCL_EXTERNAL unsigned int test_subgroup_id() {
97104
return __builtin_spirv_subgroup_id();
98105
}
99106

100107
// CHECK: @test_subgroup_local_invocation_id(
101108
// CHECK-NEXT: [[ENTRY:.*:]]
102109
// CHECK-NEXT: tail call i32 @llvm.spv.subgroup.local.invocation.id()
103110
//
104-
unsigned int test_subgroup_local_invocation_id() {
111+
SYCL_EXTERNAL unsigned int test_subgroup_local_invocation_id() {
105112
return __builtin_spirv_subgroup_local_invocation_id();
106113
}

‎clang/test/Headers/spirv_ids.cpp‎

Lines changed: 6 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -4,6 +4,11 @@
44
// 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
55
// 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
66

7+
#ifdef __SYCL_DEVICE_ONLY__
8+
#define SYCL_EXTERNAL __attribute__((sycl_device))
9+
#else
10+
#define SYCL_EXTERNAL
11+
#endif
712

813
// CHECK64: call i64 @llvm.spv.num.workgroups.i64(i32 0)
914
// CHECK64: call i64 @llvm.spv.num.workgroups.i64(i32 1)
@@ -80,7 +85,7 @@
8085
// NV: call noundef i32 @_Z18__spirv_SubgroupIdv() #2
8186
// NV: call noundef i32 @_Z33__spirv_SubgroupLocalInvocationIdv() #2
8287

83-
void test_id_and_range() {
88+
SYCL_EXTERNAL void test_id_and_range() {
8489
__spirv_NumWorkgroups(0);
8590
__spirv_NumWorkgroups(1);
8691
__spirv_NumWorkgroups(2);

0 commit comments

Comments
 (0)