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

[via flang-commits]

Author: Zhen Wang
Date: 2025-04-06T19:31:09-07:00
New Revision: 8f0d8d28ccd8a1ced82a744679c5152f90e80c77

URL: https://github.com/llvm/llvm-project/commit/8f0d8d28ccd8a1ced82a744679c5152f90e80c77
DIFF: https://github.com/llvm/llvm-project/commit/8f0d8d28ccd8a1ced82a744679c5152f90e80c77.diff

LOG: Delete duplicated hlfir.declare op of induction variables of do concurrent when inside cuf kernel directive. (#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.

Added: 
    

Modified: 
    flang/lib/Lower/Bridge.cpp
    flang/test/Lower/CUDA/cuda-doconc.cuf

Removed: 
    


################################################################################
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>