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

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

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

Managed should not map to global

>From cf91e4f0bc60a120c2e1114dff7efa36435d806c Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Wed, 15 Apr 2026 11:12:37 -0700
Subject: [PATCH] [flang][cuda] Fix managed address space mismatch on host side

---
 flang/lib/Optimizer/CodeGen/CodeGen.cpp |  2 --
 flang/test/Fir/CUDA/cuda-code-gen.mlir  | 25 +++++++++++--------------
 2 files changed, 11 insertions(+), 16 deletions(-)

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