Skip to content

Commit fad7dbb

Browse files
authored
[NFC][SYCL] Don't emit OpenCL metadata for SYCL (#17723)
`kernel_arg_addr_space`, `kernel_arg_access_qual`, `kernel_arg_type`, `kernel_arg_base_type` and `kernel_arg_type_qual` metadata are designed to provide information about kernel arguments to [clGetKernelInfo](https://registry.khronos.org/OpenCL/specs/3.0-unified/html/OpenCL_API.html#_kernel_object_queries) function. SYCL does not provide similar information through API.
1 parent 501114f commit fad7dbb

File tree

3 files changed

+32
-34
lines changed

3 files changed

+32
-34
lines changed

clang/lib/CodeGen/CodeGenModule.cpp

+29-30
Original file line numberDiff line numberDiff line change
@@ -2592,40 +2592,39 @@ void CodeGenModule::GenKernelArgMetadata(llvm::Function *Fn,
25922592
}
25932593
}
25942594

2595-
bool IsEsimdFunction = FD && FD->hasAttr<SYCLSimdAttr>();
2596-
2597-
if (LangOpts.SYCLIsDevice && !IsEsimdFunction) {
2598-
Fn->setMetadata("kernel_arg_buffer_location",
2599-
llvm::MDNode::get(VMContext, argSYCLBufferLocationAttr));
2600-
// Generate this metadata only if atleast one kernel argument is an
2601-
// accessor.
2602-
if (isKernelArgAnAccessor) {
2603-
Fn->setMetadata("kernel_arg_runtime_aligned",
2604-
llvm::MDNode::get(VMContext, argSYCLAccessorPtrs));
2605-
Fn->setMetadata("kernel_arg_exclusive_ptr",
2595+
if (getLangOpts().SYCLIsDevice) {
2596+
if (FD && FD->hasAttr<SYCLSimdAttr>()) {
2597+
Fn->setMetadata("kernel_arg_accessor_ptr",
26062598
llvm::MDNode::get(VMContext, argSYCLAccessorPtrs));
2607-
}
2608-
} else {
2609-
if (getLangOpts().OpenCL || getLangOpts().SYCLIsDevice) {
2610-
Fn->setMetadata("kernel_arg_addr_space",
2611-
llvm::MDNode::get(VMContext, addressQuals));
2612-
Fn->setMetadata("kernel_arg_access_qual",
2613-
llvm::MDNode::get(VMContext, accessQuals));
2614-
Fn->setMetadata("kernel_arg_type",
2615-
llvm::MDNode::get(VMContext, argTypeNames));
2616-
Fn->setMetadata("kernel_arg_base_type",
2617-
llvm::MDNode::get(VMContext, argBaseTypeNames));
2618-
Fn->setMetadata("kernel_arg_type_qual",
2619-
llvm::MDNode::get(VMContext, argTypeQuals));
2620-
if (IsEsimdFunction)
2621-
Fn->setMetadata("kernel_arg_accessor_ptr",
2599+
} else {
2600+
Fn->setMetadata("kernel_arg_buffer_location",
2601+
llvm::MDNode::get(VMContext, argSYCLBufferLocationAttr));
2602+
// Generate this metadata only if at least one kernel argument is an
2603+
// accessor.
2604+
if (isKernelArgAnAccessor) {
2605+
Fn->setMetadata("kernel_arg_runtime_aligned",
26222606
llvm::MDNode::get(VMContext, argSYCLAccessorPtrs));
2607+
Fn->setMetadata("kernel_arg_exclusive_ptr",
2608+
llvm::MDNode::get(VMContext, argSYCLAccessorPtrs));
2609+
}
26232610
}
2624-
if (getCodeGenOpts().EmitOpenCLArgMetadata ||
2625-
getCodeGenOpts().HIPSaveKernelArgName)
2626-
Fn->setMetadata("kernel_arg_name",
2627-
llvm::MDNode::get(VMContext, argNames));
26282611
}
2612+
2613+
if (getLangOpts().OpenCL) {
2614+
Fn->setMetadata("kernel_arg_addr_space",
2615+
llvm::MDNode::get(VMContext, addressQuals));
2616+
Fn->setMetadata("kernel_arg_access_qual",
2617+
llvm::MDNode::get(VMContext, accessQuals));
2618+
Fn->setMetadata("kernel_arg_type",
2619+
llvm::MDNode::get(VMContext, argTypeNames));
2620+
Fn->setMetadata("kernel_arg_base_type",
2621+
llvm::MDNode::get(VMContext, argBaseTypeNames));
2622+
Fn->setMetadata("kernel_arg_type_qual",
2623+
llvm::MDNode::get(VMContext, argTypeQuals));
2624+
}
2625+
if (getCodeGenOpts().EmitOpenCLArgMetadata ||
2626+
getCodeGenOpts().HIPSaveKernelArgName)
2627+
Fn->setMetadata("kernel_arg_name", llvm::MDNode::get(VMContext, argNames));
26292628
}
26302629

26312630
/// Determines whether the language options require us to model

clang/test/CodeGenSYCL/kernel-op-calls.cpp

+1-1
Original file line numberDiff line numberDiff line change
@@ -53,7 +53,7 @@ int main() {
5353

5454
Q.submit([&](sycl::handler& cgh) {
5555
ESIMDFunctor EF;
56-
// 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]+}} {
56+
// 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]+}} {
5757
cgh.parallel_for(sycl::range<1>(10), EF);
5858
});
5959

sycl/test/check_device_code/esimd/dae.cpp

+2-3
Original file line numberDiff line numberDiff line change
@@ -15,10 +15,9 @@ __attribute__((sycl_kernel)) void my_kernel(Func kernelFunc) {
1515

1616
SYCL_EXTERNAL SYCL_ESIMD_FUNCTION ESIMD_NOINLINE void callee(int x) {}
1717

18-
// CHECK: define dso_local spir_kernel {{.*}} !kernel_arg_addr_space ![[#MD:]]
19-
// 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]]
18+
// CHECK: define dso_local spir_kernel {{.*}} !sycl_kernel_omit_args ![[#MD:]]
2019
SYCL_EXTERNAL void __attribute__((noinline)) caller(int x) {
2120
my_kernel<class kernel_abc>([=]() SYCL_ESIMD_KERNEL { callee(x); });
2221
}
2322

24-
//CHECK: [[#MD]] = !{}
23+
//CHECK: [[#MD]] = !{i1 true}

0 commit comments

Comments
 (0)