llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: Arseniy Obolenskiy (aobolensk) <details> <summary>Changes</summary> Add it to `__hip_*` and `__opencl_*` atomicrmw and cmpxchg on a flat pointer, so the backend can skip the private memory check --- Full diff: https://github.com/llvm/llvm-project/pull/232350.diff 6 Files Affected: - (modified) clang/include/clang/CIR/Dialect/IR/CIRDialect.td (+1) - (modified) clang/include/clang/CIR/MissingFeatures.h (-1) - (modified) clang/lib/CIR/CodeGen/CIRGenAtomic.cpp (+38-36) - (modified) clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp (+14-1) - (modified) clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp (+16) - (modified) clang/test/CIR/CodeGenHIP/atomic-options.hip (+39-6) ``````````diff diff --git a/clang/include/clang/CIR/Dialect/IR/CIRDialect.td b/clang/include/clang/CIR/Dialect/IR/CIRDialect.td index 5fb3366d0c4e7..dff732ec40f74 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRDialect.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRDialect.td @@ -127,6 +127,7 @@ def CIR_Dialect : Dialect { static llvm::StringRef getAMDGPUNoRemoteMemoryAttrName() { return "cir.amdgpu_no_remote_memory"; } static llvm::StringRef getAMDGPUIgnoreDenormalModeAttrName() { return "cir.amdgpu_ignore_denormal_mode"; } static llvm::StringRef getAMDGPURawAtomicAttrName() { return "cir.amdgpu_raw_atomic"; } + static llvm::StringRef getAMDGPUNoPrivateAddrSpaceAttrName() { return "cir.amdgpu_no_private_addrspace"; } static llvm::StringRef getOpenCLKernelArgMetadataAttrName() { return "cir.cl.kernel_arg_metadata"; } static llvm::StringRef getMaxWorkGroupSizeAttrName() { return "cir.max_work_group_size"; } static llvm::StringRef getOpenCLVersionAttrName() { return "cir.cl.version"; } diff --git a/clang/include/clang/CIR/MissingFeatures.h b/clang/include/clang/CIR/MissingFeatures.h index cd1e43b17ec3f..709e4a8239136 100644 --- a/clang/include/clang/CIR/MissingFeatures.h +++ b/clang/include/clang/CIR/MissingFeatures.h @@ -155,7 +155,6 @@ struct MissingFeatures { static bool atomicUseLibCall() { return false; } static bool atomicMicrosoftVolatile() { return false; } static bool atomicOpenMP() { return false; } - static bool atomicAMDGPUNoaliasAddrspace() { return false; } static bool atomicAMDGPUAvailableVisibleMMRA() { return false; } // Global ctor handling diff --git a/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp b/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp index 095ebdaa04ce1..ccbb453b13673 100644 --- a/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp @@ -483,6 +483,41 @@ static void emitMemOrderCaseLabel(CIRGenBuilderTy &builder, mlir::Location loc, builder.restoreInsertionPoint(insertPoint); } +/// Attach the AMDGPU atomic metadata markers that the current atomic options +/// call for. The "no.X" metadata is emitted when the corresponding option is +/// off, since it asserts the absence of that memory kind. The clang::atomic +/// attribute is what turns the options on and off. +static void setAMDGPUAtomicMetadata(CIRGenFunction &cgf, mlir::Operation *op, + const AtomicExpr *expr) { + // TODO: AMDGPUTargetCodeGenInfo::setTargetAtomicMetadata also calls + // CGF.AddAMDGPUAvailableVisibleMMRA on every atomic instruction; this is + // tied to the AMDGPUAvailableVisible statement attribute, which + // CIRGenStmt.cpp's emitAttributedStmt currently rejects via errorNYI. + assert(!cir::MissingFeatures::atomicAMDGPUAvailableVisibleMMRA()); + + // Lowering keeps this only on a flat pointer. + mlir::UnitAttr unit = cgf.getBuilder().getUnitAttr(); + if (expr->threadPrivateMemoryAtomicsAreUndefined()) + op->setAttr(cir::CIRDialect::getAMDGPUNoPrivateAddrSpaceAttrName(), unit); + + // Only a read-modify-write instruction carries these; a plain load, store or + // cmpxchg does not. + auto fetchOp = mlir::dyn_cast<cir::AtomicFetchOp>(op); + if (!fetchOp) + return; + + clang::AtomicOptions atomicOpts = cgf.cgm.getAtomicOpts(); + if (!atomicOpts.getOption(clang::AtomicOptionKind::FineGrainedMemory)) + op->setAttr(cir::CIRDialect::getAMDGPUNoFineGrainedMemoryAttrName(), unit); + if (!atomicOpts.getOption(clang::AtomicOptionKind::RemoteMemory)) + op->setAttr(cir::CIRDialect::getAMDGPUNoRemoteMemoryAttrName(), unit); + // Denormal flushing only matters for a float add. + if (atomicOpts.getOption(clang::AtomicOptionKind::IgnoreDenormalMode) && + fetchOp.getBinop() == cir::AtomicFetchKind::Add && + mlir::isa<cir::SingleType>(fetchOp.getVal().getType())) + op->setAttr(cir::CIRDialect::getAMDGPUIgnoreDenormalModeAttrName(), unit); +} + static void emitAtomicCmpXchg(CIRGenFunction &cgf, AtomicExpr *e, bool isWeak, Address dest, Address ptr, Address val1, Address val2, Address expectedResult, @@ -505,6 +540,8 @@ static void emitAtomicCmpXchg(CIRGenFunction &cgf, AtomicExpr *e, bool isWeak, cmpxchg.setIsVolatile(e->isVolatile()); cmpxchg.setWeak(isWeak); + if (cgf.cgm.getTriple().isAMDGCN()) + setAMDGPUAtomicMetadata(cgf, cmpxchg, e); mlir::Value failed = builder.createNot(cmpxchg.getSuccess()); cir::IfOp::create( @@ -661,41 +698,6 @@ static void emitAtomicCmpXchgFailureSetCheckWeak( }); } -/// Attach the AMDGPU atomic metadata markers that the current atomic options -/// call for. The "no.X" metadata is emitted when the corresponding option is -/// off, since it asserts the absence of that memory kind. The clang::atomic -/// attribute is what turns the options on and off. -static void setAMDGPUAtomicMetadata(CIRGenFunction &cgf, mlir::Operation *op) { - // TODO: AMDGPUTargetCodeGenInfo::setTargetAtomicMetadata also emits - // !noalias.addrspace on a flat-pointer atomic when the source atomic - // expression's memory is thread-private-undefined (OpenCL / old-style HIP - // atomics), regardless of whether it is a read-modify-write or cmpxchg. - assert(!cir::MissingFeatures::atomicAMDGPUNoaliasAddrspace()); - // TODO: AMDGPUTargetCodeGenInfo::setTargetAtomicMetadata also calls - // CGF.AddAMDGPUAvailableVisibleMMRA on every atomic instruction; this is - // tied to the AMDGPUAvailableVisible statement attribute, which - // CIRGenStmt.cpp's emitAttributedStmt currently rejects via errorNYI. - assert(!cir::MissingFeatures::atomicAMDGPUAvailableVisibleMMRA()); - - // Only a read-modify-write instruction carries these; a plain load, store or - // cmpxchg does not. - auto fetchOp = mlir::dyn_cast<cir::AtomicFetchOp>(op); - if (!fetchOp) - return; - - clang::AtomicOptions atomicOpts = cgf.cgm.getAtomicOpts(); - mlir::UnitAttr unit = cgf.getBuilder().getUnitAttr(); - if (!atomicOpts.getOption(clang::AtomicOptionKind::FineGrainedMemory)) - op->setAttr(cir::CIRDialect::getAMDGPUNoFineGrainedMemoryAttrName(), unit); - if (!atomicOpts.getOption(clang::AtomicOptionKind::RemoteMemory)) - op->setAttr(cir::CIRDialect::getAMDGPUNoRemoteMemoryAttrName(), unit); - // Denormal flushing only matters for a float add. - if (atomicOpts.getOption(clang::AtomicOptionKind::IgnoreDenormalMode) && - fetchOp.getBinop() == cir::AtomicFetchKind::Add && - mlir::isa<cir::SingleType>(fetchOp.getVal().getType())) - op->setAttr(cir::CIRDialect::getAMDGPUIgnoreDenormalModeAttrName(), unit); -} - static void emitAtomicOp(CIRGenFunction &cgf, AtomicExpr *expr, Address dest, Address ptr, Address val1, Address val2, Address expectedResult, Expr *isWeakExpr, @@ -969,7 +971,7 @@ static void emitAtomicOp(CIRGenFunction &cgf, AtomicExpr *expr, Address dest, rmwOp->setAttr("fetch_first", builder.getUnitAttr()); if (cgf.cgm.getTriple().isAMDGCN()) - setAMDGPUAtomicMetadata(cgf, rmwOp); + setAMDGPUAtomicMetadata(cgf, rmwOp, expr); mlir::Value result = rmwOp->getResult(0); diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp index fb2be40281883..695deda3980b4 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -1411,6 +1411,16 @@ getLLVMSyncScope(std::optional<cir::SyncScopeKind> syncScope, return std::nullopt; } +static void forwardAMDGPUNoPrivateAddrSpace(mlir::Operation *op, + mlir::Operation *llvmOp, + mlir::Value llvmPtr) { + llvm::StringRef name = cir::CIRDialect::getAMDGPUNoPrivateAddrSpaceAttrName(); + auto ptrTy = mlir::cast<mlir::LLVM::LLVMPointerType>(llvmPtr.getType()); + if (mlir::Attribute a = op->getAttr(name); + a && ptrTy.getAddressSpace() == llvm::AMDGPUAS::FLAT_ADDRESS) + llvmOp->setAttr(name, a); +} + mlir::LogicalResult CIRToLLVMAtomicCmpXchgOpLowering::matchAndRewrite( cir::AtomicCmpXchgOp op, OpAdaptor adaptor, mlir::ConversionPatternRewriter &rewriter) const { @@ -1427,6 +1437,7 @@ mlir::LogicalResult CIRToLLVMAtomicCmpXchgOpLowering::matchAndRewrite( cmpxchg.setAlignment(adaptor.getAlignment()); cmpxchg.setWeak(adaptor.getWeak()); cmpxchg.setVolatile_(adaptor.getIsVolatile()); + forwardAMDGPUNoPrivateAddrSpace(op, cmpxchg, adaptor.getPtr()); // Check result and apply stores accordingly. auto old = mlir::LLVM::ExtractValueOp::create(rewriter, op.getLoc(), @@ -1445,9 +1456,10 @@ mlir::LogicalResult CIRToLLVMAtomicXchgOpLowering::matchAndRewrite( mlir::LLVM::AtomicOrdering llvmOrder = getLLVMMemOrder(adaptor.getMemOrder()); llvm::StringRef llvmSyncScope = getLLVMSyncScope( op.getSyncScope(), op->getParentOfType<mlir::ModuleOp>()); - rewriter.replaceOpWithNewOp<mlir::LLVM::AtomicRMWOp>( + auto xchg = rewriter.replaceOpWithNewOp<mlir::LLVM::AtomicRMWOp>( op, mlir::LLVM::AtomicBinOp::xchg, adaptor.getPtr(), adaptor.getVal(), llvmOrder, llvmSyncScope, /*alignment=*/0, op.getIsVolatile()); + forwardAMDGPUNoPrivateAddrSpace(op, xchg, adaptor.getPtr()); return mlir::success(); } @@ -1660,6 +1672,7 @@ mlir::LogicalResult CIRToLLVMAtomicFetchOpLowering::matchAndRewrite( cir::CIRDialect::getAMDGPUIgnoreDenormalModeAttrName()}) if (mlir::Attribute a = op->getAttr(marker)) rmwVal->setAttr(marker, a); + forwardAMDGPUNoPrivateAddrSpace(op, rmwVal, adaptor.getPtr()); // The AMDGPU raw hardware atomic builtins need metadata for the backend to // select the native instruction. LDS atomics are always native, so the diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp index 9fc12631d8ed5..0cd3f2a19747e 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp @@ -20,7 +20,9 @@ #include "llvm/IR/Constant.h" #include "llvm/IR/Constants.h" #include "llvm/IR/GlobalVariable.h" +#include "llvm/IR/MDBuilder.h" #include "llvm/IR/Metadata.h" +#include "llvm/Support/AMDGPUAddrSpace.h" using namespace llvm; @@ -75,6 +77,9 @@ class CIRDialectLLVMIRTranslationInterface attribute.getName() == cir::CIRDialect::getAMDGPUIgnoreDenormalModeAttrName()) { amendAMDGPUAtomicMetadata(instructions, attribute, moduleTranslation); + } else if (attribute.getName() == + cir::CIRDialect::getAMDGPUNoPrivateAddrSpaceAttrName()) { + amendAMDGPUNoPrivateAddrSpace(instructions, moduleTranslation); } return mlir::success(); } @@ -128,6 +133,17 @@ class CIRDialectLLVMIRTranslationInterface inst->setMetadata(mdName, empty); } + void amendAMDGPUNoPrivateAddrSpace( + llvm::ArrayRef<llvm::Instruction *> instructions, + mlir::LLVM::ModuleTranslation &moduleTranslation) const { + llvm::MDBuilder mdHelper(moduleTranslation.getLLVMContext()); + llvm::MDNode *asRange = mdHelper.createRange( + llvm::APInt(32, llvm::AMDGPUAS::PRIVATE_ADDRESS), + llvm::APInt(32, llvm::AMDGPUAS::PRIVATE_ADDRESS + 1)); + for (llvm::Instruction *inst : instructions) + inst->setMetadata(llvm::LLVMContext::MD_noalias_addrspace, asRange); + } + mlir::LogicalResult amendRISCVNontemporalDomain( mlir::Operation *op, llvm::ArrayRef<llvm::Instruction *> instructions, mlir::NamedAttribute attribute, diff --git a/clang/test/CIR/CodeGenHIP/atomic-options.hip b/clang/test/CIR/CodeGenHIP/atomic-options.hip index 7fa524f23de51..5dc137ed7cfc5 100644 --- a/clang/test/CIR/CodeGenHIP/atomic-options.hip +++ b/clang/test/CIR/CodeGenHIP/atomic-options.hip @@ -10,7 +10,7 @@ #define __global__ __attribute__((global)) // CIR-LABEL: @_Z9plain_addPff -// CIR: cir.atomic.fetch add relaxed syncscope(hip_agent) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.float>, !cir.float) -> !cir.float {cir.amdgpu_no_fine_grained_memory, cir.amdgpu_no_remote_memory} +// CIR: cir.atomic.fetch add relaxed syncscope(hip_agent) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.float>, !cir.float) -> !cir.float {cir.amdgpu_no_fine_grained_memory, cir.amdgpu_no_private_addrspace, cir.amdgpu_no_remote_memory} // LLVM-LABEL: @_Z9plain_addPff // LLVM: atomicrmw fadd ptr %{{.+}}, float %{{.+}} syncscope("agent") monotonic, align 4{{.*}}!amdgpu.{{(no.remote.memory|no.fine.grained.memory)}} !{{[0-9]+}}{{.*}}!amdgpu.{{(no.remote.memory|no.fine.grained.memory)}} __device__ float plain_add(float *p, float v) { @@ -18,7 +18,7 @@ __device__ float plain_add(float *p, float v) { } // CIR-LABEL: @_Z10unsafe_addPff -// CIR: cir.atomic.fetch add relaxed syncscope(hip_agent) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.float>, !cir.float) -> !cir.float {cir.amdgpu_no_remote_memory} +// CIR: cir.atomic.fetch add relaxed syncscope(hip_agent) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.float>, !cir.float) -> !cir.float {cir.amdgpu_no_private_addrspace, cir.amdgpu_no_remote_memory} // LLVM-LABEL: @_Z10unsafe_addPff // LLVM: atomicrmw fadd ptr %{{.+}}, float %{{.+}} syncscope("agent") monotonic // LLVM-NOT: amdgpu.no.fine.grained.memory @@ -29,7 +29,7 @@ __device__ float unsafe_add(float *p, float v) { } // CIR-LABEL: @_Z9no_fg_addPff -// CIR: cir.atomic.fetch add relaxed syncscope(hip_agent) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.float>, !cir.float) -> !cir.float {cir.amdgpu_ignore_denormal_mode, cir.amdgpu_no_fine_grained_memory, cir.amdgpu_no_remote_memory} +// CIR: cir.atomic.fetch add relaxed syncscope(hip_agent) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.float>, !cir.float) -> !cir.float {cir.amdgpu_ignore_denormal_mode, cir.amdgpu_no_fine_grained_memory, cir.amdgpu_no_private_addrspace, cir.amdgpu_no_remote_memory} // LLVM-LABEL: @_Z9no_fg_addPff // LLVM: atomicrmw fadd ptr %{{.+}}, float %{{.+}} syncscope("agent") monotonic, align 4{{.*}}!atomic.ignore.denormal.mode !{{[0-9]+}}, !amdgpu.no.fine.grained.memory !{{[0-9]+}}, !amdgpu.no.remote.memory __device__ float no_fg_add(float *p, float v) { @@ -39,7 +39,7 @@ __device__ float no_fg_add(float *p, float v) { } // CIR-LABEL: @_Z9no_fg_addPdd -// CIR: cir.atomic.fetch add relaxed syncscope(hip_agent) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.double>, !cir.double) -> !cir.double {cir.amdgpu_no_fine_grained_memory, cir.amdgpu_no_remote_memory} +// CIR: cir.atomic.fetch add relaxed syncscope(hip_agent) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.double>, !cir.double) -> !cir.double {cir.amdgpu_no_fine_grained_memory, cir.amdgpu_no_private_addrspace, cir.amdgpu_no_remote_memory} // LLVM-LABEL: @_Z9no_fg_addPdd // LLVM: atomicrmw fadd ptr %{{.+}}, double %{{.+}} syncscope("agent") monotonic // LLVM-NOT: atomic.ignore.denormal.mode @@ -50,7 +50,7 @@ __device__ double no_fg_add(double *p, double v) { } // CIR-LABEL: @_Z12workgroup_ddPdd -// CIR: cir.atomic.fetch add relaxed syncscope(hip_workgroup) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.double>, !cir.double) -> !cir.double {cir.amdgpu_no_fine_grained_memory, cir.amdgpu_no_remote_memory} +// CIR: cir.atomic.fetch add relaxed syncscope(hip_workgroup) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.double>, !cir.double) -> !cir.double {cir.amdgpu_no_fine_grained_memory, cir.amdgpu_no_private_addrspace, cir.amdgpu_no_remote_memory} // LLVM-LABEL: @_Z12workgroup_ddPdd // LLVM: atomicrmw fadd ptr %{{.+}}, double %{{.+}} syncscope("workgroup") monotonic __device__ double workgroup_dd(double *p, double v) { @@ -58,10 +58,43 @@ __device__ double workgroup_dd(double *p, double v) { } // CIR-LABEL: @_Z12after_scopedPff -// CIR: cir.atomic.fetch add relaxed syncscope(hip_agent) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.float>, !cir.float) -> !cir.float {cir.amdgpu_no_fine_grained_memory, cir.amdgpu_no_remote_memory} +// CIR: cir.atomic.fetch add relaxed syncscope(hip_agent) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.float>, !cir.float) -> !cir.float {cir.amdgpu_no_fine_grained_memory, cir.amdgpu_no_private_addrspace, cir.amdgpu_no_remote_memory} // LLVM-LABEL: @_Z12after_scopedPff // LLVM: atomicrmw fadd ptr %{{.+}}, float %{{.+}} syncscope("agent") monotonic, align 4{{.*}}!amdgpu.no.fine.grained.memory __device__ float after_scoped(float *p, float v) { [[clang::atomic(fine_grained_memory)]] { (void)0; } return __hip_atomic_fetch_add(p, v, __ATOMIC_RELAXED, 4); } + +// CIR-LABEL: @_Z4xchgPii +// CIR: cir.atomic.xchg relaxed syncscope(hip_agent) %{{.+}}, %{{.+}} : (!cir.ptr<!s32i>, !s32i) -> !s32i {cir.amdgpu_no_private_addrspace} +// LLVM-LABEL: @_Z4xchgPii +// LLVM: atomicrmw xchg ptr %{{.+}}, i32 %{{.+}} syncscope("agent") monotonic, align 4, !noalias.addrspace ![[#NOPRIV:]] +__device__ int xchg(int *p, int v) { + return __hip_atomic_exchange(p, v, __ATOMIC_RELAXED, 4); +} + +// CIR-LABEL: @_Z7cmpxchgPiii +// CIR: cir.atomic.cmpxchg success(relaxed) failure(relaxed) syncscope(hip_agent) %{{.+}}, %{{.+}}, %{{.+}} align(4) : (!cir.ptr<!s32i>, !s32i, !s32i) -> (!s32i, !cir.bool) {cir.amdgpu_no_private_addrspace} +// LLVM-LABEL: @_Z7cmpxchgPiii +// LLVM: cmpxchg ptr %{{.+}}, i32 %{{.+}}, i32 %{{.+}} syncscope("agent") monotonic monotonic, align 4, !noalias.addrspace ![[#NOPRIV]]{{$}} +__device__ bool cmpxchg(int *p, int e, int d) { + return __hip_atomic_compare_exchange_strong(p, &e, d, __ATOMIC_RELAXED, + __ATOMIC_RELAXED, 4); +} + +// LLVM-LABEL: @_Z10global_addPU3AS1ii +// LLVM: atomicrmw add ptr addrspace(1) %{{.+}}, i32 %{{.+}} syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory +__device__ int global_add(__attribute__((address_space(1))) int *p, int v) { + return __hip_atomic_fetch_add(p, v, __ATOMIC_RELAXED, 4); +} + +// CIR-LABEL: @_Z7gnu_addPii +// CIR: cir.atomic.fetch add relaxed syncscope(system) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!s32i>, !s32i) -> !s32i {cir.amdgpu_no_fine_grained_memory, cir.amdgpu_no_remote_memory} +// LLVM-LABEL: @_Z7gnu_addPii +// LLVM: atomicrmw add ptr %{{.+}}, i32 %{{.+}} monotonic, align 4, !amdgpu.no.fine.grained.memory +__device__ int gnu_add(int *p, int v) { + return __atomic_fetch_add(p, v, __ATOMIC_RELAXED); +} + +// LLVM: ![[#NOPRIV]] = !{i32 5, i32 6} `````````` </details> https://github.com/llvm/llvm-project/pull/232350 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
