================
@@ -367,6 +370,67 @@ void
SemaSYCL::CheckSYCLEntryPointFunctionDecl(FunctionDecl *FD) {
namespace {
+CompoundStmt *BuildSYCLKernelLaunchStmt(Sema &SemaRef, FunctionDecl *FD,
+ const std::string &KernelName) {
+ ASTContext &Ctx = SemaRef.getASTContext();
+ SmallVector<Stmt *> Stmts;
+
+ // Prepare a string literal that contains the kernel name in the ordinary
+ // literal encoding.
+ // FIXME: transcode the contents of KernelName from UTF-8 to the
+ // ordinary literal encoding.
+ QualType KernelNameCharTy = Ctx.CharTy.withConst();
+ llvm::APInt KernelNameSize(Ctx.getTypeSize(Ctx.getSizeType()),
+ KernelName.size() + 1);
+ QualType KernelNameArrayTy = Ctx.getConstantArrayType(
+ KernelNameCharTy, KernelNameSize, nullptr, ArraySizeModifier::Normal, 0);
+ StringLiteral *KernelNameExpr = StringLiteral::Create(
+ Ctx, KernelName, StringLiteralKind::Ordinary,
+ /*Pascal*/ false, KernelNameArrayTy, SourceLocation());
+
+ // FIXME: An extern variable declaration with assignment to the kernel
+ // name expression is added to Stmts as a temporary measure to see results.
+ // reflected in tests. The kernel name expression will need to be passed as
+ // the first function argument in a call to sycl_enqueue_kernel_launch.
+ QualType ExternVarType = Ctx.getPointerType(Ctx.CharTy.withConst());
+ const IdentifierInfo *ExternVarName =
+ SemaRef.getPreprocessor().getIdentifierInfo("kernel_name");
+ VarDecl *ExternVarDecl = VarDecl::Create(
+ Ctx, FD, SourceLocation(), SourceLocation(), ExternVarName,
ExternVarType,
+ /*TInfo*/ nullptr, SC_Extern);
+ DeclStmt *ExternVarDeclStmt = new (Ctx)
+ DeclStmt(DeclGroupRef(ExternVarDecl), SourceLocation(),
SourceLocation());
+ Stmts.push_back(ExternVarDeclStmt);
+ DeclRefExpr *ExternVarDeclRef = new (Ctx) DeclRefExpr(
+ Ctx, ExternVarDecl, /*RefersToEnclosingVariableOrCapture*/ false,
+ ExternVarType, VK_LValue, SourceLocation());
+ ImplicitCastExpr *KernelNameArrayDecayExpr = new (Ctx) ImplicitCastExpr(
+ ImplicitCastExpr::OnStack, ExternVarType, CK_ArrayToPointerDecay,
+ KernelNameExpr, VK_PRValue, FPOptionsOverride());
+ BinaryOperator *AssignmentExpr = BinaryOperator::Create(
+ Ctx, ExternVarDeclRef, KernelNameArrayDecayExpr, BO_Assign,
ExternVarType,
+ VK_LValue, OK_Ordinary, SourceLocation(), FPOptionsOverride());
+ Stmts.push_back(AssignmentExpr);
+
+ // Perform overload resolution for a call to an accessible (member) function
+ // template named 'sycl_enqueue_kernel_launch' from within the definition of
+ // FD where:
+ // - The kernel name type is passed as the first template argument.
+ // - Any remaining template parameters are deduced from the function
arguments
+ // or assigned by default template arguments.
+ // - 'this' is passed as the implicit function argument if 'FD' is a
+ // non-static member function.
+ // - The name of the kernel, expressed as a string literal, is passed as the
+ // first function argument.
+ // - The parameters of FD are forwarded as-if by 'std::forward()' as the
+ // remaining explicit function arguments.
+ // - Any remaining function arguments are initialized by default arguments.
----------------
tahonermann wrote:
Agreed. We'll need to ensure the diagnostic issued when overload resolution
fails is clear. I hope we can do better than what Clang does for CUDA when the
`__cudaPushCallConfiguration()` declaration is missing:
```
t.cu:6:9: error: use of undeclared identifier __cudaPushCallConfiguration
6 | kernel<<<1,1>>>([] __attribute__((device)) {});
| ^
```
https://github.com/llvm/llvm-project/pull/152403
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
