Delete duplicated hlfir.declare op of induction variables of do concurrent when inside cuf kernel directive.

[wangzpgi]

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.

[llvmbot]

@llvm/pr-subscribers-flang-fir-hlfir

Author: Zhen Wang (wangzpgi)

Changes

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.

---

Full diff: https://github.com/llvm/llvm-project/pull/134467.diff

2 Files Affected:

• (modified) flang/lib/Lower/Bridge.cpp (+3-5)
• (modified) flang/test/Lower/CUDA/cuda-doconc.cuf (+7-4)

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>

[clementval]

LGTM