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?= 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;