[flang-commits] [flang] [flang][cuda] Lower kernel launch to fir.cuda_kernel_launch (PR #81891)

via flang-commits flang-commits at lists.llvm.org
Fri Feb 16 01:08:37 PST 2024


================
@@ -2436,6 +2436,65 @@ def fir_DispatchOp : fir_Op<"dispatch", []> {
   }];
 }
 
+def fir_CUDAKernelLaunch : fir_Op<"cuda_kernel_launch", [CallOpInterface,
+    AttrSizedOperandSegments]> {
+  let summary = "call CUDA kernel";
+
+  let description = [{
+    Launch a CUDA kernel from the host.
+
+    ```
+      // launch simple kernel with no arguments. bytes and stream value are
+      // optional in the chevron notation.
+      fir.cuda_kernel_launch @kernel<<<%gx, %gy, %bx, %by, %bz>>>()
----------------
jeanPerier wrote:

Nice that you could reproduce the chevron syntax in the MLIR format!

https://github.com/llvm/llvm-project/pull/81891


More information about the flang-commits mailing list