[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