[CIR][CUDA] Generate kernel calls (#1348)

Now we could generate calls to `__global__` functions.

Most work is already done in AST. It rewrites `fn<<<2, 2>>>()` to
something like `__cudaPushCallConfiguration(dim3(2, 1, 1), dim3(2, 1,
1), 0, nullptr)`, which returns a bool. We calls the device stub as a
normal function when the call returns true.

Author: AdUhTkJm

clang/lib/CIR/CodeGen/CIRGenCUDARuntime.cpp

@@ -169,3 +169,23 @@ void CIRGenCUDARuntime::emitDeviceStub(CIRGenFunction &cgf, cir::FuncOp fn,
   else
     emitDeviceStubBodyLegacy(cgf, fn, args);
 }
+
+RValue CIRGenCUDARuntime::emitCUDAKernelCallExpr(CIRGenFunction &cgf,
+                                                 const CUDAKernelCallExpr *expr,
+                                                 ReturnValueSlot retValue) {
+  auto builder = cgm.getBuilder();
+  mlir::Location loc =
+      cgf.currSrcLoc ? cgf.currSrcLoc.value() : builder.getUnknownLoc();
+
+  cgf.emitIfOnBoolExpr(
+      expr->getConfig(),
+      [&](mlir::OpBuilder &b, mlir::Location l) {
+        CIRGenCallee callee = cgf.emitCallee(expr->getCallee());
+        cgf.emitCall(expr->getCallee()->getType(), callee, expr, retValue);
+        b.create<cir::YieldOp>(loc);
+      },
+      loc, [](mlir::OpBuilder &b, mlir::Location l) {},
+      std::optional<mlir::Location>());
+
+  return RValue::get(nullptr);
+}

clang/lib/CIR/CodeGen/CIRGenCUDARuntime.h

@@ -23,6 +23,8 @@ namespace clang::CIRGen {
 class CIRGenFunction;
 class CIRGenModule;
 class FunctionArgList;
+class RValue;
+class ReturnValueSlot;
 
 class CIRGenCUDARuntime {
 protected:
@@ -40,6 +42,10 @@ class CIRGenCUDARuntime {
 
   virtual void emitDeviceStub(CIRGenFunction &cgf, cir::FuncOp fn,
                               FunctionArgList &args);
+
+  virtual RValue emitCUDAKernelCallExpr(CIRGenFunction &cgf,
+                                        const CUDAKernelCallExpr *expr,
+                                        ReturnValueSlot retValue);
 };
 
 } // namespace clang::CIRGen

clang/lib/CIR/CodeGen/CIRGenExpr.cpp

@@ -530,7 +530,10 @@ static CIRGenCallee emitDirectCallee(CIRGenModule &CGM, GlobalDecl GD) {
 
   auto CalleePtr = emitFunctionDeclPointer(CGM, GD);
 
-  assert(!CGM.getLangOpts().CUDA && "NYI");
+  // For HIP, the device stub should be converted to handle.
+  if (CGM.getLangOpts().HIP && !CGM.getLangOpts().CUDAIsDevice &&
+      FD->hasAttr<CUDAGlobalAttr>())
+    llvm_unreachable("NYI");
 
   return CIRGenCallee::forDirect(CalleePtr, GD);
 }
@@ -1405,7 +1408,9 @@ RValue CIRGenFunction::emitCallExpr(const clang::CallExpr *E,
   if (const auto *CE = dyn_cast<CXXMemberCallExpr>(E))
     return emitCXXMemberCallExpr(CE, ReturnValue);
 
-  assert(!dyn_cast<CUDAKernelCallExpr>(E) && "CUDA NYI");
+  if (const auto *CE = dyn_cast<CUDAKernelCallExpr>(E))
+    return CGM.getCUDARuntime().emitCUDAKernelCallExpr(*this, CE, ReturnValue);
+
   if (const auto *CE = dyn_cast<CXXOperatorCallExpr>(E))
     if (const CXXMethodDecl *MD =
             dyn_cast_or_null<CXXMethodDecl>(CE->getCalleeDecl()))

clang/test/CIR/CodeGen/CUDA/simple.cu

@@ -31,3 +31,18 @@ __global__ void global_fn(int a) {}
 // CIR-HOST: cir.call @__cudaPopCallConfiguration
 // CIR-HOST: cir.get_global @_Z24__device_stub__global_fni
 // CIR-HOST: cir.call @cudaLaunchKernel
+
+int main() {
+  global_fn<<<1, 1>>>(1);
+}
+// CIR-DEVICE-NOT: cir.func @main()
+
+// CIR-HOST: cir.func @main()
+// CIR-HOST: cir.call @_ZN4dim3C1Ejjj
+// CIR-HOST: cir.call @_ZN4dim3C1Ejjj
+// CIR-HOST: [[Push:%[0-9]+]] = cir.call @__cudaPushCallConfiguration
+// CIR-HOST: [[ConfigOK:%[0-9]+]] = cir.cast(int_to_bool, [[Push]]
+// CIR-HOST: cir.if [[ConfigOK]] {
+// CIR-HOST:   [[Arg:%[0-9]+]] = cir.const #cir.int<1>
+// CIR-HOST:   cir.call @_Z24__device_stub__global_fni([[Arg]])
+// CIR-HOST: }