[flang-commits] [flang] Revert "[flang][cuda] Set kernel intent(in) as const __restrict__" (PR #203734)

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

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

Reverts llvm/llvm-project#203652

breaking some downstream cudafor code

>From a68a70a544f532ba63154d1d749bb8db971ce84f 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: Sat, 13 Jun 2026 18:06:36 -0700
Subject: [PATCH] Revert "[flang][cuda] Set kernel intent(in) as const
 __restrict__ (#203652)"

This reverts commit 7670d88e7ea753095f753e6616b66f5e74da42ef.
---
 .../CUDA/CUFDeviceFuncTransform.cpp           | 38 -------------------
 .../Fir/CUDA/cuda-device-func-transform.mlir  | 18 +--------
 2 files changed, 2 insertions(+), 54 deletions(-)

diff --git a/flang/lib/Optimizer/Transforms/CUDA/CUFDeviceFuncTransform.cpp b/flang/lib/Optimizer/Transforms/CUDA/CUFDeviceFuncTransform.cpp
index d08940de0a739..88382cafcc3bd 100644
--- a/flang/lib/Optimizer/Transforms/CUDA/CUFDeviceFuncTransform.cpp
+++ b/flang/lib/Optimizer/Transforms/CUDA/CUFDeviceFuncTransform.cpp
@@ -41,41 +41,6 @@ class CUFDeviceFuncTransform
   using CUFDeviceFuncTransformBase<
       CUFDeviceFuncTransform>::CUFDeviceFuncTransformBase;
 
-  // Decorate INTENT(IN) kernel arguments like C "const __restrict__". NVPTX
-  // only tags loads as invariant (and lowers them to ld.global.nc) when a
-  // kernel pointer parameter is both readonly and noalias; see
-  // NVPTXTagInvariantLoads.
-  static void setIntentInKernelArgAttrs(mlir::func::FuncOp funcOp,
-                                        gpu::GPUFuncOp deviceFuncOp) {
-    mlir::UnitAttr unitAttr = mlir::UnitAttr::get(funcOp.getContext());
-
-    auto markArg = [&](unsigned argIndex) {
-      if (argIndex >= deviceFuncOp.getNumArguments())
-        return;
-      deviceFuncOp.setArgAttr(
-          argIndex, mlir::LLVM::LLVMDialect::getReadonlyAttrName(), unitAttr);
-      deviceFuncOp.setArgAttr(
-          argIndex, mlir::LLVM::LLVMDialect::getNoAliasAttrName(), unitAttr);
-    };
-
-    funcOp.walk([&](fir::DeclareOp declareOp) {
-      auto var =
-          mlir::cast<fir::FortranVariableOpInterface>(declareOp.getOperation());
-      if (!var.isIntentIn())
-        return;
-      if (std::optional<uint32_t> dummyArgNo = declareOp.getDummyArgNo()) {
-        // Dummy argument numbers are 1-based in FIR.
-        markArg(*dummyArgNo - 1);
-        return;
-      }
-      if (auto blockArg =
-              mlir::dyn_cast<mlir::BlockArgument>(declareOp.getMemref()))
-        if (blockArg.getOwner()->isEntryBlock() &&
-            blockArg.getOwner()->getParentOp() == funcOp)
-          markArg(blockArg.getArgNumber());
-    });
-  }
-
   static gpu::GPUFuncOp createGPUFuncOp(mlir::func::FuncOp funcOp,
                                         bool isGlobal, int computeCap) {
     mlir::OpBuilder builder(funcOp.getContext());
@@ -102,9 +67,6 @@ class CUFDeviceFuncTransform
     auto deviceFuncOp =
         gpu::GPUFuncOp::create(builder, loc, funcOp.getName(), type,
                                mlir::TypeRange{}, mlir::TypeRange{});
-    if (mlir::ArrayAttr argAttrs = funcOp.getAllArgAttrs())
-      deviceFuncOp.setAllArgAttrs(argAttrs);
-    setIntentInKernelArgAttrs(funcOp, deviceFuncOp);
     if (isGlobal)
       deviceFuncOp.setKernel(true);
 
diff --git a/flang/test/Fir/CUDA/cuda-device-func-transform.mlir b/flang/test/Fir/CUDA/cuda-device-func-transform.mlir
index ad15b93ef43ef..cc183f0144e82 100644
--- a/flang/test/Fir/CUDA/cuda-device-func-transform.mlir
+++ b/flang/test/Fir/CUDA/cuda-device-func-transform.mlir
@@ -34,7 +34,7 @@ func.func private @_QMmod1Psub1(!fir.ref<!fir.array<10xi32>> {cuf.data_attr = #c
 
 // CHECK: gpu.func @_QPsub_device1()
 
-// CHECK: gpu.func @_QPsub_device2(%[[ARG0:.*]]: !fir.ref<f32>
+// CHECK: gpu.func @_QPsub_device2(%[[ARG0:.*]]: !fir.ref<f32>) {
 // CHECK:   %[[DECL:.*]] = fir.declare %[[ARG0]] {uniq_name = "_QFsub1Ei"} : (!fir.ref<f32>) -> !fir.ref<f32>
 // CHECK:   %[[CST:.*]] = arith.constant 2.000000e+00 : f32
 // CHECK:   fir.store %[[CST]] to %[[DECL]] : !fir.ref<f32>
@@ -151,26 +151,12 @@ func.func @_QPpartialsumshflshflr8(%arg0: !fir.ref<!fir.array<?xf64>> {cuf.data_
 }
 
 // CHECK-LABEL: gpu.module @cuda_device_mod
-// CHECK: gpu.func @_QPpartialsumshflshflr8({{.*}}) kernel
+// CHECK: gpu.func @_QPpartialsumshflshflr8(%arg0: !fir.ref<!fir.array<?xf64>>, %arg1: i32) kernel
       
 // CHECK: func.func @_QPpartialsumshflshflr8
 
 // -----
 
-func.func @_QPldg_attrs(%arg0: !fir.ref<!fir.array<?xf32>> {fir.bindc_name = "a"}, %arg1: !fir.ref<!fir.array<?xf32>> {fir.bindc_name = "b"}) attributes {cuf.proc_attr = #cuf.cuda_proc<global>} {
-  %c10 = arith.constant 10 : index
-  %scope = fir.dummy_scope : !fir.dscope
-  %shape = fir.shape %c10 : (index) -> !fir.shape<1>
-  %0 = fir.declare %arg0(%shape) dummy_scope %scope arg 1 {fortran_attrs = #fir.var_attrs<intent_out>, uniq_name = "_QFldg_attrsEa"} : (!fir.ref<!fir.array<?xf32>>, !fir.shape<1>, !fir.dscope) -> !fir.ref<!fir.array<?xf32>>
-  %1 = fir.declare %arg1(%shape) dummy_scope %scope arg 2 {fortran_attrs = #fir.var_attrs<intent_in>, uniq_name = "_QFldg_attrsEb"} : (!fir.ref<!fir.array<?xf32>>, !fir.shape<1>, !fir.dscope) -> !fir.ref<!fir.array<?xf32>>
-  return
-}
-
-// CHECK-LABEL: gpu.module @cuda_device_mod
-// CHECK: gpu.func @_QPldg_attrs(%{{.*}}: !fir.ref<!fir.array<?xf32>>{{.*}}, %{{.*}}: !fir.ref<!fir.array<?xf32>> {{{.*}}llvm.noalias, llvm.readonly}) kernel
-
-// -----
-
 func.func @_QPsub_maxtnid() attributes {cuf.launch_bounds = #cuf.launch_bounds<maxTPB = 256 : i64, minBPM = 2 : i64, upperBoundClusterSize = 3 : i64>, cuf.proc_attr = #cuf.cuda_proc<global>} {
   %cst = arith.constant 2.000000e+00 : f32
   return