fix flang

Author: fabianmcg

flang/lib/Optimizer/CodeGen/CodeGen.cpp

@@ -3202,7 +3202,8 @@ struct GlobalOpConversion : public fir::FIROpConversion<fir::GlobalOp> {
 
     if (global.getDataAttr() &&
         *global.getDataAttr() == cuf::DataAttribute::Shared)
-      g.setAddrSpace(mlir::NVVM::NVVMMemorySpace::kSharedMemorySpace);
+      g.setAddrSpace(
+          static_cast<unsigned>(mlir::NVVM::NVVMMemorySpace::Shared));
 
     rewriter.eraseOp(global);
     return mlir::success();

flang/lib/Optimizer/Transforms/CUFGPUToLLVMConversion.cpp

@@ -221,7 +221,8 @@ static mlir::Value createAddressOfOp(mlir::ConversionPatternRewriter &rewriter,
                                      gpu::GPUModuleOp gpuMod,
                                      std::string &sharedGlobalName) {
   auto llvmPtrTy = mlir::LLVM::LLVMPointerType::get(
-      rewriter.getContext(), mlir::NVVM::NVVMMemorySpace::kSharedMemorySpace);
+      rewriter.getContext(),
+      static_cast<unsigned>(mlir::NVVM::NVVMMemorySpace::Shared));
   if (auto g = gpuMod.lookupSymbol<fir::GlobalOp>(sharedGlobalName))
     return mlir::LLVM::AddressOfOp::create(rewriter, loc, llvmPtrTy,
                                            g.getSymName());

mlir/test/Dialect/LLVMIR/nvvm.mlir

@@ -620,7 +620,7 @@ func.func @prefetch_tensormap(%gen_ptr: !llvm.ptr, %const_ptr: !llvm.ptr<4>) {
   return
 }
 
-// CHECK-LABEL: @nvvm_ptr_address_space
+// CHECK-LABEL: @nvvm_address_space
 func.func private @nvvm_address_space(
     !ptr.ptr<#nvvm.memory_space<global>>,
     !ptr.ptr<#nvvm.memory_space<shared>>,