[flang-commits] [flang] Delete duplicated hlfir.declare op of induction variables of do concurrent when inside cuf kernel directive. (PR #134467)

[Zhen Wang via flang-commits]

https://github.com/wangzpgi created https://github.com/llvm/llvm-project/pull/134467

Delete duplicated creation of hlfir.declare op of do concurrent induction variables when inside cuf kernel directive. 
Obtain the correct hlfir.declare op generated from bindSymbol, and add it to ivValues.

>From 42755aeec6fc6d304ad98715179e48d4ba98a09e Mon Sep 17 00:00:00 2001
From: Zhen Wang <zhenw at nvidia.com>
Date: Fri, 4 Apr 2025 17:47:08 -0700
Subject: [PATCH] Delete duplicated hlfir.declare op of induction variables of
 do concurrent when inside cuf kernel directive.

---
 flang/lib/Lower/Bridge.cpp            |  8 +++-----
 flang/test/Lower/CUDA/cuda-doconc.cuf | 11 +++++++----
 2 files changed, 10 insertions(+), 9 deletions(-)

diff --git a/flang/lib/Lower/Bridge.cpp b/flang/lib/Lower/Bridge.cpp
index 65edf1cea8761..b4d1197822a43 100644
--- a/flang/lib/Lower/Bridge.cpp
+++ b/flang/lib/Lower/Bridge.cpp
@@ -3209,13 +3209,11 @@ class FirConverter : public Fortran::lower::AbstractConverter {
           builder->restoreInsertionPoint(insPt);
         }
 
-        // Create the hlfir.declare operation using the symbol's name
-        auto declareOp = builder->create<hlfir::DeclareOp>(
-            loc, ivValue, toStringRef(name.symbol->name()));
-        ivValue = declareOp.getResult(0);
-
         // Bind the symbol to the declared variable
         bindSymbol(*name.symbol, ivValue);
+        Fortran::lower::SymbolBox hsb = localSymbols.lookupSymbol(*name.symbol);
+        fir::ExtendedValue extIvValue = symBoxToExtendedValue(hsb);
+        ivValue = fir::getBase(extIvValue);
         ivValues.push_back(ivValue);
         ivTypes.push_back(idxTy);
         ivLocs.push_back(loc);
diff --git a/flang/test/Lower/CUDA/cuda-doconc.cuf b/flang/test/Lower/CUDA/cuda-doconc.cuf
index e240b1adc206a..ebdc0de91cf38 100644
--- a/flang/test/Lower/CUDA/cuda-doconc.cuf
+++ b/flang/test/Lower/CUDA/cuda-doconc.cuf
@@ -15,8 +15,9 @@ subroutine doconc1
 end
 
 ! CHECK: func.func @_QPdoconc1() {
-! CHECK: %[[DECL:.*]]:2 = hlfir.declare %{{.*}}#0 {uniq_name = "_QFdoconc1Ei"} : (!fir.ref<index>) -> (!fir.ref<index>, !fir.ref<index>)
-! CHECK: cuf.kernel<<<*, *>>>
+! CHECK: %[[DECL:.*]]:2 = hlfir.declare %{{.*}} {uniq_name = "_QFdoconc1Ei"} : (!fir.ref<index>) -> (!fir.ref<index>, !fir.ref<index>)
+! CHECK: cuf.kernel<<<*, *>>> (%arg0 : index)
+! CHECK: fir.store %arg0 to %[[DECL]]#0 : !fir.ref<index>
 ! CHECK: %{{.*}} = fir.load %[[DECL]]#0 : !fir.ref<index>
 
 subroutine doconc2
@@ -32,8 +33,10 @@ subroutine doconc2
 end
 
 ! CHECK: func.func @_QPdoconc2() {
-! CHECK: %[[DECLI:.*]]:2 = hlfir.declare %{{.*}}#0 {uniq_name = "_QFdoconc2Ei"} : (!fir.ref<index>) -> (!fir.ref<index>, !fir.ref<index>)
-! CHECK: %[[DECLJ:.*]]:2 = hlfir.declare %{{.*}}#0 {uniq_name = "_QFdoconc2Ej"} : (!fir.ref<index>) -> (!fir.ref<index>, !fir.ref<index>)
+! CHECK: %[[DECLI:.*]]:2 = hlfir.declare %{{.*}} {uniq_name = "_QFdoconc2Ei"} : (!fir.ref<index>) -> (!fir.ref<index>, !fir.ref<index>)
+! CHECK: %[[DECLJ:.*]]:2 = hlfir.declare %{{.*}} {uniq_name = "_QFdoconc2Ej"} : (!fir.ref<index>) -> (!fir.ref<index>, !fir.ref<index>)
 ! CHECK: cuf.kernel<<<*, *>>> (%arg0 : index, %arg1 : index) = (%{{.*}}, %{{.*}} : index, index) to (%{{.*}}, %{{.*}} : index, index)  step (%{{.*}}, %{{.*}} : index, index) {
+! CHECK: fir.store %arg0 to %[[DECLI]]#0 : !fir.ref<index>
+! CHECK: fir.store %arg1 to %[[DECLJ]]#0 : !fir.ref<index>
 ! CHECK: %{{.*}} = fir.load %[[DECLI]]#0 : !fir.ref<index>
 ! CHECK: %{{.*}} = fir.load %[[DECLJ]]#0 : !fir.ref<index>