[CINN] Explicitly set launch bounds for grid reduce (#72191)

lshpku · web-flow · commit 207e34aa1648 · 2025-04-11T19:26:42.000+08:00
diff --git a/paddle/cinn/backends/codegen_gpu_dev.cc b/paddle/cinn/backends/codegen_gpu_dev.cc
@@ -244,6 +244,15 @@ void CodeGenGpuDev::PrintFunctionDeclaration(const ir::_LoweredFunc_ *op) {
     if (!has_symbol_in_thread_num) {
       str_ += "__launch_bounds__(";
       str_ += std::to_string(thread_num);
+      // Explicitly set min_blocks_per_sm for grid reduce to prevent launch
+      // failure.
+      if (!op->temp_spaces.empty()) {
+        int min_blocks_per_sm = 1024 / thread_num;
+        if (min_blocks_per_sm > 1) {
+          str_ += ", ";
+          str_ += std::to_string(min_blocks_per_sm);
+        }
+      }
       str_ += ") ";
     }
   }
