Skip to content

Commit 65d926c

Browse files
authored
[SYCL][NativeCPU] Set target attrs for subkernels. (#19169)
This fixes LLVMIntrinsicLowering/intrinsic_cmp.cpp with -mavx. Target attributes can affect the ABI, and the subkernel and kernel need to agree on the ABI in order for the call to work.
1 parent 0df8a45 commit 65d926c

File tree

1 file changed

+7
-1
lines changed

1 file changed

+7
-1
lines changed

llvm/lib/SYCLNativeCPUUtils/PrepareSYCLNativeCPU.cpp

Lines changed: 7 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -114,10 +114,16 @@ void emitSubkernelForKernel(Function *F, Type *NativeCPUArgDescType,
114114

115115
fixCallingConv(F);
116116
fixCallingConv(SubhF);
117-
// Add sycl-module-id attribute
117+
// Add sycl-module-id and target attributes
118118
// Todo: we may want to copy other attributes to the subhandler,
119119
// but we can't simply use setAttributes(F->getAttributes) since
120120
// the function signatures are different
121+
if (F->hasFnAttribute("target-features")) {
122+
SubhF->addFnAttr(F->getFnAttribute("target-features"));
123+
}
124+
if (F->hasFnAttribute("target-cpu")) {
125+
SubhF->addFnAttr(F->getFnAttribute("target-cpu"));
126+
}
121127
if (F->hasFnAttribute(sycl::utils::ATTR_SYCL_MODULE_ID)) {
122128
Attribute MId = F->getFnAttribute(sycl::utils::ATTR_SYCL_MODULE_ID);
123129
SubhF->addFnAttr("sycl-module-id", MId.getValueAsString());

0 commit comments

Comments
 (0)