[llvm] [mlir] [LLVM][NVPTX] Add async bulk copy global to shared extensions (PR #222323)
via llvm-commits
llvm-commits at lists.llvm.org
Wed Sep 9 05:57:23 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-mlir
Author: Rajat Bajpai (rajatbajpai)
<details>
<summary>Changes</summary>
This change adds following things to bulk copy intrinsics.
1. Relaxed memory ordering semantics with a scope argument.
2. Data-validity reporting patterns (introduced in Rubin).
3. 32-bit multicast mask for global to shared::cluster variants (introduced in Rubin).
4. Ignore out of bound checks for global to shared::cta variants.
Note: MLIR lowering is updated to emit the new intrinsic signatures. Support for the new features in MLIR will be done in a separate change.
---
Patch is 85.85 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/222323.diff
23 Files Affected:
- (modified) llvm/docs/NVPTXUsage.md (+138-18)
- (modified) llvm/include/llvm/IR/IntrinsicsNVVM.td (+67-13)
- (modified) llvm/include/llvm/IR/NVVMIntrinsicUtils.h (+24)
- (modified) llvm/lib/IR/AutoUpgrade.cpp (+105-3)
- (modified) llvm/lib/IR/NVVMIntrinsicUtils.cpp (+10)
- (modified) llvm/lib/IR/Verifier.cpp (+14)
- (modified) llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp (+6)
- (modified) llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.h (+2)
- (modified) llvm/lib/Target/NVPTX/NVPTXInstrInfo.td (+10)
- (modified) llvm/lib/Target/NVPTX/NVPTXIntrinsics.td (+138-51)
- (modified) llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll (+7-2)
- (added) llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-cta-invalid.ll (+30)
- (added) llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-cta-oob.ll (+27)
- (added) llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-cta-relaxed.ll (+75)
- (added) llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-cta-validate-data.ll (+56)
- (added) llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-invalid.ll (+31)
- (added) llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-mc32.ll (+52)
- (added) llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-relaxed.ll (+59)
- (added) llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-validate-data.ll (+58)
- (added) llvm/test/Verifier/NVPTX/cp-async-bulk-g2s-cta.ll (+16)
- (modified) mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td (+1-1)
- (modified) mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp (+11-3)
- (modified) mlir/test/Target/LLVMIR/nvvm/tma_bulk_copy.mlir (+6-6)
``````````diff
diff --git a/llvm/docs/NVPTXUsage.md b/llvm/docs/NVPTXUsage.md
index 6e0ca158266ca..55f75bd7c2f89 100644
--- a/llvm/docs/NVPTXUsage.md
+++ b/llvm/docs/NVPTXUsage.md
@@ -1729,8 +1729,7 @@ For more information, refer to the [PTX ISA](https://docs.nvidia.com/cuda/parall
#### TMA Global-to-Shared Data-Validity Reporting
The '`@llvm.nvvm.cp.async.bulk.tensor.g2s.*`' and '`@llvm.nvvm.cp.async.bulk.tensor.g2s.cta.*`' intrinsics, including the property-override
-forms, take an `i32 %validate_pattern` immediate argument in the range \[0, 6). It selects the data-validity reporting mechanism applied to the
-tensor data while it is copied:
+forms, take an `i32 %validate_pattern` immediate argument in the range \[0, 6). The non-tensor '`@llvm.nvvm.cp.async.bulk.global.to.shared.cluster`', '`@llvm.nvvm.cp.async.bulk.global.to.shared.cluster.relaxed`', '`@llvm.nvvm.cp.async.bulk.global.to.shared.cta`', and '`@llvm.nvvm.cp.async.bulk.global.to.shared.cta.relaxed`' intrinsics also take the same `i32 %validate_pattern` immediate argument. It selects the data-validity reporting mechanism applied to the data while it is copied:
| Value | PTX Report Mechanism | Semantics |
|-------|---------------------------------------------------------|------------------------------------------------------------------|
@@ -1752,7 +1751,8 @@ For more information, refer to the [PTX ISA](https://docs.nvidia.com/cuda/parall
##### Syntax:
```llvm
-declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster(ptr addrspace(7) %dst, ptr addrspace(3) %mbar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 %flag_mc, i1 %flag_ch)
+declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.i16(ptr addrspace(7) %dst, ptr addrspace(3) %mbar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 %flag_mc, i1 %flag_ch, i32 %validate_pattern)
+declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.i32(..., i32 %mc, i64 %ch, i1 %flag_mc, i1 %flag_ch, i32 %validate_pattern)
```
##### Overview:
@@ -1763,16 +1763,71 @@ instructions. These instructions initiate an asynchronous copy of bulk data from
global memory to shared::cluster memory. The 32-bit operand `%size` specifies
the amount of memory to be copied and it must be a multiple of 16.
-- The last two arguments to these intrinsics are boolean flags indicating
- support for cache_hint and/or multicast modifiers. These flag arguments must
- be compile-time constants. The backend looks through these flags and lowers
- the intrinsics appropriately.
-- The Nth argument (denoted by `i1 %flag_ch`) when set, indicates a valid
- cache_hint (`i64 %ch`) and generates the `.L2::cache_hint` variant of the
- PTX instruction.
-- The [N-1]th argument (denoted by `i1 %flag_mc`) when set, indicates the
- presence of a multicast mask (`i16 %mc`) and generates the PTX instruction
- with the `.multicast::cluster` modifier.
+- The trailing `%flag_mc`, `%flag_ch`, and `%validate_pattern` arguments control
+ the multicast and cache-hint modifiers and the data-validity reporting
+ pattern. They must be compile-time constants.
+- The argument denoted by `i1 %flag_ch`, when set, indicates a valid cache_hint
+ (`i64 %ch`) and generates the `.L2::cache_hint` variant of the PTX
+ instruction.
+- The argument denoted by `i1 %flag_mc`, when set, indicates the presence of a
+ multicast mask (`%mc`) and generates the PTX instruction with the
+ `.multicast::cluster` modifier. An `i16` mask selects the default/16-bit
+ multicast form while an `i32` mask selects the `.multicast::cluster::32b`
+ variant. Only `i16` and `i32` mask types are supported. The 32-bit form
+ requires `sm_107f` or later in the same family and is supported only on PTX
+ 9.4 or later.
+- The trailing `i32 %validate_pattern` flag selects the data-validity reporting
+ mechanism; see
+ [TMA Global-to-Shared Data-Validity Reporting](#tma-global-to-shared-data-validity-reporting).
+
+For more information, refer [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk).
+
+#### '`llvm.nvvm.cp.async.bulk.global.to.shared.cluster.relaxed`'
+
+##### Syntax:
+
+```llvm
+declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.relaxed.i16(ptr addrspace(7) %dst, ptr addrspace(3) %mbar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 %flag_mc, i1 %flag_ch, i32 %scope, i32 %validate_pattern)
+declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.relaxed.i32(..., i32 %mc, i64 %ch, i1 %flag_mc, i1 %flag_ch, i32 %scope, i32 %validate_pattern)
+```
+
+##### Overview:
+
+The '`@llvm.nvvm.cp.async.bulk.global.to.shared.cluster.relaxed`' intrinsic
+corresponds to the
+`cp.async.bulk.relaxed.<scope>.shared::cluster.global.mbarrier::complete_tx::bytes{...}.b128`
+family of PTX instructions. These instructions initiate an asynchronous copy of bulk data from
+global memory to shared::cluster memory, like the weak variant above. Unlike
+the weak variant, the copy accesses memory with naturally-aligned strong
+memory operations with element-wise atomic size specified by the `.b128` type (implicit) and
+thread scope specified by `%scope`; the complete-tx operation on the mbarrier has `.release` semantics at the
+`.cluster` scope.
+
+- The trailing `%flag_mc`, `%flag_ch`, `%scope`, and `%validate_pattern`
+ arguments control the multicast and cache-hint modifiers, the scope of the
+ relaxed memory ordering semantics, and the data-validity reporting pattern.
+ They must be compile-time constants.
+- The argument denoted by `i1 %flag_ch`, when set, indicates a valid cache_hint
+ (`i64 %ch`) and generates the `.L2::cache_hint` variant of the PTX
+ instruction.
+- The argument denoted by `i1 %flag_mc`, when set, indicates the presence of a
+ multicast mask (`%mc`) and generates the PTX instruction with the
+ `.multicast::cluster` modifier. An `i16` mask selects the default/16-bit
+ multicast form while an `i32` mask selects the `.multicast::cluster::32b`
+ variant. Only `i16` and `i32` mask types are supported.
+- The `i32 %scope` argument takes values within the range \[0, 4) and selects
+ the `.scope` qualifier of the PTX instruction:
+
+| Value | `scope` |
+|-------|-------------|
+| 0 | `.cta` |
+| 1 | `.cluster` |
+| 2 | `.gpu` |
+| 3 | `.sys` |
+
+- The trailing `i32 %validate_pattern` flag selects the data-validity reporting
+ mechanism; see
+ [TMA Global-to-Shared Data-Validity Reporting](#tma-global-to-shared-data-validity-reporting).
For more information, refer [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk).
@@ -1781,7 +1836,7 @@ For more information, refer [PTX ISA](https://docs.nvidia.com/cuda/parallel-thre
##### Syntax:
```llvm
-declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cta(ptr addrspace(3) %dst, ptr addrspace(3) %mbar, ptr addrspace(1) %src, i32 %size, i64 %ch, i1 %flag_ch)
+declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cta(ptr addrspace(3) %dst, ptr addrspace(3) %mbar, ptr addrspace(1) %src, i32 %size, i32 %ibl, i32 %ibr, i64 %ch, i1 %flag_oob, i1 %flag_ch, i32 %validate_pattern)
```
##### Overview:
@@ -1790,10 +1845,75 @@ The '`@llvm.nvvm.cp.async.bulk.global.to.shared.cta`' intrinsic corresponds to
the `cp.async.bulk.shared::cta.global.*` family of PTX instructions. These
instructions initiate an asynchronous copy of bulk data from global memory to
shared::cta memory. The 32-bit operand `%size` specifies the amount of memory
-to be copied and it must be a multiple of 16. The last argument (denoted by
-`i1 %flag_ch`) is a compile-time constant. When set, it indicates a valid
-cache_hint (`i64 %ch`) and generates the `.L2::cache_hint` variant of the
-PTX instruction.
+to be copied and it must be a multiple of 16.
+
+- The trailing `%flag_oob`, `%flag_ch`, and `%validate_pattern` arguments
+ control the out-of-bounds handling, the cache-hint modifier, and the
+ data-validity reporting pattern. They must be compile-time constants.
+- The argument denoted by `i1 %flag_ch`, when set, indicates a valid cache_hint
+ (`i64 %ch`) and generates the `.L2::cache_hint` variant of the PTX
+ instruction.
+- The argument denoted by `i1 %flag_oob`, when set, generates the
+ `.ignore_oob` variant of the PTX instruction. In that case, the `i32 %ibl`
+ and `i32 %ibr` operands specify the numbers of bytes to ignore on the left
+ and right sides of the source, respectively. Both operands must be in the
+ range \[0, 15]; otherwise, the behavior is undefined. The `.ignore_oob`
+ modifier requires PTX ISA 9.2 or later. When `%flag_oob` is set,
+ `%validate_pattern` must be 0 (data-validity reporting disabled).
+- The trailing `i32 %validate_pattern` flag selects the data-validity reporting
+ mechanism; see
+ [TMA Global-to-Shared Data-Validity Reporting](#tma-global-to-shared-data-validity-reporting).
+
+For more information, refer [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk).
+
+#### '`llvm.nvvm.cp.async.bulk.global.to.shared.cta.relaxed`'
+
+##### Syntax:
+
+```llvm
+declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cta.relaxed(ptr addrspace(3) %dst, ptr addrspace(3) %mbar, ptr addrspace(1) %src, i32 %size, i32 %ibl, i32 %ibr, i64 %ch, i1 %flag_oob, i1 %flag_ch, i32 %scope, i32 %validate_pattern)
+```
+
+##### Overview:
+
+The '`@llvm.nvvm.cp.async.bulk.global.to.shared.cta.relaxed`' intrinsic
+corresponds to the
+`cp.async.bulk.relaxed.<scope>.shared::cta.global.mbarrier::complete_tx::bytes{...}.b128`
+family of PTX instructions. These instructions initiate an asynchronous copy of bulk data from
+global memory to shared::cta memory, like the weak variant above. Unlike
+the weak variant, the copy accesses memory with naturally-aligned strong
+memory operations with element-wise atomic size specified by the `.b128` type (implicit) and
+thread scope specified by `%scope`; the complete-tx operation on the mbarrier has `.release` semantics at the
+`.cluster` scope. The `.relaxed` variants require PTX ISA 9.3 or later and
+`sm_90a`, `sm_100f`, or `sm_110f` or later in the same family.
+
+- The trailing `%flag_oob`, `%flag_ch`, `%scope`, and `%validate_pattern`
+ arguments control the out-of-bounds handling, the cache-hint modifier, the
+ scope of the relaxed memory ordering semantics, and the data-validity
+ reporting pattern. They must be compile-time constants.
+- The argument denoted by `i1 %flag_ch`, when set, indicates a valid cache_hint
+ (`i64 %ch`) and generates the `.L2::cache_hint` variant of the PTX
+ instruction.
+- The argument denoted by `i1 %flag_oob`, when set, generates the
+ `.ignore_oob` variant of the PTX instruction. In that case, the `i32 %ibl`
+ and `i32 %ibr` operands specify the numbers of bytes to ignore on the left
+ and right sides of the source, respectively. Both operands must be in the
+ range \[0, 15]; otherwise, the behavior is undefined. When `%flag_oob` is
+ set, `%validate_pattern` must be 0 (data-validity reporting disabled); this
+ mutual exclusion is enforced by the verifier.
+- The `i32 %scope` argument takes values within the range \[0, 4) and selects
+ the `.scope` qualifier of the PTX instruction:
+
+| Value | `scope` |
+|-------|-------------|
+| 0 | `.cta` |
+| 1 | `.cluster` |
+| 2 | `.gpu` |
+| 3 | `.sys` |
+
+- The trailing `i32 %validate_pattern` flag selects the data-validity reporting
+ mechanism; see
+ [TMA Global-to-Shared Data-Validity Reporting](#tma-global-to-shared-data-validity-reporting).
For more information, refer [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk).
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 87ae664ce4c17..4389d8858fdb3 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -1443,6 +1443,16 @@ class NVVM_TMA_VALIDATE_PATTERN_ARGPROP<int argNo> {
];
}
+// Helper that emits the intrinsic-property set for the `scope` immediate
+// argument. The allowed values are: cta, cluster, gpu, sys. i.e. [0, 4).
+class NVVM_MEM_SCOPE_ARGPROP<int argNo> {
+ list<IntrinsicProperty> Prop = [
+ Range<ArgIndex<argNo>, 0, 4>,
+ ArgInfo<ArgIndex<argNo>, [ArgName<"scope">,
+ ImmArgPrinter<"printMemScope">]>
+ ];
+}
+
class CP_ASYNC_BULK_TENSOR_PREFETCH_OVERRIDE_INTR<int dim, string mode,
string override,
bit IsEvict = 0>
@@ -3519,31 +3529,75 @@ let IntrProperties = [NoCapture<ArgIndex<0>>, ImmArg<ArgIndex<1>>, IntrHasSideEf
}
// Intrinsics for Bulk Copy using TMA (non-tensor)
-// From Global to Shared Cluster
+// From Global to Shared Cluster (weak memory ordering).
def int_nvvm_cp_async_bulk_global_to_shared_cluster
: DefaultAttrsIntrinsicFlags<[],
[llvm_shared_cluster_ptr_ty, // dst_shared_cluster_ptr
- llvm_shared_ptr_ty, // mbarrier_ptr
- llvm_global_ptr_ty, // src_gmem_ptr
- llvm_i32_ty, // copy_size
- llvm_i16_ty, // cta_mask
- llvm_i64_ty], // cache_hint
+ llvm_shared_ptr_ty, // mbarrier_ptr
+ llvm_global_ptr_ty, // src_gmem_ptr
+ llvm_i32_ty, // copy_size
+ llvm_anyint_ty, // cta_mask (i16 or i32 overload)
+ llvm_i64_ty], // cache_hint
[llvm_i1_ty, // Flag for cta_mask
- llvm_i1_ty], // Flag for cache_hint
- [IntrConvergent, IntrArgMemOnly,
- WriteOnly<ArgIndex<0>>, ReadOnly<ArgIndex<2>>]>;
+ llvm_i1_ty, // Flag for cache_hint
+ llvm_i32_ty], // validate_pattern
+ !listconcat([IntrConvergent, IntrArgMemOnly,
+ WriteOnly<ArgIndex<0>>, ReadOnly<ArgIndex<2>>],
+ NVVM_TMA_VALIDATE_PATTERN_ARGPROP<8>.Prop)>;
+
+// From Global to Shared Cluster (relaxed memory ordering semantics).
+def int_nvvm_cp_async_bulk_global_to_shared_cluster_relaxed
+ : DefaultAttrsIntrinsicFlags<[],
+ [llvm_shared_cluster_ptr_ty, // dst_shared_cluster_ptr
+ llvm_shared_ptr_ty, // mbarrier_ptr
+ llvm_global_ptr_ty, // src_gmem_ptr
+ llvm_i32_ty, // copy_size
+ llvm_anyint_ty, // cta_mask (i16 or i32 overload)
+ llvm_i64_ty], // cache_hint
+ [llvm_i1_ty, // Flag for cta_mask
+ llvm_i1_ty, // Flag for cache_hint
+ llvm_i32_ty, // scope (cta/cluster/gpu/sys)
+ llvm_i32_ty], // validate_pattern
+ !listconcat([IntrConvergent, IntrArgMemOnly,
+ WriteOnly<ArgIndex<0>>, ReadOnly<ArgIndex<2>>],
+ NVVM_MEM_SCOPE_ARGPROP<8>.Prop,
+ NVVM_TMA_VALIDATE_PATTERN_ARGPROP<9>.Prop)>;
-// From Global to Shared CTA
+// From Global to Shared CTA (weak memory ordering).
def int_nvvm_cp_async_bulk_global_to_shared_cta
: DefaultAttrsIntrinsicFlags<[],
[llvm_shared_ptr_ty, // dst_shared_cta_ptr
llvm_shared_ptr_ty, // mbarrier_ptr
llvm_global_ptr_ty, // src_gmem_ptr
llvm_i32_ty, // copy_size
+ llvm_i32_ty, // ignore_bytes_left (controlled by flag_oob)
+ llvm_i32_ty, // ignore_bytes_right (controlled by flag_oob)
llvm_i64_ty], // cache_hint
- [llvm_i1_ty], // Flag for cache_hint
- [IntrConvergent, IntrArgMemOnly,
- WriteOnly<ArgIndex<0>>, ReadOnly<ArgIndex<2>>]>;
+ [llvm_i1_ty, // Flag for ignore_oob
+ llvm_i1_ty, // Flag for cache_hint
+ llvm_i32_ty], // validate_pattern
+ !listconcat([IntrConvergent, IntrArgMemOnly,
+ WriteOnly<ArgIndex<0>>, ReadOnly<ArgIndex<2>>],
+ NVVM_TMA_VALIDATE_PATTERN_ARGPROP<9>.Prop)>;
+
+// From Global to Shared CTA (relaxed memory ordering semantics).
+def int_nvvm_cp_async_bulk_global_to_shared_cta_relaxed
+ : DefaultAttrsIntrinsicFlags<[],
+ [llvm_shared_ptr_ty, // dst_shared_cta_ptr
+ llvm_shared_ptr_ty, // mbarrier_ptr
+ llvm_global_ptr_ty, // src_gmem_ptr
+ llvm_i32_ty, // copy_size
+ llvm_i32_ty, // ignore_bytes_left (controlled by flag_oob)
+ llvm_i32_ty, // ignore_bytes_right (controlled by flag_oob)
+ llvm_i64_ty], // cache_hint
+ [llvm_i1_ty, // Flag for ignore_oob
+ llvm_i1_ty, // Flag for cache_hint
+ llvm_i32_ty, // scope (cta/cluster/gpu/sys)
+ llvm_i32_ty], // validate_pattern
+ !listconcat([IntrConvergent, IntrArgMemOnly,
+ WriteOnly<ArgIndex<0>>, ReadOnly<ArgIndex<2>>],
+ NVVM_MEM_SCOPE_ARGPROP<9>.Prop,
+ NVVM_TMA_VALIDATE_PATTERN_ARGPROP<10>.Prop)>;
// From Shared CTA to Shared Cluster
def int_nvvm_cp_async_bulk_shared_cta_to_cluster
diff --git a/llvm/include/llvm/IR/NVVMIntrinsicUtils.h b/llvm/include/llvm/IR/NVVMIntrinsicUtils.h
index 6fdaf631a12f9..abd6725d056fc 100644
--- a/llvm/include/llvm/IR/NVVMIntrinsicUtils.h
+++ b/llvm/include/llvm/IR/NVVMIntrinsicUtils.h
@@ -103,6 +103,28 @@ inline StringRef getTMAValidateDataPatternName(TMAValidateDataPattern Pattern) {
llvm_unreachable("invalid TMA validate data pattern");
}
+// Scope of the memory ordering semantics.
+enum class MemScope : uint8_t {
+ CTA = 0,
+ CLUSTER = 1,
+ GPU = 2,
+ SYS = 3,
+};
+
+inline StringRef getMemScopeName(MemScope Scope) {
+ switch (Scope) {
+ case MemScope::CTA:
+ return "cta";
+ case MemScope::CLUSTER:
+ return "cluster";
+ case MemScope::GPU:
+ return "gpu";
+ case MemScope::SYS:
+ return "sys";
+ }
+ llvm_unreachable("invalid memory scope");
+}
+
// Eviction priorities applicable for prefetch and applypriority intrinsics.
enum class EvictPolicyType : uint8_t {
EVICT_NORMAL = 0, // default
@@ -195,6 +217,8 @@ LLVM_ABI void printTMAReductionOp(raw_ostream &OS, const Constant *ImmArgVal);
LLVM_ABI void printTMAValidateDataPattern(raw_ostream &OS,
const Constant *ImmArgVal);
+LLVM_ABI void printMemScope(raw_ostream &OS, const Constant *ImmArgVal);
+
LLVM_ABI void printTcgen05CollectorUsageOp(raw_ostream &OS,
const Constant *ImmArgVal);
diff --git a/llvm/lib/IR/AutoUpgrade.cpp b/llvm/lib/IR/AutoUpgrade.cpp
index 530ee5b7b6047..4d756aa2c25c5 100644
--- a/llvm/lib/IR/AutoUpgrade.cpp
+++ b/llvm/lib/IR/AutoUpgrade.cpp
@@ -1264,6 +1264,55 @@ static Intrinsic::ID shouldUpgradeNVPTXTMAG2SCTAIntrinsics(Function *F,
return ID;
}
+// The legacy tail of llvm.nvvm.cp.async.bulk.global.to.shared.cluster is:
+//
+// ..., i16 %mc, i64 %ch, i1 %flag_mc, i1 %flag_ch
+//
+// The current intrinsic is overloaded on the multicast-mask type and takes a
+// trailing i32 %validate_pattern; the legacy tail is recognized by an i1 at
+// parameter N-1.
+static Intrinsic::ID
+shouldUpgradeNVPTXBulkG2SClusterIntrinsic(Function *F, StringRef Name,
+ SmallVectorImpl<Type *> &OvlTys) {
+ if (!Name.consume_front("cp.async.bulk.global.to.shared.cluster"))
+ return Intrinsic::not_intrinsic;
+
+ // Parameter N-1 is i1 for the legacy tail; the current tail ends with
+ // i32 %validate_pattern, for which N-1 is i32.
+ size_t NumParams = F->getFunctionType()->getNumParams();
+ if (!F->getFunctionType()->getParamType(NumParams - 1)->isIntegerTy(1))
+ return Intrinsic::not_intrinsic;
+
+ // The multicast mask is parameter 4; legacy IR only uses i16.
+ Type *MaskTy = F->getFunctionType()->getParamType(NumParams - 4);
+ if (!MaskTy->isIntegerTy(16))
+ return Intrinsic::not_intrinsic;
+ OvlTys.push_back(MaskTy);
+
+ return Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster;
+}
+
+// The legacy tail of llvm.nvvm.cp.async.bulk.global.to.shared.cta is:
+//
+// ..., i64 %ch, i1 %flag_ch
+//
+// The current...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/222323
More information about the llvm-commits
mailing list