Skip to content

Commit 4a02954

Browse files
clementvalgithub-actions[bot]
authored andcommitted
Automerge: Revert "[flang][cuda][rt] Canonicalize block size values" (#164460)
Reverts llvm/llvm-project#164321 Align behavior with other CUDA Compiler
2 parents ece87d6 + 4a1ea3e commit 4a02954

File tree

1 file changed

+9
-9
lines changed

1 file changed

+9
-9
lines changed

flang-rt/lib/cuda/kernel.cpp

Lines changed: 9 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -23,9 +23,9 @@ void RTDEF(CUFLaunchKernel)(const void *kernel, intptr_t gridX, intptr_t gridY,
2323
gridDim.y = gridY;
2424
gridDim.z = gridZ;
2525
dim3 blockDim;
26-
blockDim.x = blockX > 1024 ? 1024 : blockX;
27-
blockDim.y = blockY > 1024 ? 1024 : blockY;
28-
blockDim.z = blockZ > 64 ? 64 : blockZ;
26+
blockDim.x = blockX;
27+
blockDim.y = blockY;
28+
blockDim.z = blockZ;
2929
unsigned nbNegGridDim{0};
3030
if (gridX < 0) {
3131
++nbNegGridDim;
@@ -88,9 +88,9 @@ void RTDEF(CUFLaunchClusterKernel)(const void *kernel, intptr_t clusterX,
8888
config.gridDim.x = gridX;
8989
config.gridDim.y = gridY;
9090
config.gridDim.z = gridZ;
91-
config.blockDim.x = blockX > 1024 ? 1024 : blockX;
92-
config.blockDim.y = blockY > 1024 ? 1024 : blockY;
93-
config.blockDim.z = blockZ > 64 ? 64 : blockZ;
91+
config.blockDim.x = blockX;
92+
config.blockDim.y = blockY;
93+
config.blockDim.z = blockZ;
9494
unsigned nbNegGridDim{0};
9595
if (gridX < 0) {
9696
++nbNegGridDim;
@@ -165,9 +165,9 @@ void RTDEF(CUFLaunchCooperativeKernel)(const void *kernel, intptr_t gridX,
165165
gridDim.y = gridY;
166166
gridDim.z = gridZ;
167167
dim3 blockDim;
168-
blockDim.x = blockX > 1024 ? 1024 : blockX;
169-
blockDim.y = blockY > 1024 ? 1024 : blockY;
170-
blockDim.z = blockZ > 64 ? 64 : blockZ;
168+
blockDim.x = blockX;
169+
blockDim.y = blockY;
170+
blockDim.z = blockZ;
171171
unsigned nbNegGridDim{0};
172172
if (gridX < 0) {
173173
++nbNegGridDim;

0 commit comments

Comments
 (0)