[flang][cuda] Do not fail if global is not found

[clementval]

No description provided.

[llvmbot]

@llvm/pr-subscribers-flang-codegen

Author: Valentin Clement (バレンタイン クレメン) (clementval)

Changes

---

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

1 Files Affected:

• (modified) flang/lib/Optimizer/CodeGen/CodeGen.cpp (+7-5)

diff --git a/flang/lib/Optimizer/CodeGen/CodeGen.cpp b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
index 3c3804f35bf30..70bb43a2510ba 100644
--- a/flang/lib/Optimizer/CodeGen/CodeGen.cpp
+++ b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
@@ -179,11 +179,13 @@ struct AddrOfOpConversion : public fir::FIROpConversion<fir::AddrOfOp> {
 
     if (auto gpuMod = addr->getParentOfType<mlir::gpu::GPUModuleOp>()) {
       auto global = gpuMod.lookupSymbol<mlir::LLVM::GlobalOp>(addr.getSymbol());
-      assert(global && "Expect global in gpu module");
-      replaceWithAddrOfOrASCast(rewriter, addr->getLoc(), global.getAddrSpace(),
-                                getProgramAddressSpace(rewriter),
-                                global.getSymName(),
-                                convertType(addr.getType()), addr);
+      replaceWithAddrOfOrASCast(
+          rewriter, addr->getLoc(),
+          global ? global.getAddrSpace() : getGlobalAddressSpace(rewriter),
+          getProgramAddressSpace(rewriter),
+          global ? global.getSymName()
+                 : addr.getSymbol().getRootReference().getValue(),
+          convertType(addr.getType()), addr);
       return mlir::success();
     }
 

[llvmbot]

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

Author: Valentin Clement (バレンタイン クレメン) (clementval)

Changes

---

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

1 Files Affected:

• (modified) flang/lib/Optimizer/CodeGen/CodeGen.cpp (+7-5)

diff --git a/flang/lib/Optimizer/CodeGen/CodeGen.cpp b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
index 3c3804f35bf30..70bb43a2510ba 100644
--- a/flang/lib/Optimizer/CodeGen/CodeGen.cpp
+++ b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
@@ -179,11 +179,13 @@ struct AddrOfOpConversion : public fir::FIROpConversion<fir::AddrOfOp> {
 
     if (auto gpuMod = addr->getParentOfType<mlir::gpu::GPUModuleOp>()) {
       auto global = gpuMod.lookupSymbol<mlir::LLVM::GlobalOp>(addr.getSymbol());
-      assert(global && "Expect global in gpu module");
-      replaceWithAddrOfOrASCast(rewriter, addr->getLoc(), global.getAddrSpace(),
-                                getProgramAddressSpace(rewriter),
-                                global.getSymName(),
-                                convertType(addr.getType()), addr);
+      replaceWithAddrOfOrASCast(
+          rewriter, addr->getLoc(),
+          global ? global.getAddrSpace() : getGlobalAddressSpace(rewriter),
+          getProgramAddressSpace(rewriter),
+          global ? global.getSymName()
+                 : addr.getSymbol().getRootReference().getValue(),
+          convertType(addr.getType()), addr);
       return mlir::success();
     }