From 0644ad28b0d6b48ded929db19e6985a17f5109ac Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval@gmail.com>
Date: Mon, 20 Oct 2025 14:08:01 -0700
Subject: [PATCH] [flang][cuda][rt] Canonicalize block size values

---
 flang-rt/lib/cuda/kernel.cpp | 18 +++++++++---------
 1 file changed, 9 insertions(+), 9 deletions(-)

diff --git a/flang-rt/lib/cuda/kernel.cpp b/flang-rt/lib/cuda/kernel.cpp
index e299a114ed7eb..c52d039ce1075 100644
--- a/flang-rt/lib/cuda/kernel.cpp
+++ b/flang-rt/lib/cuda/kernel.cpp
@@ -23,9 +23,9 @@ void RTDEF(CUFLaunchKernel)(const void *kernel, intptr_t gridX, intptr_t gridY,
   gridDim.y = gridY;
   gridDim.z = gridZ;
   dim3 blockDim;
-  blockDim.x = blockX;
-  blockDim.y = blockY;
-  blockDim.z = blockZ;
+  blockDim.x = blockX > 1024 ? 1024 : blockX;
+  blockDim.y = blockY > 1024 ? 1024 : blockY;
+  blockDim.z = blockZ > 64 ? 64 : blockZ;
   unsigned nbNegGridDim{0};
   if (gridX < 0) {
     ++nbNegGridDim;
@@ -88,9 +88,9 @@ void RTDEF(CUFLaunchClusterKernel)(const void *kernel, intptr_t clusterX,
   config.gridDim.x = gridX;
   config.gridDim.y = gridY;
   config.gridDim.z = gridZ;
-  config.blockDim.x = blockX;
-  config.blockDim.y = blockY;
-  config.blockDim.z = blockZ;
+  config.blockDim.x = blockX > 1024 ? 1024 : blockX;
+  config.blockDim.y = blockY > 1024 ? 1024 : blockY;
+  config.blockDim.z = blockZ > 64 ? 64 : blockZ;
   unsigned nbNegGridDim{0};
   if (gridX < 0) {
     ++nbNegGridDim;
@@ -165,9 +165,9 @@ void RTDEF(CUFLaunchCooperativeKernel)(const void *kernel, intptr_t gridX,
   gridDim.y = gridY;
   gridDim.z = gridZ;
   dim3 blockDim;
-  blockDim.x = blockX;
-  blockDim.y = blockY;
-  blockDim.z = blockZ;
+  blockDim.x = blockX > 1024 ? 1024 : blockX;
+  blockDim.y = blockY > 1024 ? 1024 : blockY;
+  blockDim.z = blockZ > 64 ? 64 : blockZ;
   unsigned nbNegGridDim{0};
   if (gridX < 0) {
     ++nbNegGridDim;