Skip to content

Resolve sycl-web test failures caused by https://github.com/llvm/llvm-project/pull/140282 #19895

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 4 commits into from
Sep 2, 2025
+19 −13
Merged

Resolve sycl-web test failures caused by https://github.com/llvm/llvm-project/pull/140282 #19895

Show file tree
Hide file tree
Changes from 3 commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
5e620a1
This is a test; does this open a PR against the fork?
schittir Aug 26, 2025
4865cbe
Add sycl_device attribute to pass failing test
schittir Aug 27, 2025
5ce1b9f
Add sycl_external attribute to the check in CodeGenModule
schittir Sep 1, 2025
c126b29
Remove sycl_device attribute from the tests and change || to &&
schittir Sep 2, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion clang/lib/CodeGen/CodeGenModule.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6901,7 +6901,7 @@ CodeGenModule::getLLVMLinkageForDeclarator(const DeclaratorDecl *D,
// have this, linkonce_odr suffices. If -fno-sycl-rdc is passed, we know there
// is only one translation unit and can so mark them internal.
if (getLangOpts().SYCLIsDevice && !D->hasAttr<DeviceKernelAttr>() &&
!D->hasAttr<SYCLDeviceAttr>() &&
!(D->hasAttr<SYCLDeviceAttr>() || D->hasAttr<SYCLExternalAttr>()) &&
!SemaSYCL::isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
D->getType()))
return getLangOpts().GPURelocatableDeviceCode
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We'll need to do something similar to this in upstream as part of fixing emit of inline functions in device code.

Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSPIRV/Builtins/generic_cast_to_ptr_explicit.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,7 +3,7 @@
// RUN: %clang_cc1 -O1 -triple spirv32 -cl-std=CL3.0 -x cl %s -emit-llvm -o - | FileCheck %s

#ifdef __SYCL_DEVICE_ONLY__
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Something seems off here. Why would __SYCL_DEVICE_ONLY__ be defined for the runs that don't include -fsycl-is-device? -x cl doesn't somehow imply SYCL, does it?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The very first line runs clang with -fsycl-is-device.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yes, and this makes sense for that RUN line. My question is about the other two RUN lines that pass -x cl and don't pass -fsycl-is-device. The code presumes that __SYCL_DEVICE_ONLY__ is defined for those RUN lines as well and I don't understand why that would be the case. I also find it odd that an attribute named sycl_device is expected to be used in a non-SYCL OpenCL compilation. I would expect an attribute named opencl_device or similar instead.

#define SYCL_EXTERNAL [[clang::sycl_external]]
#define SYCL_EXTERNAL [[clang::sycl_external]] __attribute__((sycl_device))
#else
#define SYCL_EXTERNAL
#endif
Expand Down
30 changes: 18 additions & 12 deletions clang/test/CodeGenSPIRV/Builtins/ids_and_ranges.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,12 +2,18 @@
// RUN: %clang_cc1 -O1 -triple spirv64 -cl-std=CL3.0 -x cl %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,CHECK64
// RUN: %clang_cc1 -O1 -triple spirv32 -cl-std=CL3.0 -x cl %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,CHECK32

#ifdef __SYCL_DEVICE_ONLY__
#define SYCL_EXTERNAL [[clang::sycl_external]] __attribute__((sycl_device))
#else
#define SYCL_EXTERNAL
#endif

// CHECK: @{{.*}}test_num_workgroups{{.*}}(
// CHECK-NEXT: [[ENTRY:.*:]]
// CHECK64-NEXT: tail call i64 @llvm.spv.num.workgroups.i64(i32 0)
// CHECK32-NEXT: tail call i32 @llvm.spv.num.workgroups.i32(i32 0)
//
[[clang::sycl_external]] unsigned int test_num_workgroups() {
SYCL_EXTERNAL unsigned int test_num_workgroups() {
return __builtin_spirv_num_workgroups(0);
}

Expand All @@ -16,7 +22,7 @@
// CHECK64-NEXT: tail call i64 @llvm.spv.workgroup.size.i64(i32 0)
// CHECK32-NEXT: tail call i32 @llvm.spv.workgroup.size.i32(i32 0)
//
[[clang::sycl_external]] unsigned int test_workgroup_size() {
SYCL_EXTERNAL unsigned int test_workgroup_size() {
return __builtin_spirv_workgroup_size(0);
}

Expand All @@ -25,7 +31,7 @@
// CHECK64-NEXT: tail call i64 @llvm.spv.group.id.i64(i32 0)
// CHECK32-NEXT: tail call i32 @llvm.spv.group.id.i32(i32 0)
//
[[clang::sycl_external]] unsigned int test_workgroup_id() {
SYCL_EXTERNAL unsigned int test_workgroup_id() {
return __builtin_spirv_workgroup_id(0);
}

Expand All @@ -34,7 +40,7 @@
// CHECK64-NEXT: tail call i64 @llvm.spv.thread.id.in.group.i64(i32 0)
// CHECK32-NEXT: tail call i32 @llvm.spv.thread.id.in.group.i32(i32 0)
//
[[clang::sycl_external]] unsigned int test_local_invocation_id() {
SYCL_EXTERNAL unsigned int test_local_invocation_id() {
return __builtin_spirv_local_invocation_id(0);
}

Expand All @@ -43,7 +49,7 @@
// CHECK64-NEXT: tail call i64 @llvm.spv.thread.id.i64(i32 0)
// CHECK32-NEXT: tail call i32 @llvm.spv.thread.id.i32(i32 0)
//
[[clang::sycl_external]] unsigned int test_global_invocation_id() {
SYCL_EXTERNAL unsigned int test_global_invocation_id() {
return __builtin_spirv_global_invocation_id(0);
}

Expand All @@ -52,7 +58,7 @@
// CHECK64-NEXT: tail call i64 @llvm.spv.global.size.i64(i32 0)
// CHECK32-NEXT: tail call i32 @llvm.spv.global.size.i32(i32 0)
//
[[clang::sycl_external]] unsigned int test_global_size() {
SYCL_EXTERNAL unsigned int test_global_size() {
return __builtin_spirv_global_size(0);
}

Expand All @@ -61,46 +67,46 @@
// CHECK64-NEXT: tail call i64 @llvm.spv.global.offset.i64(i32 0)
// CHECK32-NEXT: tail call i32 @llvm.spv.global.offset.i32(i32 0)
//
[[clang::sycl_external]] unsigned int test_global_offset() {
SYCL_EXTERNAL unsigned int test_global_offset() {
return __builtin_spirv_global_offset(0);
}

// CHECK: @{{.*}}test_subgroup_size{{.*}}(
// CHECK-NEXT: [[ENTRY:.*:]]
// CHECK-NEXT: tail call i32 @llvm.spv.subgroup.size()
//
[[clang::sycl_external]] unsigned int test_subgroup_size() {
SYCL_EXTERNAL unsigned int test_subgroup_size() {
return __builtin_spirv_subgroup_size();
}

// CHECK: @{{.*}}test_subgroup_max_size{{.*}}(
// CHECK-NEXT: [[ENTRY:.*:]]
// CHECK-NEXT: tail call i32 @llvm.spv.subgroup.max.size()
//
[[clang::sycl_external]] unsigned int test_subgroup_max_size() {
SYCL_EXTERNAL unsigned int test_subgroup_max_size() {
return __builtin_spirv_subgroup_max_size();
}

// CHECK: @{{.*}}test_num_subgroups{{.*}}(
// CHECK-NEXT: [[ENTRY:.*:]]
// CHECK-NEXT: tail call i32 @llvm.spv.num.subgroups()
//
[[clang::sycl_external]] unsigned int test_num_subgroups() {
SYCL_EXTERNAL unsigned int test_num_subgroups() {
return __builtin_spirv_num_subgroups();
}

// CHECK: @{{.*}}test_subgroup_id{{.*}}(
// CHECK-NEXT: [[ENTRY:.*:]]
// CHECK-NEXT: tail call i32 @llvm.spv.subgroup.id()
//
[[clang::sycl_external]] unsigned int test_subgroup_id() {
SYCL_EXTERNAL unsigned int test_subgroup_id() {
return __builtin_spirv_subgroup_id();
}

// CHECK: @{{.*}}test_subgroup_local_invocation_id{{.*}}(
// CHECK-NEXT: [[ENTRY:.*:]]
// CHECK-NEXT: tail call i32 @llvm.spv.subgroup.local.invocation.id()
//
[[clang::sycl_external]] unsigned int test_subgroup_local_invocation_id() {
SYCL_EXTERNAL unsigned int test_subgroup_local_invocation_id() {
return __builtin_spirv_subgroup_local_invocation_id();
}