Skip to content
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.

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

[SYCL-MLIR][clang] Attach SYCLKernelObjFunc to kernel body functions #12345

Merged
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
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
10 changes: 10 additions & 0 deletions clang/include/clang/Basic/Attr.td
Original file line number Diff line number Diff line change
Expand Up @@ -1499,6 +1499,16 @@ def SYCLAccessorReadonly : Attr {
let Documentation = [InternalOnly];
}

// Used to mark SYCL kernel object functions implementing SYCL kernels.
// The Kernels arguments corresponds to the kernels this function implements.
def SYCLKernelObjFunc : Attr {
// This attribute has no spellings as it is only ever created implicitly.
let Spellings = [];
let Args = [VariadicStringArgument<"Kernels">];
let SemaHandler = 0;
let Documentation = [InternalOnly];
}

// The attribute denotes that it is a function written in a scalar fashion, which
// is used in ESIMD context and needs to be vectorized by a vector backend compiler.
// For now, this attribute will be used only in internal implementation of
Expand Down
20 changes: 20 additions & 0 deletions clang/lib/Sema/SemaSYCL.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -846,6 +846,26 @@ class SingleDeviceFunctionTracker {
llvm::SmallVector<FunctionDecl *> CallStack;
VisitCallNode(KernelNode, GetFDFromNode(KernelNode), CallStack);

// Mark kernel body function as such. Note a single function might have
// several kernel callers, e.g., when a range rounding kernel is added.
// Note the calling kernel name is already mangled.
if (KernelBody) {
llvm::SmallVector<llvm::StringRef, 2> KernelNames;
if (auto *Attr = KernelBody->getAttr<SYCLKernelObjFuncAttr>()) {
KernelNames.assign(Attr->kernels_begin(), Attr->kernels_end());
}
std::string KernelName = SYCLKernel->getNameAsString();
// Sorted insert to keep duplicates
auto *Pos = llvm::upper_bound(KernelNames, KernelName);
if (Pos == KernelNames.end() || *Pos != KernelName) {
KernelNames.insert(Pos, KernelName);
KernelBody->dropAttr<SYCLKernelObjFuncAttr>();
KernelBody->addAttr(SYCLKernelObjFuncAttr::CreateImplicit(
KernelBody->getASTContext(), KernelNames.data(),
KernelNames.size()));
}
}

// Always inline the KernelBody in the kernel entry point. For ESIMD
// inlining is handled later down the pipeline.
if (KernelBody &&
Expand Down
6 changes: 6 additions & 0 deletions clang/test/SemaSYCL/basic-kernel-wrapper.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -18,6 +18,12 @@ int main() {
});
}

// Check kernel body function has SYCLKernelObjFuncAttr attribute attached

// CHECK: CXXMethodDecl {{.*}} used constexpr operator() 'void () const'
// CHECK: SYCLKernelObjFuncAttr
// CHECK-SAME: Implicit _ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E14kernel_wrapper

// Check declaration of the kernel

// CHECK: FunctionDecl {{.*}}kernel_wrapper{{.*}} 'void (__global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
Expand Down