[llvm] Revert "[flang][cuda][rt] Canonicalize block size values" (PR #164460)

[Valentin Clement バレンタイン クレメン via llvm-commits]

https://github.com/clementval created https://github.com/llvm/llvm-project/pull/164460

Reverts llvm/llvm-project#164321

Align behavior with other CUDA Compiler

>From 0f49018c90e60276e2a6c5629a6164e823a639b2 Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Valentin=20Clement=20=28=E3=83=90=E3=83=AC=E3=83=B3?=
 =?UTF-8?q?=E3=82=BF=E3=82=A4=E3=83=B3=20=E3=82=AF=E3=83=AC=E3=83=A1?=
 =?UTF-8?q?=E3=83=B3=29?= <clementval at gmail.com>
Date: Tue, 21 Oct 2025 10:23:23 -0700
Subject: [PATCH] Revert "[flang][cuda][rt] Canonicalize block size values
 (#164321)"

This reverts commit 803883c6622685f342b546165ddce412cb057a8b.
---
 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 c52d039ce1075..e299a114ed7eb 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 > 1024 ? 1024 : blockX;
-  blockDim.y = blockY > 1024 ? 1024 : blockY;
-  blockDim.z = blockZ > 64 ? 64 : blockZ;
+  blockDim.x = blockX;
+  blockDim.y = blockY;
+  blockDim.z = 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 > 1024 ? 1024 : blockX;
-  config.blockDim.y = blockY > 1024 ? 1024 : blockY;
-  config.blockDim.z = blockZ > 64 ? 64 : blockZ;
+  config.blockDim.x = blockX;
+  config.blockDim.y = blockY;
+  config.blockDim.z = 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 > 1024 ? 1024 : blockX;
-  blockDim.y = blockY > 1024 ? 1024 : blockY;
-  blockDim.z = blockZ > 64 ? 64 : blockZ;
+  blockDim.x = blockX;
+  blockDim.y = blockY;
+  blockDim.z = blockZ;
   unsigned nbNegGridDim{0};
   if (gridX < 0) {
     ++nbNegGridDim;