Skip to content

[NFC][SYCL] Don't emit OpenCL metadata for SYCL #17723

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 3 commits into from
Mar 31, 2025
Merged

[NFC][SYCL] Don't emit OpenCL metadata for SYCL #17723

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
99c7644
[NFC][SYCL] Don't emit OpenCL metadata for SYCL
bader Mar 28, 2025
a348614
Update metadata checks in LIT tests.
bader Mar 28, 2025
5c428ee
Code clean-up and formatting.
bader Mar 28, 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
59 changes: 29 additions & 30 deletions clang/lib/CodeGen/CodeGenModule.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2592,40 +2592,39 @@ void CodeGenModule::GenKernelArgMetadata(llvm::Function *Fn,
}
}

bool IsEsimdFunction = FD && FD->hasAttr<SYCLSimdAttr>();

if (LangOpts.SYCLIsDevice && !IsEsimdFunction) {
Fn->setMetadata("kernel_arg_buffer_location",
llvm::MDNode::get(VMContext, argSYCLBufferLocationAttr));
// Generate this metadata only if atleast one kernel argument is an
// accessor.
if (isKernelArgAnAccessor) {
Fn->setMetadata("kernel_arg_runtime_aligned",
llvm::MDNode::get(VMContext, argSYCLAccessorPtrs));
Fn->setMetadata("kernel_arg_exclusive_ptr",
if (getLangOpts().SYCLIsDevice) {
if (FD && FD->hasAttr<SYCLSimdAttr>()) {
Fn->setMetadata("kernel_arg_accessor_ptr",
llvm::MDNode::get(VMContext, argSYCLAccessorPtrs));
}
} else {
if (getLangOpts().OpenCL || getLangOpts().SYCLIsDevice) {
Fn->setMetadata("kernel_arg_addr_space",
llvm::MDNode::get(VMContext, addressQuals));
Fn->setMetadata("kernel_arg_access_qual",
llvm::MDNode::get(VMContext, accessQuals));
Fn->setMetadata("kernel_arg_type",
llvm::MDNode::get(VMContext, argTypeNames));
Fn->setMetadata("kernel_arg_base_type",
llvm::MDNode::get(VMContext, argBaseTypeNames));
Fn->setMetadata("kernel_arg_type_qual",
llvm::MDNode::get(VMContext, argTypeQuals));
if (IsEsimdFunction)
Fn->setMetadata("kernel_arg_accessor_ptr",
} else {
Fn->setMetadata("kernel_arg_buffer_location",
llvm::MDNode::get(VMContext, argSYCLBufferLocationAttr));
// Generate this metadata only if at least one kernel argument is an
// accessor.
if (isKernelArgAnAccessor) {
Fn->setMetadata("kernel_arg_runtime_aligned",
llvm::MDNode::get(VMContext, argSYCLAccessorPtrs));
Fn->setMetadata("kernel_arg_exclusive_ptr",
llvm::MDNode::get(VMContext, argSYCLAccessorPtrs));
}
}
if (getCodeGenOpts().EmitOpenCLArgMetadata ||
getCodeGenOpts().HIPSaveKernelArgName)
Fn->setMetadata("kernel_arg_name",
llvm::MDNode::get(VMContext, argNames));
}

if (getLangOpts().OpenCL) {
Fn->setMetadata("kernel_arg_addr_space",
llvm::MDNode::get(VMContext, addressQuals));
Fn->setMetadata("kernel_arg_access_qual",
llvm::MDNode::get(VMContext, accessQuals));
Fn->setMetadata("kernel_arg_type",
llvm::MDNode::get(VMContext, argTypeNames));
Fn->setMetadata("kernel_arg_base_type",
llvm::MDNode::get(VMContext, argBaseTypeNames));
Fn->setMetadata("kernel_arg_type_qual",
llvm::MDNode::get(VMContext, argTypeQuals));
}
if (getCodeGenOpts().EmitOpenCLArgMetadata ||
getCodeGenOpts().HIPSaveKernelArgName)
Fn->setMetadata("kernel_arg_name", llvm::MDNode::get(VMContext, argNames));
}

/// Determines whether the language options require us to model
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/kernel-op-calls.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -53,7 +53,7 @@ int main() {

Q.submit([&](sycl::handler& cgh) {
ESIMDFunctor EF;
// CHECK: define dso_local spir_kernel void @_ZTS12ESIMDFunctor() {{.*}} !intel_reqd_sub_group_size !{{[0-9]+}} !work_group_size_hint !{{[0-9]+}} !kernel_arg_addr_space !{{[0-9]+}} !kernel_arg_access_qual !{{[0-9]+}} !kernel_arg_type !{{[0-9]+}} !kernel_arg_base_type !{{[0-9]+}} !kernel_arg_type_qual !{{[0-9]+}} !kernel_arg_accessor_ptr !{{[0-9]+}} !sycl_explicit_simd !{{[0-9]+}} !sycl_fixed_targets !{{[0-9]+}} {
// CHECK: define dso_local spir_kernel void @_ZTS12ESIMDFunctor() {{.*}} !intel_reqd_sub_group_size !{{[0-9]+}} !work_group_size_hint !{{[0-9]+}} !kernel_arg_accessor_ptr !{{[0-9]+}} !sycl_explicit_simd !{{[0-9]+}} !sycl_fixed_targets !{{[0-9]+}} {
cgh.parallel_for(sycl::range<1>(10), EF);
});

Expand Down
5 changes: 2 additions & 3 deletions sycl/test/check_device_code/esimd/dae.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -15,10 +15,9 @@ __attribute__((sycl_kernel)) void my_kernel(Func kernelFunc) {

SYCL_EXTERNAL SYCL_ESIMD_FUNCTION ESIMD_NOINLINE void callee(int x) {}

// CHECK: define dso_local spir_kernel {{.*}} !kernel_arg_addr_space ![[#MD:]]
// CHECK: !kernel_arg_access_qual ![[#MD]] !kernel_arg_type ![[#MD]] !kernel_arg_base_type ![[#MD]] !kernel_arg_type_qual ![[#MD]] !kernel_arg_accessor_ptr ![[#MD]]
// CHECK: define dso_local spir_kernel {{.*}} !sycl_kernel_omit_args ![[#MD:]]
Copy link
Contributor Author

@bader bader Mar 28, 2025 •
edited
Loading

Choose a reason for hiding this comment

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

@intel/dpcpp-esimd-reviewers, I'm sure I understand the purpose of this test, but I tried to check that "dead argument elimination" optimization attaches metadata correctly.
It looks like original test verifies that other metadata do not refer to deleted argument. Should I check for !kernel_arg_accessor_ptr?
I removed it because "dead argument elimination" should not change the value of this metadata. It will be empty in any case.

Copy link
Contributor

@sarnex sarnex Mar 31, 2025

Choose a reason for hiding this comment

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

your change looks fine thx

SYCL_EXTERNAL void __attribute__((noinline)) caller(int x) {
my_kernel<class kernel_abc>([=]() SYCL_ESIMD_KERNEL { callee(x); });
}

//CHECK: [[#MD]] = !{}
//CHECK: [[#MD]] = !{i1 true}
Loading