[llvm] [flang][cuda][rt] Do not check error on kernel launch (PR #164463)
    Valentin Clement バレンタイン クレメン via llvm-commits 
    llvm-commits at lists.llvm.org
       
    Tue Oct 21 10:26:37 PDT 2025
    
    
  
https://github.com/clementval created https://github.com/llvm/llvm-project/pull/164463
Align behavior with other CUDA compiler. 
>From 4e98dce3d3183844bc8cbaf4b3a4cd6690440bf2 Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Tue, 21 Oct 2025 10:25:31 -0700
Subject: [PATCH] [flang][cuda][rt] Do not check error on kernel launch
---
 flang-rt/lib/cuda/kernel.cpp | 10 +++++-----
 1 file changed, 5 insertions(+), 5 deletions(-)
diff --git a/flang-rt/lib/cuda/kernel.cpp b/flang-rt/lib/cuda/kernel.cpp
index c52d039ce1075..8b9d3b87fe263 100644
--- a/flang-rt/lib/cuda/kernel.cpp
+++ b/flang-rt/lib/cuda/kernel.cpp
@@ -76,8 +76,8 @@ void RTDEF(CUFLaunchKernel)(const void *kernel, intptr_t gridX, intptr_t gridY,
     terminator.Crash("Too many invalid grid dimensions");
   }
   cudaStream_t defaultStream = 0;
-  CUDA_REPORT_IF_ERROR(cudaLaunchKernel(kernel, gridDim, blockDim, params, smem,
-      stream != nullptr ? (cudaStream_t)(*stream) : defaultStream));
+  cudaLaunchKernel(kernel, gridDim, blockDim, params, smem,
+      stream != nullptr ? (cudaStream_t)(*stream) : defaultStream);
 }
 
 void RTDEF(CUFLaunchClusterKernel)(const void *kernel, intptr_t clusterX,
@@ -153,7 +153,7 @@ void RTDEF(CUFLaunchClusterKernel)(const void *kernel, intptr_t clusterX,
   launchAttr[0].val.clusterDim.z = clusterZ;
   config.numAttrs = 1;
   config.attrs = launchAttr;
-  CUDA_REPORT_IF_ERROR(cudaLaunchKernelExC(&config, kernel, params));
+  cudaLaunchKernelExC(&config, kernel, params);
 }
 
 void RTDEF(CUFLaunchCooperativeKernel)(const void *kernel, intptr_t gridX,
@@ -218,8 +218,8 @@ void RTDEF(CUFLaunchCooperativeKernel)(const void *kernel, intptr_t gridX,
     terminator.Crash("Too many invalid grid dimensions");
   }
   cudaStream_t defaultStream = 0;
-  CUDA_REPORT_IF_ERROR(cudaLaunchCooperativeKernel(kernel, gridDim, blockDim,
-      params, smem, stream != nullptr ? (cudaStream_t)*stream : defaultStream));
+  cudaLaunchCooperativeKernel(kernel, gridDim, blockDim, params, smem,
+      stream != nullptr ? (cudaStream_t)*stream : defaultStream);
 }
 
 } // extern "C"
    
    
More information about the llvm-commits
mailing list