diff --git a/flang/lib/Optimizer/CodeGen/CodeGen.cpp b/flang/lib/Optimizer/CodeGen/CodeGen.cpp index 0afb295e58e54..3c3804f35bf30 100644 --- a/flang/lib/Optimizer/CodeGen/CodeGen.cpp +++ b/flang/lib/Optimizer/CodeGen/CodeGen.cpp @@ -176,6 +176,17 @@ struct AddrOfOpConversion : public fir::FIROpConversion { llvm::LogicalResult matchAndRewrite(fir::AddrOfOp addr, OpAdaptor adaptor, mlir::ConversionPatternRewriter &rewriter) const override { + + if (auto gpuMod = addr->getParentOfType()) { + auto global = gpuMod.lookupSymbol(addr.getSymbol()); + assert(global && "Expect global in gpu module"); + replaceWithAddrOfOrASCast(rewriter, addr->getLoc(), global.getAddrSpace(), + getProgramAddressSpace(rewriter), + global.getSymName(), + convertType(addr.getType()), addr); + return mlir::success(); + } + auto global = addr->getParentOfType() .lookupSymbol(addr.getSymbol()); replaceWithAddrOfOrASCast( @@ -3231,7 +3242,8 @@ struct GlobalOpConversion : public fir::FIROpConversion { if (global.getDataAttr() && *global.getDataAttr() == cuf::DataAttribute::Constant) - TODO(global.getLoc(), "CUDA Fortran CONSTANT variable code generation"); + g.setAddrSpace( + static_cast(mlir::NVVM::NVVMMemorySpace::Constant)); rewriter.eraseOp(global); return mlir::success(); diff --git a/flang/test/Fir/CUDA/cuda-code-gen.mlir b/flang/test/Fir/CUDA/cuda-code-gen.mlir index bbd3f9fbd351b..60cda9e98c7d8 100644 --- a/flang/test/Fir/CUDA/cuda-code-gen.mlir +++ b/flang/test/Fir/CUDA/cuda-code-gen.mlir @@ -284,3 +284,31 @@ module attributes {gpu.container_module, dlti.dl_spec = #dlti.dl_spec<#dlti.dl_e // CHECK-LABEL: llvm.func @_QQxxx() // CHECK: llvm.alloca %{{.*}} x !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<2 x array<3 x i64>>)> {alignment = 8 : i64} : (i32) -> !llvm.ptr // CHECK-NOT: llvm.call @_FortranACUFAllocDescriptor + +// ----- + +module attributes {gpu.container_module, dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry : vector<2xi64>>, #dlti.dl_entry : vector<2xi64>>, #dlti.dl_entry : vector<2xi64>>, #dlti.dl_entry, dense<64> : vector<4xi64>>, #dlti.dl_entry, dense<32> : vector<4xi64>>, #dlti.dl_entry, dense<32> : vector<4xi64>>, #dlti.dl_entry : vector<2xi64>>, #dlti.dl_entry : vector<2xi64>>, #dlti.dl_entry : vector<2xi64>>, #dlti.dl_entry : vector<2xi64>>, #dlti.dl_entry : vector<2xi64>>, #dlti.dl_entry : vector<2xi64>>, #dlti.dl_entry : vector<2xi64>>, #dlti.dl_entry : vector<4xi64>>, #dlti.dl_entry<"dlti.endianness", "little">, #dlti.dl_entry<"dlti.stack_alignment", 128 : i64>>} { + gpu.module @cuda_device_mod { + fir.global @_QMkernelsEinitial_val {data_attr = #cuf.cuda} : i32 { + %0 = fir.zero_bits i32 + fir.has_value %0 : i32 + } + gpu.func @_QMkernelsPassign(%arg0: !fir.ref>) kernel { + %c-1 = arith.constant -1 : index + %c1_i32 = arith.constant 1 : i32 + %0 = arith.constant 1 : i32 + %1 = arith.addi %0, %c1_i32 : i32 + %2 = fir.address_of(@_QMkernelsEinitial_val) : !fir.ref + %4 = fir.load %2 : !fir.ref + %5 = fir.convert %1 : (i32) -> i64 + %6 = fircg.ext_array_coor %arg0(%c-1)<%5> : (!fir.ref>, index, i64) -> !fir.ref + fir.store %4 to %6 : !fir.ref + gpu.return + } + } +} + +// CHECK: llvm.mlir.global external @_QMkernelsEinitial_val() {addr_space = 4 : i32} : i32 +// CHECK-LABEL: gpu.func @_QMkernelsPassign +// CHECK: %[[ADDROF:.*]] = llvm.mlir.addressof @_QMkernelsEinitial_val : !llvm.ptr<4> +// CHECK: %{{.*}} = llvm.addrspacecast %[[ADDROF]] : !llvm.ptr<4> to !llvm.ptr