[flang-commits] [flang] c3a2251 - [flang][cuda] Fix managed address space mismatch on host side (#192304)

[via flang-commits]

Author: Valentin Clement (バレンタイン クレメン)
Date: 2026-04-15T18:27:15Z
New Revision: c3a22516888b94a4ad9919d0c224540fe1eb771c

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

LOG: [flang][cuda] Fix managed address space mismatch on host side (#192304)

Managed should not map to global

Added: 
    

Modified: 
    flang/lib/Optimizer/CodeGen/CodeGen.cpp
    flang/test/Fir/CUDA/cuda-code-gen.mlir

Removed: 
    


################################################################################
diff  --git a/flang/lib/Optimizer/CodeGen/CodeGen.cpp b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
index e0193019d9b2d..3837bd8aaa648 100644
--- a/flang/lib/Optimizer/CodeGen/CodeGen.cpp
+++ b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
@@ -205,8 +205,6 @@ static std::optional<unsigned> getCUFAddrSpace(fir::GlobalOp global) {
       return static_cast<unsigned>(mlir::NVVM::NVVMMemorySpace::Constant);
     if (*dataAttr == cuf::DataAttribute::Shared)
       return static_cast<unsigned>(mlir::NVVM::NVVMMemorySpace::Shared);
-    if (*dataAttr == cuf::DataAttribute::Managed)
-      return static_cast<unsigned>(mlir::NVVM::NVVMMemorySpace::Global);
   }
   return std::nullopt;
 }

diff  --git a/flang/test/Fir/CUDA/cuda-code-gen.mlir b/flang/test/Fir/CUDA/cuda-code-gen.mlir
index 6d2efda67b50d..ec066fad724c6 100644
--- a/flang/test/Fir/CUDA/cuda-code-gen.mlir
+++ b/flang/test/Fir/CUDA/cuda-code-gen.mlir
@@ -385,20 +385,17 @@ func.func @_QQhost_shared() {
 // Test that fir.address_of inside gpu.module referencing a managed fir.global
 // produces an addressof with ptr<1> and an addrspacecast.
 
-module attributes {gpu.container_module} {
-  gpu.module @cuda_device_mod {
-    fir.global @_QMmodEmval {data_attr = #cuf.cuda<managed>} : i32 {
-      %0 = fir.zero_bits i32
-      fir.has_value %0 : i32
-    }
-    gpu.func @_QMkernelsPuse_managed() kernel {
-      %0 = fir.address_of(@_QMmodEmval) : !fir.ref<i32>
-      gpu.return
-    }
+module {
+  func.func @_QMkernelsPuse_managed() {
+    %0 = fir.address_of(@_QMmodEmval) : !fir.ref<i32>
+    return
+  }
+  fir.global @_QMmodEmval {data_attr = #cuf.cuda<managed>} : i32 {
+    %0 = fir.zero_bits i32
+    fir.has_value %0 : i32
   }
 }
 
-// CHECK: llvm.mlir.global external @_QMmodEmval() {addr_space = 1 : i32, nvvm.managed} : i32
-// CHECK-LABEL: gpu.func @_QMkernelsPuse_managed()
-// CHECK: %[[ADDR:.*]] = llvm.mlir.addressof @_QMmodEmval : !llvm.ptr<1>
-// CHECK: %{{.*}} = llvm.addrspacecast %[[ADDR]] : !llvm.ptr<1> to !llvm.ptr
+// CHECK-LABEL: llvm.func @_QMkernelsPuse_managed()
+// CHECK: %{{.*}} = llvm.mlir.addressof @_QMmodEmval : !llvm.ptr
+// CHECK: llvm.mlir.global external @_QMmodEmval() {addr_space = 0 : i32} : i32