https://github.com/rajatbajpai updated https://github.com/llvm/llvm-project/pull/229665
>From 18e3515253e7a161f34ab25f53939b6925b920df Mon Sep 17 00:00:00 2001 From: Rajat Bajpai <[email protected]> Date: Mon, 5 Oct 2026 12:28:40 +0000 Subject: [PATCH] [LLVM][NVPTX] Add fabric.try_put intrinsic support This change introduces fabric intrinsics design and try_put intrinsic support in NVPTX backend. --- clang/test/CodeGen/target-data.c | 6 +- llvm/docs/NVPTXUsage.md | 89 ++++++ llvm/include/llvm/IR/IntrinsicsNVVM.td | 42 +++ llvm/include/llvm/Support/NVPTXAddrSpace.h | 1 + llvm/lib/Target/NVPTX/NVPTX.h | 1 + llvm/lib/Target/NVPTX/NVPTXAliasAnalysis.cpp | 5 + llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp | 4 +- llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp | 70 +++++ llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 64 ++++- llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 54 ++++ .../Target/NVPTX/NVPTXPromoteParamAlign.cpp | 4 +- llvm/lib/Target/NVPTX/NVPTXUtilities.h | 10 +- llvm/lib/TargetParser/TargetDataLayout.cpp | 9 +- .../NVPTX/fabric-try-put-cache-hint.ll | 171 ++++++++++++ .../NVPTX/fabric-try-put-memory-effects.ll | 100 +++++++ llvm/test/CodeGen/NVPTX/fabric-try-put.ll | 262 ++++++++++++++++++ llvm/unittests/TargetParser/TripleTest.cpp | 12 +- 17 files changed, 883 insertions(+), 21 deletions(-) create mode 100644 llvm/test/CodeGen/NVPTX/fabric-try-put-cache-hint.ll create mode 100644 llvm/test/CodeGen/NVPTX/fabric-try-put-memory-effects.ll create mode 100644 llvm/test/CodeGen/NVPTX/fabric-try-put.ll diff --git a/clang/test/CodeGen/target-data.c b/clang/test/CodeGen/target-data.c index d738d5faf18617..62e2d86944bd71 100644 --- a/clang/test/CodeGen/target-data.c +++ b/clang/test/CodeGen/target-data.c @@ -144,15 +144,15 @@ // RUN: %clang_cc1 -triple nvptx-unknown -o - -emit-llvm %s | \ // RUN: FileCheck %s -check-prefix=NVPTX -// NVPTX: target datalayout = "e-p:32:32-i64:64-i128:128-i256:256-v16:16-v32:32-n16:32:64" +// NVPTX: target datalayout = "e-p:32:32-p8:128:128-ni:8-i64:64-i128:128-i256:256-v16:16-v32:32-n16:32:64" // RUN: %clang_cc1 -triple nvptx64-unknown -o - -emit-llvm %s | \ // RUN: FileCheck %s -check-prefix=NVPTX64 -// NVPTX64: target datalayout = "e-p6:32:32-i64:64-i128:128-i256:256-v16:16-v32:32-n16:32:64" +// NVPTX64: target datalayout = "e-p6:32:32-p8:128:128-ni:8-i64:64-i128:128-i256:256-v16:16-v32:32-n16:32:64" // RUN: %clang_cc1 -triple nvptx64-unknown -target-abi shortptr -o - \ // RUN: -emit-llvm %s | FileCheck %s -check-prefix=NVPTX64-SHORTPTR -// NVPTX64-SHORTPTR: target datalayout = "e-p3:32:32-p4:32:32-p5:32:32-p6:32:32-p7:32:32-p101:32:32-i64:64-i128:128-i256:256-v16:16-v32:32-n16:32:64" +// NVPTX64-SHORTPTR: target datalayout = "e-p3:32:32-p4:32:32-p5:32:32-p6:32:32-p7:32:32-p8:128:128-ni:8-p101:32:32-i64:64-i128:128-i256:256-v16:16-v32:32-n16:32:64" // RUN: %clang_cc1 -triple r600-unknown -o - -emit-llvm %s | \ // RUN: FileCheck %s -check-prefix=R600 diff --git a/llvm/docs/NVPTXUsage.md b/llvm/docs/NVPTXUsage.md index 088035ee83f1f2..8a76e0abfd85ac 100644 --- a/llvm/docs/NVPTXUsage.md +++ b/llvm/docs/NVPTXUsage.md @@ -119,6 +119,11 @@ The NVPTX back-end uses the following address space mapping: | 4 | Constant | | 5 | Local | | 7 | Shared Cluster | +| 8 | Fabric Handle | + +Pointers in address space 8 are opaque fabric handles; ordinary LLVM memory +instructions may not dereference them. See +[Fabric family of Intrinsics](#fabric-family-of-intrinsics). Every global variable and pointer type is assigned to one of these address spaces, with 0 being the default address space. Intrinsics are provided which @@ -3206,6 +3211,90 @@ The `override.addr` operands are described in For more information, refer [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-reduce-async-bulk-tensor). +### Fabric family of Intrinsics + +#### Overview: + +The '`llvm.nvvm.fabric.*`' intrinsics correspond to the fabric operations of +the PTX ISA. + +In LLVM IR, a fabric handle is an opaque `ptr addrspace(8)` value with a +128-bit, non-integral representation and 16-byte ABI alignment. Address space 8 +represents resource handles, not a PTX state space. Handles may be passed through +function arguments and returns, and through `phi` and `select` instructions. +However, ordinary dereferences, pointer arithmetic, address-space casts, and +integer conversions of handles are not supported. + +Memory accesses through fabric handles may alias accesses through global +(`ptr addrspace(1)`) pointers. + +The trailing `%flag_*` arguments must be compile-time constants. Fabric +operations require `sm_100` or higher and PTX ISA 9.3 or later; the +`.L2::cache_hint` variants require PTX ISA 9.4 or later. + +For execution semantics, synchronization, resource alignment, and error-reporting +requirements, refer to the [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#fabric-instructions). + +#### '`llvm.nvvm.fabric.handle_pair`' + +##### Syntax: + +```llvm +declare ptr addrspace(8) @llvm.nvvm.fabric.handle_pair(i32 %le_id, i64 %offset) +``` + +##### Overview: + +The '`@llvm.nvvm.fabric.handle_pair`' intrinsic constructs a fabric handle from +the 32-bit logical-endpoint identifier `%le_id` and the 64-bit byte offset +`%offset`. It has no dedicated PTX instruction. + +#### '`llvm.nvvm.fabric.try_put`' + +##### Syntax: + +```llvm +declare void @llvm.nvvm.fabric.try_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 %byte_mask, i64 %cache_hint, i1 %flag_cache_hint, i1 %flag_cp_mask, i32 %flag_multimem) +``` + +##### Overview: + +The '`@llvm.nvvm.fabric.try_put`' intrinsic corresponds to the +`fabric.try_put.async{.multimem}.*` family of PTX instructions. It copies +`%size` bytes from shared memory at `%src` to the resource referenced by +`%handle`, using `%mbar` as the completion barrier. + +- `i1 %flag_cache_hint`, when set, generates the `.L2::cache_hint` variant + with `%cache_hint` as the cache-policy operand; otherwise, `%cache_hint` + is ignored. +- `i1 %flag_cp_mask`, when set, generates the `.cp_mask` variant with + `%byte_mask` as the mask operand; otherwise, `%byte_mask` is ignored. +- `i32 %flag_multimem` accepts `0` (unicast) or `1` (`.multimem`). + +For more information, refer to the [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#fabric-instructions-try-put). + +#### '`llvm.nvvm.fabric.try_put.counted_writes`' + +##### Syntax: + +```llvm +declare void @llvm.nvvm.fabric.try_put.counted_writes(ptr addrspace(8) %data_handle, ptr addrspace(8) %counter_handle, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i64 %cache_hint, i1 %flag_cache_hint, i32 %flag_multimem) +``` + +##### Overview: + +The '`@llvm.nvvm.fabric.try_put.counted_writes`' intrinsic corresponds to the +`.counted::bytes` variants of `fabric.try_put.async{.multimem}.*`. +`%data_handle` and `%counter_handle` specify the data and counter resource +offsets. The remaining operands have the same meanings as for +'`@llvm.nvvm.fabric.try_put`'. + +Both handles map to the single endpoint identifier and the two resource +offsets of the PTX instruction. They must therefore have the same +logical-endpoint identifier; otherwise, the behavior is undefined. + +For more information, refer to the [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#fabric-instructions-try-put). + ### Warp Group Intrinsics #### '`llvm.nvvm.wgmma.fence.sync.aligned`' diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td index 08b71536ca9afb..e0df38c25f7ca1 100644 --- a/llvm/include/llvm/IR/IntrinsicsNVVM.td +++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td @@ -150,6 +150,7 @@ def llvm_constant_ptr_ty: LLVMQualPointerType<4>; // (const)ptr def llvm_local_ptr_ty : LLVMQualPointerType<5>; // (local)ptr def llvm_tmem_ptr_ty : LLVMQualPointerType<6>; // (tensor memory)ptr def llvm_shared_cluster_ptr_ty : LLVMQualPointerType<7>; // (shared_cluster)ptr +def llvm_fabric_handle_ptr_ty : LLVMQualPointerType<8>; // (fabric_handle)ptr // // MISC @@ -4268,4 +4269,45 @@ def int_nvvm_tensormap_replace_fill_mode : ArgInfo<ArgIndex<1>, [ArgName<"fill_mode">, ImmArgPrinter<"printTensormapFillMode">]>]>; +// Create an opaque 128-bit fabric handle from a 32-bit logical-endpoint id and +// a 64-bit offset. +def int_nvvm_fabric_handle_pair + : NVVMPureIntrinsic<[llvm_fabric_handle_ptr_ty], + [llvm_i32_ty, llvm_i64_ty], [], + "llvm.nvvm.fabric.handle_pair">; + +def int_nvvm_fabric_try_put + : DefaultAttrsIntrinsicFlags<[], + [llvm_fabric_handle_ptr_ty, // fabric handle + llvm_shared_ptr_ty, // src_smem_ptr + llvm_shared_ptr_ty, // mbarrier_ptr + llvm_i32_ty, // size (bytes) + llvm_i16_ty, // byte_mask (16-bit bytemask) + llvm_i64_ty], // cache_hint + [llvm_i1_ty, // flag_cache_hint + llvm_i1_ty, // flag_cp_mask + llvm_i32_ty], // flag_multimem (unitcast or multicast) + [IntrConvergent, IntrArgMemOnly, ReadOnly<ArgIndex<1>>, + NVVM_ARG_NAME_PROP<6, "flag_cache_hint">.Prop, + ArgInfo<ArgIndex<7>, [ArgName<"flag_cp_mask">]>, + Range<ArgIndex<8>, 0, 2>, + ArgInfo<ArgIndex<8>, [ArgName<"flag_multimem">]>], + "llvm.nvvm.fabric.try_put">; + +def int_nvvm_fabric_try_put_counted_writes + : DefaultAttrsIntrinsicFlags<[], + [llvm_fabric_handle_ptr_ty, // fabric handle (data) + llvm_fabric_handle_ptr_ty, // fabric handle (counter) + llvm_shared_ptr_ty, // src_smem_ptr + llvm_shared_ptr_ty, // mbarrier_ptr + llvm_i32_ty, // size (bytes) + llvm_i64_ty], // cache_hint + [llvm_i1_ty, // flag_cache_hint + llvm_i32_ty], // flag_multimem (unitcast or multicast) + [IntrConvergent, IntrArgMemOnly, ReadOnly<ArgIndex<2>>, + NVVM_ARG_NAME_PROP<6, "flag_cache_hint">.Prop, + Range<ArgIndex<7>, 0, 2>, + ArgInfo<ArgIndex<7>, [ArgName<"flag_multimem">]>], + "llvm.nvvm.fabric.try_put.counted_writes">; + } // let TargetPrefix = "nvvm" diff --git a/llvm/include/llvm/Support/NVPTXAddrSpace.h b/llvm/include/llvm/Support/NVPTXAddrSpace.h index 12c493b6568fb4..f4562a61012b6b 100644 --- a/llvm/include/llvm/Support/NVPTXAddrSpace.h +++ b/llvm/include/llvm/Support/NVPTXAddrSpace.h @@ -26,6 +26,7 @@ enum AddressSpace : unsigned { ADDRESS_SPACE_LOCAL = 5, ADDRESS_SPACE_TENSOR = 6, ADDRESS_SPACE_SHARED_CLUSTER = 7, + ADDRESS_SPACE_FABRIC_HANDLE = 8, ADDRESS_SPACE_ENTRY_PARAM = 101, }; diff --git a/llvm/lib/Target/NVPTX/NVPTX.h b/llvm/lib/Target/NVPTX/NVPTX.h index 6a58b55acd29d4..284ce1c3330342 100644 --- a/llvm/lib/Target/NVPTX/NVPTX.h +++ b/llvm/lib/Target/NVPTX/NVPTX.h @@ -324,6 +324,7 @@ enum AddressSpace : AddressSpaceUnderlyingType { Const = NVPTXAS::ADDRESS_SPACE_CONST, Local = NVPTXAS::ADDRESS_SPACE_LOCAL, SharedCluster = NVPTXAS::ADDRESS_SPACE_SHARED_CLUSTER, + FabricHandle = NVPTXAS::ADDRESS_SPACE_FABRIC_HANDLE, EntryParam = NVPTXAS::ADDRESS_SPACE_ENTRY_PARAM, // DeviceParam is not a real address space, as it does not support pointers diff --git a/llvm/lib/Target/NVPTX/NVPTXAliasAnalysis.cpp b/llvm/lib/Target/NVPTX/NVPTXAliasAnalysis.cpp index f86b5975add740..3ee4057b8b5788 100644 --- a/llvm/lib/Target/NVPTX/NVPTXAliasAnalysis.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXAliasAnalysis.cpp @@ -92,6 +92,11 @@ static AliasResult::Kind getAliasResult(unsigned AS1, unsigned AS2) { ((AS1 == ADDRESS_SPACE_SHARED_CLUSTER) && (AS2 == ADDRESS_SPACE_SHARED))) return AliasResult::MayAlias; + // Fabric endpoint resources may also be accessed through global pointers. + if (((AS1 == ADDRESS_SPACE_FABRIC_HANDLE) && (AS2 == ADDRESS_SPACE_GLOBAL)) || + ((AS1 == ADDRESS_SPACE_GLOBAL) && (AS2 == ADDRESS_SPACE_FABRIC_HANDLE))) + return AliasResult::MayAlias; + return (AS1 == AS2 ? AliasResult::MayAlias : AliasResult::NoAlias); } diff --git a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp index ab90036722ca46..32930806668261 100644 --- a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp @@ -680,7 +680,7 @@ static void printParam(const OwnerT *Owner, Type *Ty, unsigned AttrIdx, const DataLayout &DL, raw_ostream &O) { O << ".param "; - if (IsByVal || shouldPassAsArray(Ty)) { + if (IsByVal || shouldPassAsArray(Ty, DL)) { const Align ParamAlign = IsByVal && !IsKernel ? getDeviceByValParamAlign(Owner, Ty, AttrIdx, DL) : getPTXParamAlign(Owner, Ty, AttrIdx, DL); @@ -1782,7 +1782,7 @@ void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, raw_ostream &O) { // A byval param is passed as a copy of the pointee and an aggregate is // passed as a blob of bytes; both are declared as a byte array. const bool IsByVal = Arg.hasByValAttr(); - const bool AsArray = IsByVal || shouldPassAsArray(Ty); + const bool AsArray = IsByVal || shouldPassAsArray(Ty, DL); // Kernels declare image/sampler handles and the address space of a // pointee. Both of those are scalar handles, so a byte-array param is diff --git a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp index b952caf3a7775f..90054d7f9d65c4 100644 --- a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp @@ -117,6 +117,8 @@ class NVPTXDAGToDAGISel : public SelectionDAGISel { void Select(SDNode *N) override; bool tryIntrinsicChain(SDNode *N); bool tryIntrinsicVoid(SDNode *N); + /// Selects the machine instruction for fabric try put intrinsic. + void SelectFabricTryPut(SDNode *N, unsigned IID); bool tryLoad(SDNode *N); bool tryLoadVector(SDNode *N); bool tryLDU(SDNode *N); @@ -636,6 +638,7 @@ NVPTX::AddressSpace NVPTXDAGToDAGISel::getAddrSpace(const MemSDNode *N) { case NVPTX::AddressSpace::Const: case NVPTX::AddressSpace::Local: case NVPTX::AddressSpace::SharedCluster: + case NVPTX::AddressSpace::FabricHandle: case NVPTX::AddressSpace::EntryParam: case NVPTX::AddressSpace::DeviceParam: return AS; @@ -2351,7 +2354,74 @@ bool NVPTXDAGToDAGISel::tryIntrinsicVoid(SDNode *N) { SelectTcgen05St(N, /* hasOffset */ true); return true; } + case Intrinsic::nvvm_fabric_try_put: + case Intrinsic::nvvm_fabric_try_put_counted_writes: + SelectFabricTryPut(N, IID); + return true; + } +} + +// After ISelLowering, operand layout is: +// try_put: +// [chain, IID, leId, offset, src, mbar, size, byte_mask, cache_hint, +// flag_cache_hint, flag_cp_mask, flag_multimem] +// try_put.counted_writes: +// [chain, IID, leId, offset_data, offset_counter, src, mbar, size, +// cache_hint, flag_cache_hint, flag_multimem] +void NVPTXDAGToDAGISel::SelectFabricTryPut(SDNode *N, unsigned IID) { + if (Subtarget->getSmVersion() < 100 || Subtarget->getPTXVersion() < 93) + report_fatal_error( + "fabric.try_put requires sm_100 or higher and PTX 9.3 or higher"); + + const bool IsCounted = IID == Intrinsic::nvvm_fabric_try_put_counted_writes; + // The try_put carries one fabric handle (one offset after leId); + // counted_writes carries two (two offsets). src follows the leId/offset + // args in both layouts. + const unsigned SrcIdx = IsCounted ? 5 : 4; + const unsigned FlagCacheHintIdx = N->getNumOperands() - (IsCounted ? 2 : 3); + const unsigned FlagCpMaskIdx = N->getNumOperands() - 2; // try_put only + + bool HasCacheHint = N->getConstantOperandVal(FlagCacheHintIdx); + if (HasCacheHint && Subtarget->getPTXVersion() < 94) + report_fatal_error("fabric.try_put with .L2::cache_hint requires PTX 9.4 " + "or higher"); + + bool HasCpMask = !IsCounted && N->getConstantOperandVal(FlagCpMaskIdx); + + unsigned Opcode; + switch (IID) { + default: + llvm_unreachable("Unexpected fabric.try_put intrinsic"); + case Intrinsic::nvvm_fabric_try_put: + Opcode = HasCacheHint ? (HasCpMask ? NVPTX::FABRIC_TRY_PUT_ASYNC_S2F_MASK_CH + : NVPTX::FABRIC_TRY_PUT_ASYNC_S2F_CH) + : (HasCpMask ? NVPTX::FABRIC_TRY_PUT_ASYNC_S2F_MASK + : NVPTX::FABRIC_TRY_PUT_ASYNC_S2F); + break; + case Intrinsic::nvvm_fabric_try_put_counted_writes: + Opcode = HasCacheHint ? NVPTX::FABRIC_TRY_PUT_ASYNC_S2F_COUNTED_WRITES_CH + : NVPTX::FABRIC_TRY_PUT_ASYNC_S2F_COUNTED_WRITES; + break; } + + SDLoc DL(N); + SmallVector<SDValue, 12> Ops; + // Handle args (leId, offset [, offset_counter]). + Ops.append(N->op_begin() + 2, N->op_begin() + SrcIdx); + const auto [SrcBase, SrcOffset] = selectADDR(N->getOperand(SrcIdx), CurDAG); + const auto [MbarBase, MbarOffset] = + selectADDR(N->getOperand(SrcIdx + 1), CurDAG); + Ops.append({SrcBase, SrcOffset, MbarBase, MbarOffset}); + + Ops.push_back(N->getOperand(SrcIdx + 2)); + if (HasCpMask) + Ops.push_back(N->getOperand(SrcIdx + 3)); + if (HasCacheHint) + Ops.push_back(N->getOperand(IsCounted ? SrcIdx + 3 : SrcIdx + 4)); + Ops.push_back(N->getOperand(N->getNumOperands() - 1)); + + Ops.push_back(N->getOperand(0)); // chain + ReplaceNode(N, CurDAG->getMachineNode(Opcode, DL, N->getVTList(), Ops)); } void NVPTXDAGToDAGISel::selectAtomicSwap128(SDNode *N) { diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp index dab77d1cf3c7e9..87cc20bb2939c9 100644 --- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp @@ -1423,7 +1423,7 @@ SDValue NVPTXTargetLowering::LowerCall(TargetLowering::CallLoweringInfo &CLI, if (IsVAArg) return VADeclareParam; - if (IsByVal || shouldPassAsArray(Arg.Ty)) + if (IsByVal || shouldPassAsArray(Arg.Ty, DL)) return MakeDeclareArrayParam(ParamSymbol, ArgAlign, TySize); assert(ArgOuts.size() == 1 && "We must pass only one value as non-array"); @@ -1558,7 +1558,7 @@ SDValue NVPTXTargetLowering::LowerCall(TargetLowering::CallLoweringInfo &CLI, if (!Ins.empty()) { const SDValue RetSymbol = getSymbolNode(DAG, "retval0", MVT::i32); const unsigned ResultSize = DL.getTypeAllocSize(RetTy); - if (shouldPassAsArray(RetTy)) { + if (shouldPassAsArray(RetTy, DL)) { const Align RetAlign = getPTXParamAlign(CB, RetTy, AttributeList::ReturnIndex, DL); MakeDeclareArrayParam(RetSymbol, RetAlign, ResultSize); @@ -2871,6 +2871,41 @@ static SDValue lowerTensormapReplaceSwizzleMode(SDValue Op, SelectionDAG &DAG) { return Op; } +// Each i128 fabric handle is split into i32 leId + i64 offset. +static SDValue lowerFabricHandles(SDValue Op, SelectionDAG &DAG) { + SDNode *N = Op.getNode(); + SDLoc DL(N); + const unsigned IID = N->getConstantOperandVal(1); + const unsigned NumHandles = + IID == Intrinsic::nvvm_fabric_try_put_counted_writes ? 2 : 1; + + const unsigned FirstHandleIdx = 2; + const unsigned HandleEnd = FirstHandleIdx + NumHandles; + + // The custom hook fires every time the legalizer revisits this + // INTRINSIC_VOID. After we rewrite it once, operand FirstHandleIdx is the + // i32 leId we just emitted, not an i128 — bail out so we don't re-lower. + SDValue FirstHandle = N->getOperand(FirstHandleIdx); + if (FirstHandle.getValueType() != MVT::i128) + return Op; + + SmallVector<SDValue, 8> Ops = {N->getOperand(0), N->getOperand(1)}; + + SDValue LeId = DAG.getNode(ISD::EXTRACT_ELEMENT, DL, MVT::i64, FirstHandle, + DAG.getIntPtrConstant(0, DL)); + Ops.push_back(DAG.getNode(ISD::TRUNCATE, DL, MVT::i32, LeId)); + + // Add each handle's offset. + for (unsigned I = FirstHandleIdx; I != HandleEnd; ++I) { + SDValue Handle = N->getOperand(I); + Ops.push_back(DAG.getNode(ISD::EXTRACT_ELEMENT, DL, MVT::i64, Handle, + DAG.getIntPtrConstant(1, DL))); + } + + Ops.append(N->op_begin() + HandleEnd, N->op_end()); + return DAG.getNode(ISD::INTRINSIC_VOID, DL, N->getVTList(), Ops); +} + static SDValue lowerIntrinsicVoid(SDValue Op, SelectionDAG &DAG) { SDNode *N = Op.getNode(); SDValue Intrin = N->getOperand(1); @@ -2967,6 +3002,9 @@ static SDValue lowerIntrinsicVoid(SDValue Op, SelectionDAG &DAG) { return lowerTensormapReplaceElemtype(Op, DAG); case Intrinsic::nvvm_tensormap_replace_swizzle_mode: return lowerTensormapReplaceSwizzleMode(Op, DAG); + case Intrinsic::nvvm_fabric_try_put: + case Intrinsic::nvvm_fabric_try_put_counted_writes: + return lowerFabricHandles(Op, DAG); } return Op; } @@ -7717,6 +7755,25 @@ static void replaceAtomicSwap128(SDNode *N, SelectionDAG &DAG, Results.push_back(Result.getValue(2)); } +// Replace fabric handle with two i64 operands (leId and offset). +static void replaceFabricHandlePair(SDNode *N, SelectionDAG &DAG, + SmallVectorImpl<SDValue> &Results) { + SDLoc DL(N); + SDValue LeId = N->getOperand(1); // i32 + SDValue Offset = N->getOperand(2); // i64 + SDValue LeIdExt = DAG.getNode(ISD::ZERO_EXTEND, DL, MVT::i64, LeId); + Results.push_back( + DAG.getNode(ISD::BUILD_PAIR, DL, MVT::i128, LeIdExt, Offset)); +} + +static void ReplaceINTRINSIC_WO_CHAIN(SDNode *N, SelectionDAG &DAG, + SmallVectorImpl<SDValue> &Results) { + unsigned IID = N->getConstantOperandVal(0); + if (IID != Intrinsic::nvvm_fabric_handle_pair) + return; + return replaceFabricHandlePair(N, DAG, Results); +} + void NVPTXTargetLowering::ReplaceNodeResults( SDNode *N, SmallVectorImpl<SDValue> &Results, SelectionDAG &DAG) const { switch (N->getOpcode()) { @@ -7732,6 +7789,9 @@ void NVPTXTargetLowering::ReplaceNodeResults( case ISD::INTRINSIC_W_CHAIN: ReplaceINTRINSIC_W_CHAIN(N, DAG, Results); return; + case ISD::INTRINSIC_WO_CHAIN: + ReplaceINTRINSIC_WO_CHAIN(N, DAG, Results); + return; case ISD::CopyFromReg: ReplaceCopyFromReg_128(N, DAG, Results); return; diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td index d098e78d6edbfa..561274b6e73dc1 100644 --- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td +++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td @@ -7182,6 +7182,60 @@ foreach dim = ["x", "y", "z"] in { CLUSTERLAUNCHCONTROL_QUERY_CANCEL_GET_FIRST_CTAID<dim>; } +// +// Fabric Instructions +// + +def FabricMultimemFlag : Operand<i32> { + let PrintMethod = "printMultimem"; +} + +multiclass FABRIC_TRY_PUT_ASYNC_S2F_INTR { + defvar ins_base = (ins B32:$leId, B64:$offset, ADDR:$src, ADDR:$mbar, + B32:$size); + defvar ins_multimem = (ins FabricMultimemFlag:$is_multimem); + defvar asm_head = "fabric.try_put.async.${is_multimem}shared::cta" + ".mbarrier::complete_tx::16B.mbarrier::report::fabric"; + defvar asm_body = ".relaxed.sys.b128" + "\t[$leId, $offset], [$src], $size, [$mbar]"; + defvar args_mask = ", $mask"; + defvar args_ch = ", $cache_hint"; + + def "" : NVPTXInst<(outs), !con(ins_base, ins_multimem), + asm_head # asm_body # ";", []>; + def _MASK : NVPTXInst<(outs), !con(ins_base, (ins B16:$mask), ins_multimem), + asm_head # ".cp_mask" # asm_body # args_mask # ";", []>; + def _CH : NVPTXInst<(outs), + !con(ins_base, (ins B64:$cache_hint), ins_multimem), + asm_head # ".L2::cache_hint" # asm_body # args_ch # ";", []>; + def _MASK_CH : NVPTXInst<(outs), + !con(ins_base, (ins B16:$mask, B64:$cache_hint), ins_multimem), + asm_head # ".L2::cache_hint.cp_mask" # asm_body # args_mask # args_ch + # ";", []>; +} +defm FABRIC_TRY_PUT_ASYNC_S2F : FABRIC_TRY_PUT_ASYNC_S2F_INTR; + +// .counted::bytes form (counter-based completion; two offsets from the shared +// leId: data offset then counter offset. +multiclass FABRIC_TRY_PUT_ASYNC_S2F_COUNTED_WRITES_INTR { + defvar ins_base = (ins B32:$leId, B64:$offset1, B64:$offset2, ADDR:$src, + ADDR:$mbar, B32:$size); + defvar ins_multimem = (ins FabricMultimemFlag:$is_multimem); + defvar asm_head = "fabric.try_put.async.${is_multimem}shared::cta" + ".mbarrier::complete_tx::16B.mbarrier::report::fabric" + ".counted::bytes"; + defvar asm_body = ".relaxed.sys.b128" + "\t[$leId, $offset1, $offset2], [$src], $size, [$mbar]"; + defvar args_ch = ", $cache_hint"; + + def "" : NVPTXInst<(outs), !con(ins_base, ins_multimem), + asm_head # asm_body # ";", []>; + def _CH : NVPTXInst<(outs), + !con(ins_base, (ins B64:$cache_hint), ins_multimem), + asm_head # ".L2::cache_hint" # asm_body # args_ch # ";", []>; +} +defm FABRIC_TRY_PUT_ASYNC_S2F_COUNTED_WRITES : FABRIC_TRY_PUT_ASYNC_S2F_COUNTED_WRITES_INTR; + // // tcgen05.mma Helpers // diff --git a/llvm/lib/Target/NVPTX/NVPTXPromoteParamAlign.cpp b/llvm/lib/Target/NVPTX/NVPTXPromoteParamAlign.cpp index ca0405743c0b12..540c23368444c0 100644 --- a/llvm/lib/Target/NVPTX/NVPTXPromoteParamAlign.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXPromoteParamAlign.cpp @@ -182,7 +182,7 @@ static bool promoteParamAlign(Function &F) { for (Argument &Arg : F.args()) { const bool IsByVal = Arg.hasByValAttr(); Type *ArgTy = IsByVal ? Arg.getParamByValType() : Arg.getType(); - if (ArgTy->isEmptyTy() || (!IsByVal && !shouldPassAsArray(ArgTy))) + if (ArgTy->isEmptyTy() || (!IsByVal && !shouldPassAsArray(ArgTy, DL))) continue; // An explicit stackalign already wins at emission time, nothing to promote. @@ -206,7 +206,7 @@ static bool promoteParamAlign(Function &F) { // Promote an aggregate return value. Type *RetTy = F.getReturnType(); - if (shouldPassAsArray(RetTy) && !RetTy->isEmptyTy() && + if (shouldPassAsArray(RetTy, DL) && !RetTy->isEmptyTy() && !F.getAttributes().getRetStackAlignment()) { const MaybeAlign PromotedAlign = getPromotedParamAlign(getPTXParamTypeAlign(RetTy, DL)); diff --git a/llvm/lib/Target/NVPTX/NVPTXUtilities.h b/llvm/lib/Target/NVPTX/NVPTXUtilities.h index 2567ea793095b0..a8d68dd7a258e3 100644 --- a/llvm/lib/Target/NVPTX/NVPTXUtilities.h +++ b/llvm/lib/Target/NVPTX/NVPTXUtilities.h @@ -16,6 +16,7 @@ #include "NVPTX.h" #include "llvm/ADT/StringExtras.h" #include "llvm/CodeGen/ValueTypes.h" +#include "llvm/IR/DataLayout.h" #include "llvm/IR/Function.h" #include "llvm/IR/InstrTypes.h" #include "llvm/IR/Value.h" @@ -26,7 +27,6 @@ namespace llvm { -class DataLayout; class MemSDNode; Function *getMaybeBitcastedCallee(const CallBase *CB); @@ -68,9 +68,9 @@ inline unsigned promoteScalarKernelArgumentSize(unsigned Size) { return PowerOf2Ceil(std::max(Size, 8U)); } -inline bool shouldPassAsArray(Type *Ty) { - return Ty->isAggregateType() || Ty->isVectorTy() || - Ty->getScalarSizeInBits() >= 128 || Ty->isHalfTy() || Ty->isBFloatTy(); +inline bool shouldPassAsArray(Type *Ty, const DataLayout &DL) { + return Ty->isAggregateType() || Ty->isVectorTy() || Ty->isHalfTy() || + Ty->isBFloatTy() || (Ty->isSized() && DL.getTypeSizeInBits(Ty) >= 128); } namespace NVPTX { @@ -179,6 +179,8 @@ inline const char *addressSpaceToString(AddressSpace A, return UseParamSubqualifiers ? "param::func" : "param"; case AddressSpace::Local: return "local"; + case AddressSpace::FabricHandle: + return "fabric"; } report_fatal_error(formatv("Unknown NVPTX::AddressSpace \"{}\".", static_cast<AddressSpaceUnderlyingType>(A))); diff --git a/llvm/lib/TargetParser/TargetDataLayout.cpp b/llvm/lib/TargetParser/TargetDataLayout.cpp index 83bc4a624bf99b..4cbe72b9e19a5f 100644 --- a/llvm/lib/TargetParser/TargetDataLayout.cpp +++ b/llvm/lib/TargetParser/TargetDataLayout.cpp @@ -495,7 +495,6 @@ static std::string computeNVPTXDataLayout(const Triple &T, StringRef ABIName) { // - constant (addrspace:4) // - local (addrspace:5) // - shared cluster (addrspace:7) - // - entry parameter (addrspace:101) if (IsShortPtr) Ret += "-p3:32:32-p4:32:32-p5:32:32"; @@ -503,9 +502,15 @@ static std::string computeNVPTXDataLayout(const Triple &T, StringRef ABIName) { Ret += "-p6:32:32"; if (IsShortPtr) - Ret += "-p7:32:32-p101:32:32"; + Ret += "-p7:32:32"; } + // Fabric handles (addrspace:8) + Ret += "-p8:128:128-ni:8"; + + if (!Is32Bit && IsShortPtr) + Ret += "-p101:32:32"; + Ret += "-i64:64-i128:128-i256:256-v16:16-v32:32-n16:32:64"; return Ret; diff --git a/llvm/test/CodeGen/NVPTX/fabric-try-put-cache-hint.ll b/llvm/test/CodeGen/NVPTX/fabric-try-put-cache-hint.ll new file mode 100644 index 00000000000000..ead96a81b2d479 --- /dev/null +++ b/llvm/test/CodeGen/NVPTX/fabric-try-put-cache-hint.ll @@ -0,0 +1,171 @@ +; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6 +; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx94 | FileCheck %s +; RUN: llvm-as < %s | llvm-dis | FileCheck --check-prefix=FORMAT %s +; RUN: %if ptxas-sm_100 && ptxas-isa-9.4 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx94 | %ptxas-verify -arch=sm_100 %} + +target triple = "nvptx64-nvidia-cuda" + +declare ptr addrspace(8) @llvm.nvvm.fabric.handle_pair(i32, i64) +declare void @llvm.nvvm.fabric.try_put(ptr addrspace(8), ptr addrspace(3), ptr addrspace(3), i32, i16, i64, i1, i1, i32) +declare void @llvm.nvvm.fabric.try_put.counted_writes(ptr addrspace(8), ptr addrspace(8), ptr addrspace(3), ptr addrspace(3), i32, i64, i1, i32) + +define void @test_fabric_try_put_cache(i32 %leId, i64 %offset, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i64 %cacheHint) { +; CHECK-LABEL: test_fabric_try_put_cache( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-NEXT: .reg .b64 %rd<5>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [test_fabric_try_put_cache_param_0]; +; CHECK-NEXT: ld.param::func.b64 %rd1, [test_fabric_try_put_cache_param_1]; +; CHECK-NEXT: ld.param::func.b64 %rd2, [test_fabric_try_put_cache_param_2]; +; CHECK-NEXT: ld.param::func.b64 %rd3, [test_fabric_try_put_cache_param_3]; +; CHECK-NEXT: ld.param::func.b32 %r2, [test_fabric_try_put_cache_param_4]; +; CHECK-NEXT: ld.param::func.b64 %rd4, [test_fabric_try_put_cache_param_5]; +; CHECK-NEXT: fabric.try_put.async.shared::cta.mbarrier::complete_tx::16B.mbarrier::report::fabric.L2::cache_hint.relaxed.sys.b128 [%r1, %rd1], [%rd2], %r2, [%rd3], %rd4; +; CHECK-NEXT: ret; +; FORMAT-LABEL: define void @test_fabric_try_put_cache( +; FORMAT: call void @llvm.nvvm.fabric.try_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 0, i64 %cacheHint, /* flag_cache_hint= */ i1 true, /* flag_cp_mask= */ i1 false, /* flag_multimem= */ i32 0) + %handle = call ptr addrspace(8) @llvm.nvvm.fabric.handle_pair(i32 %leId, i64 %offset) + call void @llvm.nvvm.fabric.try_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 0, i64 %cacheHint, i1 true, i1 false, i32 0) + ret void +} + +define void @test_fabric_try_put_multimem_cache(i32 %leId, i64 %offset, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i64 %cacheHint) { +; CHECK-LABEL: test_fabric_try_put_multimem_cache( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-NEXT: .reg .b64 %rd<5>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [test_fabric_try_put_multimem_cache_param_0]; +; CHECK-NEXT: ld.param::func.b64 %rd1, [test_fabric_try_put_multimem_cache_param_1]; +; CHECK-NEXT: ld.param::func.b64 %rd2, [test_fabric_try_put_multimem_cache_param_2]; +; CHECK-NEXT: ld.param::func.b64 %rd3, [test_fabric_try_put_multimem_cache_param_3]; +; CHECK-NEXT: ld.param::func.b32 %r2, [test_fabric_try_put_multimem_cache_param_4]; +; CHECK-NEXT: ld.param::func.b64 %rd4, [test_fabric_try_put_multimem_cache_param_5]; +; CHECK-NEXT: fabric.try_put.async.multimem.shared::cta.mbarrier::complete_tx::16B.mbarrier::report::fabric.L2::cache_hint.relaxed.sys.b128 [%r1, %rd1], [%rd2], %r2, [%rd3], %rd4; +; CHECK-NEXT: ret; +; FORMAT-LABEL: define void @test_fabric_try_put_multimem_cache( +; FORMAT: call void @llvm.nvvm.fabric.try_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 0, i64 %cacheHint, /* flag_cache_hint= */ i1 true, /* flag_cp_mask= */ i1 false, /* flag_multimem= */ i32 1) + %handle = call ptr addrspace(8) @llvm.nvvm.fabric.handle_pair(i32 %leId, i64 %offset) + call void @llvm.nvvm.fabric.try_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 0, i64 %cacheHint, i1 true, i1 false, i32 1) + ret void +} + +define void @test_fabric_try_put_mask_cache(i32 %leId, i64 %offset, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 %mask, i64 %cacheHint) { +; CHECK-LABEL: test_fabric_try_put_mask_cache( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-NEXT: .reg .b64 %rd<5>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [test_fabric_try_put_mask_cache_param_0]; +; CHECK-NEXT: ld.param::func.b64 %rd1, [test_fabric_try_put_mask_cache_param_1]; +; CHECK-NEXT: ld.param::func.b64 %rd2, [test_fabric_try_put_mask_cache_param_2]; +; CHECK-NEXT: ld.param::func.b64 %rd3, [test_fabric_try_put_mask_cache_param_3]; +; CHECK-NEXT: ld.param::func.b32 %r2, [test_fabric_try_put_mask_cache_param_4]; +; CHECK-NEXT: ld.param::func.b16 %rs1, [test_fabric_try_put_mask_cache_param_5]; +; CHECK-NEXT: ld.param::func.b64 %rd4, [test_fabric_try_put_mask_cache_param_6]; +; CHECK-NEXT: fabric.try_put.async.shared::cta.mbarrier::complete_tx::16B.mbarrier::report::fabric.L2::cache_hint.cp_mask.relaxed.sys.b128 [%r1, %rd1], [%rd2], %r2, [%rd3], %rs1, %rd4; +; CHECK-NEXT: ret; +; FORMAT-LABEL: define void @test_fabric_try_put_mask_cache( +; FORMAT: call void @llvm.nvvm.fabric.try_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 %mask, i64 %cacheHint, /* flag_cache_hint= */ i1 true, /* flag_cp_mask= */ i1 true, /* flag_multimem= */ i32 0) + %handle = call ptr addrspace(8) @llvm.nvvm.fabric.handle_pair(i32 %leId, i64 %offset) + call void @llvm.nvvm.fabric.try_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 %mask, i64 %cacheHint, i1 true, i1 true, i32 0) + ret void +} + +define void @test_fabric_try_put_mask_multimem_cache(i32 %leId, i64 %offset, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 %mask, i64 %cacheHint) { +; CHECK-LABEL: test_fabric_try_put_mask_multimem_cache( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-NEXT: .reg .b64 %rd<5>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [test_fabric_try_put_mask_multimem_cache_param_0]; +; CHECK-NEXT: ld.param::func.b64 %rd1, [test_fabric_try_put_mask_multimem_cache_param_1]; +; CHECK-NEXT: ld.param::func.b64 %rd2, [test_fabric_try_put_mask_multimem_cache_param_2]; +; CHECK-NEXT: ld.param::func.b64 %rd3, [test_fabric_try_put_mask_multimem_cache_param_3]; +; CHECK-NEXT: ld.param::func.b32 %r2, [test_fabric_try_put_mask_multimem_cache_param_4]; +; CHECK-NEXT: ld.param::func.b16 %rs1, [test_fabric_try_put_mask_multimem_cache_param_5]; +; CHECK-NEXT: ld.param::func.b64 %rd4, [test_fabric_try_put_mask_multimem_cache_param_6]; +; CHECK-NEXT: fabric.try_put.async.multimem.shared::cta.mbarrier::complete_tx::16B.mbarrier::report::fabric.L2::cache_hint.cp_mask.relaxed.sys.b128 [%r1, %rd1], [%rd2], %r2, [%rd3], %rs1, %rd4; +; CHECK-NEXT: ret; +; FORMAT-LABEL: define void @test_fabric_try_put_mask_multimem_cache( +; FORMAT: call void @llvm.nvvm.fabric.try_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 %mask, i64 %cacheHint, /* flag_cache_hint= */ i1 true, /* flag_cp_mask= */ i1 true, /* flag_multimem= */ i32 1) + %handle = call ptr addrspace(8) @llvm.nvvm.fabric.handle_pair(i32 %leId, i64 %offset) + call void @llvm.nvvm.fabric.try_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 %mask, i64 %cacheHint, i1 true, i1 true, i32 1) + ret void +} + +define void @test_fabric_try_put_counted_cache(i32 %leId, i64 %dataOffset, i64 %counterOffset, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i64 %cacheHint) { +; CHECK-LABEL: test_fabric_try_put_counted_cache( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-NEXT: .reg .b64 %rd<6>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [test_fabric_try_put_counted_cache_param_0]; +; CHECK-NEXT: ld.param::func.b64 %rd1, [test_fabric_try_put_counted_cache_param_1]; +; CHECK-NEXT: ld.param::func.b64 %rd2, [test_fabric_try_put_counted_cache_param_2]; +; CHECK-NEXT: ld.param::func.b64 %rd3, [test_fabric_try_put_counted_cache_param_3]; +; CHECK-NEXT: ld.param::func.b64 %rd4, [test_fabric_try_put_counted_cache_param_4]; +; CHECK-NEXT: ld.param::func.b32 %r2, [test_fabric_try_put_counted_cache_param_5]; +; CHECK-NEXT: ld.param::func.b64 %rd5, [test_fabric_try_put_counted_cache_param_6]; +; CHECK-NEXT: fabric.try_put.async.shared::cta.mbarrier::complete_tx::16B.mbarrier::report::fabric.counted::bytes.L2::cache_hint.relaxed.sys.b128 [%r1, %rd1, %rd2], [%rd3], %r2, [%rd4], %rd5; +; CHECK-NEXT: ret; +; FORMAT-LABEL: define void @test_fabric_try_put_counted_cache( +; FORMAT: call void @llvm.nvvm.fabric.try_put.counted_writes(ptr addrspace(8) %handle_data, ptr addrspace(8) %handle_counted, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i64 %cacheHint, /* flag_cache_hint= */ i1 true, /* flag_multimem= */ i32 0) + %handle_data = call ptr addrspace(8) @llvm.nvvm.fabric.handle_pair(i32 %leId, i64 %dataOffset) + %handle_counted = call ptr addrspace(8) @llvm.nvvm.fabric.handle_pair(i32 %leId, i64 %counterOffset) + call void @llvm.nvvm.fabric.try_put.counted_writes(ptr addrspace(8) %handle_data, ptr addrspace(8) %handle_counted, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i64 %cacheHint, i1 true, i32 0) + ret void +} + +define void @test_fabric_try_put_counted_multimem_cache(i32 %leId, i64 %dataOffset, i64 %counterOffset, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i64 %cacheHint) { +; CHECK-LABEL: test_fabric_try_put_counted_multimem_cache( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-NEXT: .reg .b64 %rd<6>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [test_fabric_try_put_counted_multimem_cache_param_0]; +; CHECK-NEXT: ld.param::func.b64 %rd1, [test_fabric_try_put_counted_multimem_cache_param_1]; +; CHECK-NEXT: ld.param::func.b64 %rd2, [test_fabric_try_put_counted_multimem_cache_param_2]; +; CHECK-NEXT: ld.param::func.b64 %rd3, [test_fabric_try_put_counted_multimem_cache_param_3]; +; CHECK-NEXT: ld.param::func.b64 %rd4, [test_fabric_try_put_counted_multimem_cache_param_4]; +; CHECK-NEXT: ld.param::func.b32 %r2, [test_fabric_try_put_counted_multimem_cache_param_5]; +; CHECK-NEXT: ld.param::func.b64 %rd5, [test_fabric_try_put_counted_multimem_cache_param_6]; +; CHECK-NEXT: fabric.try_put.async.multimem.shared::cta.mbarrier::complete_tx::16B.mbarrier::report::fabric.counted::bytes.L2::cache_hint.relaxed.sys.b128 [%r1, %rd1, %rd2], [%rd3], %r2, [%rd4], %rd5; +; CHECK-NEXT: ret; +; FORMAT-LABEL: define void @test_fabric_try_put_counted_multimem_cache( +; FORMAT: call void @llvm.nvvm.fabric.try_put.counted_writes(ptr addrspace(8) %handle_data, ptr addrspace(8) %handle_counted, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i64 %cacheHint, /* flag_cache_hint= */ i1 true, /* flag_multimem= */ i32 1) + %handle_data = call ptr addrspace(8) @llvm.nvvm.fabric.handle_pair(i32 %leId, i64 %dataOffset) + %handle_counted = call ptr addrspace(8) @llvm.nvvm.fabric.handle_pair(i32 %leId, i64 %counterOffset) + call void @llvm.nvvm.fabric.try_put.counted_writes(ptr addrspace(8) %handle_data, ptr addrspace(8) %handle_counted, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i64 %cacheHint, i1 true, i32 1) + ret void +} + +define void @test_fabric_try_put_cache_hint_ignored(i32 %leId, i64 %offset, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i64 %cacheHint) { +; CHECK-LABEL: test_fabric_try_put_cache_hint_ignored( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-NEXT: .reg .b64 %rd<4>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [test_fabric_try_put_cache_hint_ignored_param_0]; +; CHECK-NEXT: ld.param::func.b64 %rd1, [test_fabric_try_put_cache_hint_ignored_param_1]; +; CHECK-NEXT: ld.param::func.b64 %rd2, [test_fabric_try_put_cache_hint_ignored_param_2]; +; CHECK-NEXT: ld.param::func.b64 %rd3, [test_fabric_try_put_cache_hint_ignored_param_3]; +; CHECK-NEXT: ld.param::func.b32 %r2, [test_fabric_try_put_cache_hint_ignored_param_4]; +; CHECK-NEXT: fabric.try_put.async.shared::cta.mbarrier::complete_tx::16B.mbarrier::report::fabric.relaxed.sys.b128 [%r1, %rd1], [%rd2], %r2, [%rd3]; +; CHECK-NEXT: ret; +; FORMAT-LABEL: define void @test_fabric_try_put_cache_hint_ignored( +; FORMAT: call void @llvm.nvvm.fabric.try_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 0, i64 %cacheHint, /* flag_cache_hint= */ i1 false, /* flag_cp_mask= */ i1 false, /* flag_multimem= */ i32 0) + %handle = call ptr addrspace(8) @llvm.nvvm.fabric.handle_pair(i32 %leId, i64 %offset) + call void @llvm.nvvm.fabric.try_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 0, i64 %cacheHint, i1 false, i1 false, i32 0) + ret void +} diff --git a/llvm/test/CodeGen/NVPTX/fabric-try-put-memory-effects.ll b/llvm/test/CodeGen/NVPTX/fabric-try-put-memory-effects.ll new file mode 100644 index 00000000000000..9ec240baae5459 --- /dev/null +++ b/llvm/test/CodeGen/NVPTX/fabric-try-put-memory-effects.ll @@ -0,0 +1,100 @@ +; Test the IR memory-effects contract of the fabric put intrinsics. +; +; RUN: opt -passes=aa-eval -aa-pipeline=nvptx-aa,basic-aa -print-all-alias-modref-info -disable-output < %s 2>&1 | FileCheck %s --check-prefix=AA +; RUN: opt -passes=aa-eval -aa-pipeline=basic-aa,nvptx-aa -print-all-alias-modref-info -disable-output < %s 2>&1 | FileCheck %s --check-prefix=AA +; RUN: opt -passes=gvn -aa-pipeline=nvptx-aa,basic-aa -S < %s | FileCheck %s --check-prefix=GVN +; RUN: opt -passes=gvn -aa-pipeline=basic-aa,nvptx-aa -S < %s | FileCheck %s --check-prefix=GVN +; RUN: llvm-as < %s | llvm-dis | FileCheck %s --check-prefix=ATTR + +target triple = "nvptx64-nvidia-cuda" +target datalayout = "e-p6:32:32-p8:128:128-ni:8-i64:64-i128:128-i256:256-v16:16-v32:32-n16:32:64" + +declare ptr addrspace(8) @llvm.nvvm.fabric.handle_pair(i32, i64) +declare void @llvm.nvvm.fabric.try_put(ptr addrspace(8), ptr addrspace(3), ptr addrspace(3), i32, i16, i64, i1, i1, i32) +declare void @llvm.nvvm.fabric.try_put.counted_writes(ptr addrspace(8), ptr addrspace(8), ptr addrspace(3), ptr addrspace(3), i32, i64, i1, i32) + +define i32 @put_effects(ptr addrspace(1) %global, i32 %id, i64 %offset, + ptr addrspace(3) %src, ptr addrspace(3) %bar) { +; AA-LABEL: Function: put_effects: +; The constructor stays memory-free. +; AA-DAG: NoModRef: Ptr: {{.*}}%global{{[[:space:]]*}}<->{{[[:space:]]*}}{{.*}}@llvm.nvvm.fabric.handle_pair +; The put may access global memory through the handle. +; AA-DAG: Both ModRef: Ptr: {{.*}}%global{{[[:space:]]*}}<->{{[[:space:]]*}}call void @llvm.nvvm.fabric.try_put( +; GVN-LABEL: define i32 @put_effects( +; GVN: %before = load i32, ptr addrspace(1) %global +; GVN: %after = load i32, ptr addrspace(1) %global +; GVN: %difference = sub i32 %after, %before + %handle = call ptr addrspace(8) @llvm.nvvm.fabric.handle_pair(i32 %id, i64 %offset) + %before = load i32, ptr addrspace(1) %global + call void @llvm.nvvm.fabric.try_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %bar, + i32 16, i16 0, i64 0, i1 false, i1 false, i32 0) + %after = load i32, ptr addrspace(1) %global + %difference = sub i32 %after, %before + ret i32 %difference +} + +; Both counted handles share one endpoint ID and use distinct offsets. +define i32 @counted_put_effects(ptr addrspace(1) %global, i32 %id, + i64 %data_offset, i64 %counter_offset, + ptr addrspace(3) %src, ptr addrspace(3) %bar) { +; AA-LABEL: Function: counted_put_effects: +; AA-DAG: NoModRef: Ptr: {{.*}}%global{{[[:space:]]*}}<->{{[[:space:]]*}}{{.*}}@llvm.nvvm.fabric.handle_pair +; AA-DAG: Both ModRef: Ptr: {{.*}}%global{{[[:space:]]*}}<->{{[[:space:]]*}}call void @llvm.nvvm.fabric.try_put.counted_writes( +; GVN-LABEL: define i32 @counted_put_effects( +; GVN: %before = load i32, ptr addrspace(1) %global +; GVN: %after = load i32, ptr addrspace(1) %global +; GVN: %difference = sub i32 %after, %before + %data_handle = call ptr addrspace(8) @llvm.nvvm.fabric.handle_pair(i32 %id, i64 %data_offset) + %counter_handle = call ptr addrspace(8) @llvm.nvvm.fabric.handle_pair(i32 %id, i64 %counter_offset) + %before = load i32, ptr addrspace(1) %global + call void @llvm.nvvm.fabric.try_put.counted_writes(ptr addrspace(8) %data_handle, ptr addrspace(8) %counter_handle, + ptr addrspace(3) %src, ptr addrspace(3) %bar, i32 16, i64 0, i1 false, i32 0) + %after = load i32, ptr addrspace(1) %global + %difference = sub i32 %after, %before + ret i32 %difference +} + +; The handle arrives as an argument; the resource access is based on it. +define i32 @put_handle_argument_effects(ptr addrspace(8) %handle, + ptr addrspace(1) %global, + ptr addrspace(3) %src, + ptr addrspace(3) %bar) { +; AA-LABEL: Function: put_handle_argument_effects: +; AA: Both ModRef: Ptr: {{.*}}%global{{[[:space:]]*}}<->{{[[:space:]]*}}call void @llvm.nvvm.fabric.try_put( +; GVN-LABEL: define i32 @put_handle_argument_effects( +; GVN: %before = load i32, ptr addrspace(1) %global +; GVN: %after = load i32, ptr addrspace(1) %global +; GVN: %difference = sub i32 %after, %before + %before = load i32, ptr addrspace(1) %global + call void @llvm.nvvm.fabric.try_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %bar, i32 16, + i16 0, i64 0, i1 false, i1 false, i32 0) + %after = load i32, ptr addrspace(1) %global + %difference = sub i32 %after, %before + ret i32 %difference +} + +; A locally constructed handle versus an ordinary generic (AS0) pointer. +define i32 @put_generic_effects(ptr %generic, i32 %id, i64 %offset, + ptr addrspace(3) %src, ptr addrspace(3) %bar) { +; AA-LABEL: Function: put_generic_effects: +; AA-DAG: NoModRef: Ptr: {{.*}}%generic{{[[:space:]]*}}<->{{[[:space:]]*}}{{.*}}@llvm.nvvm.fabric.handle_pair +; AA-DAG: Both ModRef: Ptr: {{.*}}%generic{{[[:space:]]*}}<->{{[[:space:]]*}}call void @llvm.nvvm.fabric.try_put( +; GVN-LABEL: define i32 @put_generic_effects( +; GVN: %before = load i32, ptr %generic +; GVN: %after = load i32, ptr %generic +; GVN: %difference = sub i32 %after, %before + %handle = call ptr addrspace(8) @llvm.nvvm.fabric.handle_pair(i32 %id, i64 %offset) + %before = load i32, ptr %generic + call void @llvm.nvvm.fabric.try_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %bar, i32 16, + i16 0, i64 0, i1 false, i1 false, i32 0) + %after = load i32, ptr %generic + %difference = sub i32 %after, %before + ret i32 %difference +} + +; ATTR: declare ptr addrspace(8) @llvm.nvvm.fabric.handle_pair(i32, i64) #[[HANDLE:[0-9]+]] +; The put declarations are argument-memory-only: their resource accesses are modeled through the handle arguments. +; ATTR: declare void @llvm.nvvm.fabric.try_put(ptr addrspace(8), ptr addrspace(3) readonly, ptr addrspace(3), i32, i16, i64, i1 immarg, i1 immarg, i32 immarg range(i32 0, 2)) #[[PUT:[0-9]+]] +; ATTR: declare void @llvm.nvvm.fabric.try_put.counted_writes(ptr addrspace(8), ptr addrspace(8), ptr addrspace(3) readonly, ptr addrspace(3), i32, i64, i1 immarg, i32 immarg range(i32 0, 2)) #[[PUT]] +; ATTR: attributes #[[HANDLE]] = { {{.*}}memory(none){{.*}} } +; ATTR: attributes #[[PUT]] = { {{.*}}memory(argmem: readwrite){{.*}} } diff --git a/llvm/test/CodeGen/NVPTX/fabric-try-put.ll b/llvm/test/CodeGen/NVPTX/fabric-try-put.ll new file mode 100644 index 00000000000000..b0e02a444f6748 --- /dev/null +++ b/llvm/test/CodeGen/NVPTX/fabric-try-put.ll @@ -0,0 +1,262 @@ +; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6 +; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx93 | FileCheck %s +; RUN: llvm-as < %s | llvm-dis | FileCheck --check-prefix=FORMAT %s +; RUN: %if ptxas-sm_100 && ptxas-isa-9.3 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx93 | %ptxas-verify -arch=sm_100 %} + +target triple = "nvptx64-nvidia-cuda" + +declare i32 @llvm.nvvm.read.ptx.sreg.cluster.ctarank() +declare ptr addrspace(8) @llvm.nvvm.fabric.handle_pair(i32, i64) +declare void @llvm.nvvm.fabric.try_put(ptr addrspace(8), ptr addrspace(3), ptr addrspace(3), i32, i16, i64, i1, i1, i32) +declare void @llvm.nvvm.fabric.try_put.counted_writes(ptr addrspace(8), ptr addrspace(8), ptr addrspace(3), ptr addrspace(3), i32, i64, i1, i32) + +define void @test_fabric_try_put(i32 %leId, i64 %offset, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size) { +; CHECK-LABEL: test_fabric_try_put( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-NEXT: .reg .b64 %rd<4>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [test_fabric_try_put_param_0]; +; CHECK-NEXT: ld.param::func.b64 %rd1, [test_fabric_try_put_param_1]; +; CHECK-NEXT: ld.param::func.b64 %rd2, [test_fabric_try_put_param_2]; +; CHECK-NEXT: ld.param::func.b64 %rd3, [test_fabric_try_put_param_3]; +; CHECK-NEXT: ld.param::func.b32 %r2, [test_fabric_try_put_param_4]; +; CHECK-NEXT: fabric.try_put.async.shared::cta.mbarrier::complete_tx::16B.mbarrier::report::fabric.relaxed.sys.b128 [%r1, %rd1], [%rd2], %r2, [%rd3]; +; CHECK-NEXT: ret; +; FORMAT-LABEL: define void @test_fabric_try_put( +; FORMAT: call void @llvm.nvvm.fabric.try_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 0, i64 0, /* flag_cache_hint= */ i1 false, /* flag_cp_mask= */ i1 false, /* flag_multimem= */ i32 0) + %handle = call ptr addrspace(8) @llvm.nvvm.fabric.handle_pair(i32 %leId, i64 %offset) + call void @llvm.nvvm.fabric.try_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 0, i64 0, i1 false, i1 false, i32 0) + ret void +} + +define void @test_fabric_try_put_mask(i32 %leId, i64 %offset, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 %mask) { +; CHECK-LABEL: test_fabric_try_put_mask( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-NEXT: .reg .b64 %rd<4>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [test_fabric_try_put_mask_param_0]; +; CHECK-NEXT: ld.param::func.b64 %rd1, [test_fabric_try_put_mask_param_1]; +; CHECK-NEXT: ld.param::func.b64 %rd2, [test_fabric_try_put_mask_param_2]; +; CHECK-NEXT: ld.param::func.b64 %rd3, [test_fabric_try_put_mask_param_3]; +; CHECK-NEXT: ld.param::func.b32 %r2, [test_fabric_try_put_mask_param_4]; +; CHECK-NEXT: ld.param::func.b16 %rs1, [test_fabric_try_put_mask_param_5]; +; CHECK-NEXT: fabric.try_put.async.shared::cta.mbarrier::complete_tx::16B.mbarrier::report::fabric.cp_mask.relaxed.sys.b128 [%r1, %rd1], [%rd2], %r2, [%rd3], %rs1; +; CHECK-NEXT: ret; +; FORMAT-LABEL: define void @test_fabric_try_put_mask( +; FORMAT: call void @llvm.nvvm.fabric.try_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 %mask, i64 0, /* flag_cache_hint= */ i1 false, /* flag_cp_mask= */ i1 true, /* flag_multimem= */ i32 0) + %handle = call ptr addrspace(8) @llvm.nvvm.fabric.handle_pair(i32 %leId, i64 %offset) + call void @llvm.nvvm.fabric.try_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 %mask, i64 0, i1 false, i1 true, i32 0) + ret void +} + +define void @test_fabric_try_put_byte_mask_ignored(i32 %leId, i64 %offset, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 %mask) { +; CHECK-LABEL: test_fabric_try_put_byte_mask_ignored( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-NEXT: .reg .b64 %rd<4>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [test_fabric_try_put_byte_mask_ignored_param_0]; +; CHECK-NEXT: ld.param::func.b64 %rd1, [test_fabric_try_put_byte_mask_ignored_param_1]; +; CHECK-NEXT: ld.param::func.b64 %rd2, [test_fabric_try_put_byte_mask_ignored_param_2]; +; CHECK-NEXT: ld.param::func.b64 %rd3, [test_fabric_try_put_byte_mask_ignored_param_3]; +; CHECK-NEXT: ld.param::func.b32 %r2, [test_fabric_try_put_byte_mask_ignored_param_4]; +; CHECK-NEXT: fabric.try_put.async.shared::cta.mbarrier::complete_tx::16B.mbarrier::report::fabric.relaxed.sys.b128 [%r1, %rd1], [%rd2], %r2, [%rd3]; +; CHECK-NEXT: ret; +; FORMAT-LABEL: define void @test_fabric_try_put_byte_mask_ignored( +; FORMAT: call void @llvm.nvvm.fabric.try_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 %mask, i64 0, /* flag_cache_hint= */ i1 false, /* flag_cp_mask= */ i1 false, /* flag_multimem= */ i32 0) + %handle = call ptr addrspace(8) @llvm.nvvm.fabric.handle_pair(i32 %leId, i64 %offset) + call void @llvm.nvvm.fabric.try_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 %mask, i64 0, i1 false, i1 false, i32 0) + ret void +} + +; Verify that the cross-BB case keeps the harmless cvt.u32.u64 for the leId +; operand (the handle i128 is built in the entry block and consumed in the +; conditional successor). +define void @test_fabric_try_put_cross_bb(i64 %offset, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i1 %cond) { +; CHECK-LABEL: test_fabric_try_put_cross_bb( +; CHECK: { +; CHECK-NEXT: .reg .pred %p<3>; +; CHECK-NEXT: .reg .b16 %rs<3>; +; CHECK-NEXT: .reg .b32 %r<4>; +; CHECK-NEXT: .reg .b64 %rd<5>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: // %entry +; CHECK-NEXT: ld.param::func.b8 %rs1, [test_fabric_try_put_cross_bb_param_4]; +; CHECK-NEXT: and.b16 %rs2, %rs1, 1; +; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0; +; CHECK-NEXT: not.pred %p2, %p1; +; CHECK-NEXT: @%p2 bra $L__BB3_2; +; CHECK-NEXT: // %bb.1: // %do_put +; CHECK-NEXT: ld.param::func.b32 %r1, [test_fabric_try_put_cross_bb_param_3]; +; CHECK-NEXT: ld.param::func.b64 %rd4, [test_fabric_try_put_cross_bb_param_2]; +; CHECK-NEXT: ld.param::func.b64 %rd3, [test_fabric_try_put_cross_bb_param_1]; +; CHECK-NEXT: ld.param::func.b64 %rd2, [test_fabric_try_put_cross_bb_param_0]; +; CHECK-NEXT: mov.u32 %r2, %cluster_ctarank; +; CHECK-NEXT: cvt.u64.u32 %rd1, %r2; +; CHECK-NEXT: cvt.u32.u64 %r3, %rd1; +; CHECK-NEXT: fabric.try_put.async.shared::cta.mbarrier::complete_tx::16B.mbarrier::report::fabric.relaxed.sys.b128 [%r3, %rd2], [%rd3], %r1, [%rd4]; +; CHECK-NEXT: $L__BB3_2: // %exit +; CHECK-NEXT: ret; +; FORMAT-LABEL: define void @test_fabric_try_put_cross_bb( +; FORMAT: call void @llvm.nvvm.fabric.try_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 0, i64 0, /* flag_cache_hint= */ i1 false, /* flag_cp_mask= */ i1 false, /* flag_multimem= */ i32 0) +entry: + %leId = call i32 @llvm.nvvm.read.ptx.sreg.cluster.ctarank() + %handle = call ptr addrspace(8) @llvm.nvvm.fabric.handle_pair(i32 %leId, i64 %offset) + br i1 %cond, label %do_put, label %exit + +do_put: + call void @llvm.nvvm.fabric.try_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 0, i64 0, i1 false, i1 false, i32 0) + br label %exit + +exit: + ret void +} + +; Handle created in kernel, passed to device function which issues the try_put. +; The i128 handle is passed as a 16-byte param across the call boundary. +define ptx_device void @device_do_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size) { +; CHECK-LABEL: device_do_put( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-NEXT: .reg .b64 %rd<5>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.v2.b64 {%rd1, %rd2}, [device_do_put_param_0]; +; CHECK-NEXT: cvt.u32.u64 %r1, %rd1; +; CHECK-NEXT: ld.param::func.b64 %rd3, [device_do_put_param_1]; +; CHECK-NEXT: ld.param::func.b64 %rd4, [device_do_put_param_2]; +; CHECK-NEXT: ld.param::func.b32 %r2, [device_do_put_param_3]; +; CHECK-NEXT: fabric.try_put.async.shared::cta.mbarrier::complete_tx::16B.mbarrier::report::fabric.relaxed.sys.b128 [%r1, %rd2], [%rd3], %r2, [%rd4]; +; CHECK-NEXT: ret; +; FORMAT-LABEL: define ptx_device void @device_do_put( +; FORMAT: call void @llvm.nvvm.fabric.try_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 0, i64 0, /* flag_cache_hint= */ i1 false, /* flag_cp_mask= */ i1 false, /* flag_multimem= */ i32 0) + call void @llvm.nvvm.fabric.try_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 0, i64 0, i1 false, i1 false, i32 0) + ret void +} + +define ptx_kernel void @kernel_caller(i64 %offset, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size) { +; CHECK-LABEL: kernel_caller( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-NEXT: .reg .b64 %rd<5>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::entry.b64 %rd1, [kernel_caller_param_0]; +; CHECK-NEXT: ld.param::entry.b64 %rd2, [kernel_caller_param_1]; +; CHECK-NEXT: mov.u32 %r1, %cluster_ctarank; +; CHECK-NEXT: cvt.u64.u32 %rd3, %r1; +; CHECK-NEXT: ld.param::entry.b64 %rd4, [kernel_caller_param_2]; +; CHECK-NEXT: { // callseq 0, 0 +; CHECK-NEXT: .param .align 16 .b8 param0[16]; +; CHECK-NEXT: .param .b64 param1; +; CHECK-NEXT: .param .b64 param2; +; CHECK-NEXT: .param .b32 param3; +; CHECK-NEXT: st.param::func.v2.b64 [param0], {%rd3, %rd1}; +; CHECK-NEXT: ld.param::entry.b32 %r2, [kernel_caller_param_3]; +; CHECK-NEXT: st.param::func.b32 [param3], %r2; +; CHECK-NEXT: st.param::func.b64 [param2], %rd4; +; CHECK-NEXT: st.param::func.b64 [param1], %rd2; +; CHECK-NEXT: call.uni device_do_put, (param0, param1, param2, param3); +; CHECK-NEXT: } // callseq 0 +; CHECK-NEXT: ret; + %leId = call i32 @llvm.nvvm.read.ptx.sreg.cluster.ctarank() + %handle = call ptr addrspace(8) @llvm.nvvm.fabric.handle_pair(i32 %leId, i64 %offset) + call void @device_do_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size) + ret void +} + +define void @test_fabric_try_put_counted_writes(i32 %leId0, i64 %offset0, i32 %leId1, i64 %offset1, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size) { +; CHECK-LABEL: test_fabric_try_put_counted_writes( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-NEXT: .reg .b64 %rd<5>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [test_fabric_try_put_counted_writes_param_0]; +; CHECK-NEXT: ld.param::func.b64 %rd1, [test_fabric_try_put_counted_writes_param_1]; +; CHECK-NEXT: ld.param::func.b64 %rd2, [test_fabric_try_put_counted_writes_param_3]; +; CHECK-NEXT: ld.param::func.b64 %rd3, [test_fabric_try_put_counted_writes_param_4]; +; CHECK-NEXT: ld.param::func.b64 %rd4, [test_fabric_try_put_counted_writes_param_5]; +; CHECK-NEXT: ld.param::func.b32 %r2, [test_fabric_try_put_counted_writes_param_6]; +; CHECK-NEXT: fabric.try_put.async.shared::cta.mbarrier::complete_tx::16B.mbarrier::report::fabric.counted::bytes.relaxed.sys.b128 [%r1, %rd1, %rd2], [%rd3], %r2, [%rd4]; +; CHECK-NEXT: ret; +; FORMAT-LABEL: define void @test_fabric_try_put_counted_writes( +; FORMAT: call void @llvm.nvvm.fabric.try_put.counted_writes(ptr addrspace(8) %handle_data, ptr addrspace(8) %handle_counted, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i64 0, /* flag_cache_hint= */ i1 false, /* flag_multimem= */ i32 0) + %handle_data = call ptr addrspace(8) @llvm.nvvm.fabric.handle_pair(i32 %leId0, i64 %offset0) + %handle_counted = call ptr addrspace(8) @llvm.nvvm.fabric.handle_pair(i32 %leId1, i64 %offset1) + call void @llvm.nvvm.fabric.try_put.counted_writes(ptr addrspace(8) %handle_data, ptr addrspace(8) %handle_counted, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i64 0, i1 false, i32 0) + ret void +} + +define void @test_fabric_try_put_multimem(i32 %leId, i64 %offset, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size) { +; CHECK-LABEL: test_fabric_try_put_multimem( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-NEXT: .reg .b64 %rd<4>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [test_fabric_try_put_multimem_param_0]; +; CHECK-NEXT: ld.param::func.b64 %rd1, [test_fabric_try_put_multimem_param_1]; +; CHECK-NEXT: ld.param::func.b64 %rd2, [test_fabric_try_put_multimem_param_2]; +; CHECK-NEXT: ld.param::func.b64 %rd3, [test_fabric_try_put_multimem_param_3]; +; CHECK-NEXT: ld.param::func.b32 %r2, [test_fabric_try_put_multimem_param_4]; +; CHECK-NEXT: fabric.try_put.async.multimem.shared::cta.mbarrier::complete_tx::16B.mbarrier::report::fabric.relaxed.sys.b128 [%r1, %rd1], [%rd2], %r2, [%rd3]; +; CHECK-NEXT: ret; +; FORMAT-LABEL: define void @test_fabric_try_put_multimem( +; FORMAT: call void @llvm.nvvm.fabric.try_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 0, i64 0, /* flag_cache_hint= */ i1 false, /* flag_cp_mask= */ i1 false, /* flag_multimem= */ i32 1) + %handle = call ptr addrspace(8) @llvm.nvvm.fabric.handle_pair(i32 %leId, i64 %offset) + call void @llvm.nvvm.fabric.try_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 0, i64 0, i1 false, i1 false, i32 1) + ret void +} + +define void @test_fabric_try_put_mask_multimem(i32 %leId, i64 %offset, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 %mask) { +; CHECK-LABEL: test_fabric_try_put_mask_multimem( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-NEXT: .reg .b64 %rd<4>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [test_fabric_try_put_mask_multimem_param_0]; +; CHECK-NEXT: ld.param::func.b64 %rd1, [test_fabric_try_put_mask_multimem_param_1]; +; CHECK-NEXT: ld.param::func.b64 %rd2, [test_fabric_try_put_mask_multimem_param_2]; +; CHECK-NEXT: ld.param::func.b64 %rd3, [test_fabric_try_put_mask_multimem_param_3]; +; CHECK-NEXT: ld.param::func.b32 %r2, [test_fabric_try_put_mask_multimem_param_4]; +; CHECK-NEXT: ld.param::func.b16 %rs1, [test_fabric_try_put_mask_multimem_param_5]; +; CHECK-NEXT: fabric.try_put.async.multimem.shared::cta.mbarrier::complete_tx::16B.mbarrier::report::fabric.cp_mask.relaxed.sys.b128 [%r1, %rd1], [%rd2], %r2, [%rd3], %rs1; +; CHECK-NEXT: ret; +; FORMAT-LABEL: define void @test_fabric_try_put_mask_multimem( +; FORMAT: call void @llvm.nvvm.fabric.try_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 %mask, i64 0, /* flag_cache_hint= */ i1 false, /* flag_cp_mask= */ i1 true, /* flag_multimem= */ i32 1) + %handle = call ptr addrspace(8) @llvm.nvvm.fabric.handle_pair(i32 %leId, i64 %offset) + call void @llvm.nvvm.fabric.try_put(ptr addrspace(8) %handle, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i16 %mask, i64 0, i1 false, i1 true, i32 1) + ret void +} + +define void @test_fabric_try_put_counted_writes_multimem(i32 %leId0, i64 %offset0, i32 %leId1, i64 %offset1, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size) { +; CHECK-LABEL: test_fabric_try_put_counted_writes_multimem( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-NEXT: .reg .b64 %rd<5>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [test_fabric_try_put_counted_writes_multimem_param_0]; +; CHECK-NEXT: ld.param::func.b64 %rd1, [test_fabric_try_put_counted_writes_multimem_param_1]; +; CHECK-NEXT: ld.param::func.b64 %rd2, [test_fabric_try_put_counted_writes_multimem_param_3]; +; CHECK-NEXT: ld.param::func.b64 %rd3, [test_fabric_try_put_counted_writes_multimem_param_4]; +; CHECK-NEXT: ld.param::func.b64 %rd4, [test_fabric_try_put_counted_writes_multimem_param_5]; +; CHECK-NEXT: ld.param::func.b32 %r2, [test_fabric_try_put_counted_writes_multimem_param_6]; +; CHECK-NEXT: fabric.try_put.async.multimem.shared::cta.mbarrier::complete_tx::16B.mbarrier::report::fabric.counted::bytes.relaxed.sys.b128 [%r1, %rd1, %rd2], [%rd3], %r2, [%rd4]; +; CHECK-NEXT: ret; +; FORMAT-LABEL: define void @test_fabric_try_put_counted_writes_multimem( +; FORMAT: call void @llvm.nvvm.fabric.try_put.counted_writes(ptr addrspace(8) %handle_data, ptr addrspace(8) %handle_counted, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i64 0, /* flag_cache_hint= */ i1 false, /* flag_multimem= */ i32 1) + %handle_data = call ptr addrspace(8) @llvm.nvvm.fabric.handle_pair(i32 %leId0, i64 %offset0) + %handle_counted = call ptr addrspace(8) @llvm.nvvm.fabric.handle_pair(i32 %leId1, i64 %offset1) + call void @llvm.nvvm.fabric.try_put.counted_writes(ptr addrspace(8) %handle_data, ptr addrspace(8) %handle_counted, ptr addrspace(3) %src, ptr addrspace(3) %mbar, i32 %size, i64 0, i1 false, i32 1) + ret void +} diff --git a/llvm/unittests/TargetParser/TripleTest.cpp b/llvm/unittests/TargetParser/TripleTest.cpp index cc19b14827db56..81ece71eb5c7c3 100644 --- a/llvm/unittests/TargetParser/TripleTest.cpp +++ b/llvm/unittests/TargetParser/TripleTest.cpp @@ -4036,22 +4036,22 @@ TEST(DataLayoutTest, NVPTX) { return Specs; }; - // The 32-bit target uses a single 32-bit pointer specification and is - // unaffected by the ABI name. + // Fabric handles (addrspace:8) are 128-bit and non-integral. EXPECT_THAT(PointerLayoutSpecs(TT32.computeDataLayout("")), - testing::ElementsAre("p:32:32")); + testing::ElementsAre("p:32:32", "p8:128:128")); EXPECT_THAT(PointerLayoutSpecs(TT32.computeDataLayout("shortptr")), - testing::ElementsAre("p:32:32")); + testing::ElementsAre("p:32:32", "p8:128:128")); // The default 64-bit target only shrinks Tensor Memory (addrspace:6). EXPECT_THAT(PointerLayoutSpecs(TT64.computeDataLayout("")), - testing::ElementsAre("p6:32:32")); + testing::ElementsAre("p6:32:32", "p8:128:128")); // In shortptr mode the extra address spaces become 32-bit. The pointer // specifications must remain sorted by address space. EXPECT_THAT(PointerLayoutSpecs(TT64.computeDataLayout("shortptr")), testing::ElementsAre("p3:32:32", "p4:32:32", "p5:32:32", - "p6:32:32", "p7:32:32", "p101:32:32")); + "p6:32:32", "p7:32:32", "p8:128:128", + "p101:32:32")); } TEST(DataLayoutTest, CheriRISCV32) { _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
