[llvm] [mlir] [LLVM][NVPTX] Add async bulk copy global to shared extensions (PR #222323)
Rajat Bajpai via llvm-commits
llvm-commits at lists.llvm.org
Wed Sep 9 05:56:46 PDT 2026
https://github.com/rajatbajpai created https://github.com/llvm/llvm-project/pull/222323
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.
>From b6cf19b30fbd02c7d705a6effadefddd49318f9f Mon Sep 17 00:00:00 2001
From: Rajat Bajpai <rbajpai at nvidia.com>
Date: Tue, 8 Sep 2026 10:17:24 +0000
Subject: [PATCH] [LLVM][NVPTX] Add async bulk copy global to shared extensions
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.
---
llvm/docs/NVPTXUsage.md | 156 +++++++++++++--
llvm/include/llvm/IR/IntrinsicsNVVM.td | 80 ++++++--
llvm/include/llvm/IR/NVVMIntrinsicUtils.h | 24 +++
llvm/lib/IR/AutoUpgrade.cpp | 108 +++++++++-
llvm/lib/IR/NVVMIntrinsicUtils.cpp | 10 +
llvm/lib/IR/Verifier.cpp | 14 ++
.../NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp | 6 +
.../NVPTX/MCTargetDesc/NVPTXInstPrinter.h | 2 +
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td | 10 +
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 189 +++++++++++++-----
.../Assembler/auto_upgrade_nvvm_intrinsics.ll | 9 +-
.../NVPTX/cp-async-bulk-g2s-cta-invalid.ll | 30 +++
.../NVPTX/cp-async-bulk-g2s-cta-oob.ll | 27 +++
.../NVPTX/cp-async-bulk-g2s-cta-relaxed.ll | 75 +++++++
.../cp-async-bulk-g2s-cta-validate-data.ll | 56 ++++++
.../NVPTX/cp-async-bulk-g2s-invalid.ll | 31 +++
.../CodeGen/NVPTX/cp-async-bulk-g2s-mc32.ll | 52 +++++
.../NVPTX/cp-async-bulk-g2s-relaxed.ll | 59 ++++++
.../NVPTX/cp-async-bulk-g2s-validate-data.ll | 58 ++++++
.../Verifier/NVPTX/cp-async-bulk-g2s-cta.ll | 16 ++
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 2 +-
mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp | 14 +-
.../Target/LLVMIR/nvvm/tma_bulk_copy.mlir | 12 +-
23 files changed, 943 insertions(+), 97 deletions(-)
create mode 100644 llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-cta-invalid.ll
create mode 100644 llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-cta-oob.ll
create mode 100644 llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-cta-relaxed.ll
create mode 100644 llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-cta-validate-data.ll
create mode 100644 llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-invalid.ll
create mode 100644 llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-mc32.ll
create mode 100644 llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-relaxed.ll
create mode 100644 llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-validate-data.ll
create mode 100644 llvm/test/Verifier/NVPTX/cp-async-bulk-g2s-cta.ll
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 intrinsic adds %ignore_bytes_left/%ignore_bytes_right before
+// %ch and trailing %flag_oob/%validate_pattern; the legacy tail is
+// recognized by an i1 at parameter N-1, whereas the current tail ends
+// with an i32.
+static Intrinsic::ID shouldUpgradeNVPTXBulkG2SCTAIntrinsic(Function *F,
+ StringRef Name) {
+ if (!Name.consume_front("cp.async.bulk.global.to.shared.cta"))
+ 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.
+ if (!F->getFunctionType()->getParamType(5)->isIntegerTy(1))
+ return Intrinsic::not_intrinsic;
+
+ return Intrinsic::nvvm_cp_async_bulk_global_to_shared_cta;
+}
+
// The legacy TMA reduction intrinsics encode the reduction operator in their
// name, while the current ones take it as an immediate argument. Map the
// operator part of a legacy name to the corresponding immediate value.
@@ -1310,8 +1359,6 @@ static Intrinsic::ID shouldUpgradeNVPTXSharedClusterIntrinsic(Function *F,
if (Name.consume_front("cp.async.bulk.")) {
Intrinsic::ID ID =
StringSwitch<Intrinsic::ID>(Name)
- .Case("global.to.shared.cluster",
- Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster)
.Case("shared.cta.to.cluster",
Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster)
.Default(Intrinsic::not_intrinsic);
@@ -2081,6 +2128,26 @@ static bool upgradeIntrinsicFunction1(Function *F, Function *&NewFn,
return true;
}
+ // Upgrade the legacy cp.async.bulk.global.to.shared.cluster signature
+ // (multicast-mask overloading + trailing validate_pattern).
+ SmallVector<Type *, 1> BulkG2SOvlTys;
+ IID = shouldUpgradeNVPTXBulkG2SClusterIntrinsic(F, Name, BulkG2SOvlTys);
+ if (IID != Intrinsic::not_intrinsic) {
+ rename(F);
+ NewFn = Intrinsic::getOrInsertDeclaration(F->getParent(), IID,
+ BulkG2SOvlTys);
+ return true;
+ }
+
+ // Upgrade the legacy cp.async.bulk.global.to.shared.cta signature
+ // (no ignore_bytes_left/right + trailing validate_pattern).
+ IID = shouldUpgradeNVPTXBulkG2SCTAIntrinsic(F, Name);
+ if (IID != Intrinsic::not_intrinsic) {
+ rename(F);
+ NewFn = Intrinsic::getOrInsertDeclaration(F->getParent(), IID);
+ return true;
+ }
+
// Upgrade tcgen05.mma intrinsics missing collector_usage_b.
IID = shouldUpgradeNVPTXTcgen05MMAIntrinsic(F, Name);
if (IID != Intrinsic::not_intrinsic) {
@@ -6049,7 +6116,42 @@ void llvm::UpgradeIntrinsicCall(CallBase *CI, Function *NewFn) {
CI->eraseFromParent();
return;
}
- case Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster:
+ case Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster: {
+ SmallVector<Value *, 4> Args(CI->args());
+ unsigned AS = Args[0]->getType()->getPointerAddressSpace();
+ if (AS == NVPTXAS::ADDRESS_SPACE_SHARED)
+ Args[0] = Builder.CreateAddrSpaceCast(
+ Args[0], Builder.getPtrTy(NVPTXAS::ADDRESS_SPACE_SHARED_CLUSTER));
+
+ // Append the missing trailing validate_pattern (0 = disabled).
+ Args.push_back(Builder.getInt32(0));
+
+ NewCall = Builder.CreateCall(NewFn, Args);
+ NewCall->takeName(CI);
+ CI->replaceAllUsesWith(NewCall);
+ CI->eraseFromParent();
+ return;
+ }
+ case Intrinsic::nvvm_cp_async_bulk_global_to_shared_cta: {
+ // (dst, mbar, src, size, ch, flag_ch)
+ // -> (dst, mbar, src, size, i32 0, i32 0, ch, i1 false, flag_ch,
+ // i32 0 /* validate_pattern=disabled */)
+ SmallVector<Value *, 10> Args;
+ for (unsigned I = 0; I < 4; ++I)
+ Args.push_back(CI->getArgOperand(I));
+ Args.push_back(Builder.getInt32(0)); // ignore_bytes_left
+ Args.push_back(Builder.getInt32(0)); // ignore_bytes_right
+ Args.push_back(CI->getArgOperand(4)); // cache_hint
+ Args.push_back(Builder.getInt1(false)); // flag_oob
+ Args.push_back(CI->getArgOperand(5)); // flag_ch
+ Args.push_back(Builder.getInt32(0)); // validate_pattern
+
+ NewCall = Builder.CreateCall(NewFn, Args);
+ NewCall->takeName(CI);
+ CI->replaceAllUsesWith(NewCall);
+ CI->eraseFromParent();
+ return;
+ }
case Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster: {
// Create a new call with the correct address space.
SmallVector<Value *, 4> Args(CI->args());
diff --git a/llvm/lib/IR/NVVMIntrinsicUtils.cpp b/llvm/lib/IR/NVVMIntrinsicUtils.cpp
index f6816fa0e706b..387e555f8899e 100644
--- a/llvm/lib/IR/NVVMIntrinsicUtils.cpp
+++ b/llvm/lib/IR/NVVMIntrinsicUtils.cpp
@@ -49,6 +49,16 @@ void nvvm::printTMAValidateDataPattern(raw_ostream &OS,
static_cast<TMAValidateDataPattern>(CI->getZExtValue()));
}
+void nvvm::printMemScope(raw_ostream &OS, const Constant *ImmArgVal) {
+ const auto *CI = dyn_cast<ConstantInt>(ImmArgVal);
+ if (!CI || CI->getZExtValue() > static_cast<uint64_t>(MemScope::SYS)) {
+ OS << "Unknown memory scope";
+ return;
+ }
+
+ OS << getMemScopeName(static_cast<MemScope>(CI->getZExtValue()));
+}
+
void nvvm::printTcgen05MMAKind(raw_ostream &OS, const Constant *ImmArgVal) {
if (const auto *CI = dyn_cast<ConstantInt>(ImmArgVal)) {
uint64_t Val = CI->getZExtValue();
diff --git a/llvm/lib/IR/Verifier.cpp b/llvm/lib/IR/Verifier.cpp
index 26c34b4221fde..5f3316e2fef8b 100644
--- a/llvm/lib/IR/Verifier.cpp
+++ b/llvm/lib/IR/Verifier.cpp
@@ -7165,6 +7165,20 @@ void Verifier::visitIntrinsicCall(Intrinsic::ID ID, CallBase &Call) {
"reg_count argument to nvvm.setmaxnreg must be in multiples of 8");
break;
}
+ case Intrinsic::nvvm_cp_async_bulk_global_to_shared_cta:
+ case Intrinsic::nvvm_cp_async_bulk_global_to_shared_cta_relaxed: {
+ const unsigned ArgSize = Call.arg_size();
+ const unsigned ValidatePatternIndex = ArgSize - 1;
+ const unsigned IgnoreOOBFlagIndex = 7;
+ bool IgnoreOOB =
+ cast<ConstantInt>(Call.getArgOperand(IgnoreOOBFlagIndex))->isOne();
+ const auto *ValidatePattern =
+ cast<ConstantInt>(Call.getArgOperand(ValidatePatternIndex));
+ Check(!IgnoreOOB || ValidatePattern->isZero(),
+ "validate_pattern must be 0 (disabled) when ignore_oob is enabled",
+ &Call);
+ break;
+ }
case Intrinsic::experimental_convergence_entry:
case Intrinsic::experimental_convergence_anchor:
break;
diff --git a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp
index ce6ab63eebb64..dc250c67c3bc9 100644
--- a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp
+++ b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp
@@ -590,6 +590,12 @@ void NVPTXInstPrinter::printTMAValidateDataFlags(const MCInst *MI, int OpNum,
<< nvvm::getTMAValidateDataPatternName(Pattern);
}
+void NVPTXInstPrinter::printMemScope(const MCInst *MI, int OpNum,
+ const MCSubtargetInfo &, raw_ostream &O) {
+ const MCOperand &MO = MI->getOperand(OpNum);
+ O << "." << nvvm::getMemScopeName(static_cast<nvvm::MemScope>(MO.getImm()));
+}
+
void NVPTXInstPrinter::printEvictPolicy(const MCInst *MI, int OpNum,
const MCSubtargetInfo &, raw_ostream &O,
StringRef Modifier) {
diff --git a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.h b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.h
index 17438cb874744..dc3e229364d0e 100644
--- a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.h
+++ b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.h
@@ -69,6 +69,8 @@ class NVPTXInstPrinter : public MCInstPrinter {
raw_ostream &O);
void printTMAValidateDataFlags(const MCInst *MI, int OpNum,
const MCSubtargetInfo &STI, raw_ostream &O);
+ void printMemScope(const MCInst *MI, int OpNum, const MCSubtargetInfo &STI,
+ raw_ostream &O);
void printEvictPolicy(const MCInst *MI, int OpNum, const MCSubtargetInfo &STI,
raw_ostream &O, StringRef Modifier = {});
void printCallOperand(const MCInst *MI, int OpNum, const MCSubtargetInfo &STI,
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index 846bcc0557464..d1ebb56943bf3 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -208,6 +208,16 @@ def hasMMASparseWithMXF4NVF4Scale4xE8M0
// - tile_scatter4 mode support
def hasTMABlackwellSupport : PredOr<[SM100f, SM110f]>;
+
+def hasCpAsyncBulkClusterSupport : PredAnd<[PTX80, SM90]>;
+def hasCpAsyncBulkCTASupport : PredAnd<[PTX86, SM90]>;
+
+// Checks support for relaxed memory ordering semantics support in cp.async.bulk family.
+def hasCpAsyncBulkRelaxedSupport : PredAnd<[PTX93, PredOr<[SM90a, SM100f, SM110f]>]>;
+
+// Checks for out of bounds (OOB) support in cp.async.bulk family.
+def hasCpAsyncBulkOOBSupport : PredAnd<[PTX92, SM90]>;
+
// Checks support for conversions involving e4m3x2 and e5m2x2. These arrived
// with sm_90 in PTX 7.8, which SM90 already implies, and were extended to
// sm_89 in PTX 8.1.
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 220ef64732830..8c0151e81545b 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -524,20 +524,35 @@ let Predicates = [PTX80, SM90] in {
// TMA Async Bulk Copy Functions
//------------------------------
-class CpAsyncBulkStr<bit mc, bit ch, bit mask = 0> {
+class CpAsyncBulkStr<bits<2> mc, bit ch, bit mask = 0, bit is_relaxed = 0,
+ bit oob = 0> {
+ // Multicast kind: 0 = none, 1 = 16-bit mask, 2 = 32-bit mask.
+ string mc_str = !cond(!eq(mc, 1): ".multicast::cluster",
+ !eq(mc, 2): ".multicast::cluster::32b",
+ true: "");
+
// Shared to Global memory
string S2G = "cp.async.bulk.global.shared::cta.bulk_group"
# !if(ch, ".L2::cache_hint", "")
# !if(mask, ".cp_mask", "");
- // Global to Shared cluster memory
- string G2S = "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes"
- # !if(mc, ".multicast::cluster", "")
- # !if(ch, ".L2::cache_hint", "");
+ // Global to Shared cluster memory. The relaxed form carries the "${scope}"
+ // qualifier and an explicit ".b128" suffix; the weak form has neither.
+ defvar prefix = !if(is_relaxed,
+ "cp.async.bulk.relaxed${scope}.shared::cluster.global.mbarrier::complete_tx::bytes$validate",
+ "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes$validate");
+ string G2S = prefix # mc_str
+ # !if(ch, ".L2::cache_hint", "")
+ # !if(is_relaxed, ".b128", "");
- // Global to Shared CTA memory
- string G2S_CTA = "cp.async.bulk.shared::cta.global.mbarrier::complete_tx::bytes"
- # !if(ch, ".L2::cache_hint", "");
+ // Global to Shared CTA memory.
+ defvar prefix_cta = !if(is_relaxed,
+ "cp.async.bulk.relaxed${scope}.shared::cta.global.mbarrier::complete_tx::bytes$validate",
+ "cp.async.bulk.shared::cta.global.mbarrier::complete_tx::bytes$validate");
+ string G2S_CTA = prefix_cta
+ # !if(ch, ".L2::cache_hint", "")
+ # !if(oob, ".ignore_oob", "")
+ # !if(is_relaxed, ".b128", "");
// Shared CTA to Cluster memory
string C2C = "cp.async.bulk.shared::cluster.shared::cta.mbarrier::complete_tx::bytes";
@@ -561,44 +576,127 @@ multiclass CP_ASYNC_BULK_S2G_INTR<bit has_ch> {
defm CP_ASYNC_BULK_S2G : CP_ASYNC_BULK_S2G_INTR<has_ch = 0>;
defm CP_ASYNC_BULK_S2G_CH : CP_ASYNC_BULK_S2G_INTR<has_ch = 1>;
-multiclass CP_ASYNC_BULK_G2S_INTR<bit has_ch> {
- defvar Intr = int_nvvm_cp_async_bulk_global_to_shared_cluster;
+def TMAValidateDataFlags : Operand<i32> {
+ let PrintMethod = "printTMAValidateDataFlags";
+}
- def "" : NVPTXInst<(outs),
- (ins ADDR:$dst, ADDR:$mbar, ADDR:$src,
- B32:$size, B16:$mask, B64:$ch),
- !if(has_ch,
- CpAsyncBulkStr<0, 1>.G2S # " [$dst], [$src], $size, [$mbar], $ch;",
- CpAsyncBulkStr<0, 0>.G2S # " [$dst], [$src], $size, [$mbar];"),
- [(Intr addr:$dst, addr:$mbar, addr:$src, i32:$size, i16:$mask, i64:$ch, 0, !if(has_ch, -1, 0))]>,
- Requires<[PTX80, SM90]>;
+// Matches validate pattern values 0-5. The 1-5 values are Rubin-only
+// validate patterns. 0 denotes disabled validate pattern (default).
+def tma_validate_data_imm : TImmLeaf<i32, [{
+ return Imm == 0 || (Imm >= 1 && Imm <= 5 &&
+ Subtarget->hasFeature(NVPTX::SM107f));
+}]>;
- def _MC : NVPTXInst<(outs),
- (ins ADDR:$dst, ADDR:$mbar, ADDR:$src,
- B32:$size, B16:$mask, B64:$ch),
- !if(has_ch,
- CpAsyncBulkStr<1, 1>.G2S # " [$dst], [$src], $size, [$mbar], $mask, $ch;",
- CpAsyncBulkStr<1, 0>.G2S # " [$dst], [$src], $size, [$mbar], $mask;"),
- [(Intr addr:$dst, addr:$mbar, addr:$src, i32:$size, i16:$mask, i64:$ch, -1, !if(has_ch, -1, 0))]>,
- Requires<[PTX80, SM90]>;
+// Prints the scope for relaxed memory ordering semantics.
+def MemScopeFlags : Operand<i32> {
+ let PrintMethod = "printMemScope";
+}
+
+// Matches scope values 0-3 (cta/cluster/gpu/sys) for relaxed memory ordering semantics.
+def async_bulk_scope_imm : TImmLeaf<i32, [{ return Imm >= 0 && Imm <= 3; }]>;
+
+multiclass CP_ASYNC_BULK_G2S_INTR<bit has_ch, bit is_relaxed,
+ list<Predicate> preds = [hasCpAsyncBulkClusterSupport]> {
+ defvar Intr =
+ !cast<Intrinsic>("int_nvvm_cp_async_bulk_global_to_shared_cluster" #
+ !if(is_relaxed, "_relaxed", ""));
+
+ defvar asm_nomc = CpAsyncBulkStr<0, has_ch, 0, is_relaxed>.G2S;
+ defvar asm_mc16 = CpAsyncBulkStr<1, has_ch, 0, is_relaxed>.G2S;
+ defvar asm_mc32 = CpAsyncBulkStr<2, has_ch, 0, is_relaxed>.G2S;
+
+ // Relaxed form carries the extra scope operand.
+ defvar ins_tail = !if(is_relaxed,
+ (ins MemScopeFlags:$scope, TMAValidateDataFlags:$validate),
+ (ins TMAValidateDataFlags:$validate));
+ defvar ins_i16 = !con((ins ADDR:$dst, ADDR:$mbar, ADDR:$src, B32:$size,
+ B16:$mask, B64:$ch), ins_tail);
+ defvar ins_i32 = !con((ins ADDR:$dst, ADDR:$mbar, ADDR:$src, B32:$size,
+ B32:$mask, B64:$ch), ins_tail);
+
+ // Relaxed form carries the extra scope operand.
+ defvar p_tail = !if(is_relaxed,
+ (Intr async_bulk_scope_imm:$scope, tma_validate_data_imm:$validate),
+ (Intr tma_validate_data_imm:$validate));
+ defvar p_no_mc16 = !con((Intr addr:$dst, addr:$mbar, addr:$src, i32:$size,
+ i16:$mask, i64:$ch, 0, !if(has_ch, -1, 0)), p_tail);
+ defvar p_mc16 = !con((Intr addr:$dst, addr:$mbar, addr:$src, i32:$size,
+ i16:$mask, i64:$ch, -1, !if(has_ch, -1, 0)), p_tail);
+ defvar p_no_mc32 = !con((Intr addr:$dst, addr:$mbar, addr:$src, i32:$size,
+ i32:$mask, i64:$ch, 0, !if(has_ch, -1, 0)), p_tail);
+ defvar p_mc32 = !con((Intr addr:$dst, addr:$mbar, addr:$src, i32:$size,
+ i32:$mask, i64:$ch, -1, !if(has_ch, -1, 0)), p_tail);
+
+ defvar args_ch = !if(has_ch, ", $ch", "");
+
+ // 16-bit multicast mask variants
+ def "" : NVPTXInst<(outs), ins_i16,
+ asm_nomc # " [$dst], [$src], $size, [$mbar]" # args_ch # ";",
+ [p_no_mc16]>, Requires<preds>;
+ def _MC : NVPTXInst<(outs), ins_i16,
+ asm_mc16 # " [$dst], [$src], $size, [$mbar], $mask" # args_ch # ";",
+ [p_mc16]>, Requires<preds>;
+
+ // 32-bit multicast mask variants
+ let Predicates = [hasRubinFamilySupport] in {
+ def _NO_MC32 : NVPTXInst<(outs), ins_i32,
+ asm_nomc # " [$dst], [$src], $size, [$mbar]" # args_ch # ";",
+ [p_no_mc32]>;
+ def _MC32 : NVPTXInst<(outs), ins_i32,
+ asm_mc32 # " [$dst], [$src], $size, [$mbar], $mask" # args_ch # ";",
+ [p_mc32]>;
+ }
}
-defm CP_ASYNC_BULK_G2S : CP_ASYNC_BULK_G2S_INTR<has_ch = 0>;
-defm CP_ASYNC_BULK_G2S_CH : CP_ASYNC_BULK_G2S_INTR<has_ch = 1>;
-multiclass CP_ASYNC_BULK_G2S_CTA_INTR<bit has_ch> {
- defvar Intr = int_nvvm_cp_async_bulk_global_to_shared_cta;
+defm CP_ASYNC_BULK_G2S : CP_ASYNC_BULK_G2S_INTR</*has_ch=*/0, /*is_relaxed=*/0>;
+defm CP_ASYNC_BULK_G2S_CH : CP_ASYNC_BULK_G2S_INTR</*has_ch=*/1, /*is_relaxed=*/0>;
+defm CP_ASYNC_BULK_G2S_RELAXED : CP_ASYNC_BULK_G2S_INTR</*has_ch=*/0, /*is_relaxed=*/1, [hasCpAsyncBulkRelaxedSupport]>;
+defm CP_ASYNC_BULK_G2S_RELAXED_CH : CP_ASYNC_BULK_G2S_INTR</*has_ch=*/1, /*is_relaxed=*/1, [hasCpAsyncBulkRelaxedSupport]>;
- def "" : NVPTXInst<(outs),
- (ins ADDR:$dst, ADDR:$mbar, ADDR:$src,
- B32:$size, B64:$ch),
- !if(has_ch,
- CpAsyncBulkStr<0, 1>.G2S_CTA # " [$dst], [$src], $size, [$mbar], $ch;",
- CpAsyncBulkStr<0, 0>.G2S_CTA # " [$dst], [$src], $size, [$mbar];"),
- [(Intr addr:$dst, addr:$mbar, addr:$src, i32:$size, i64:$ch, !if(has_ch, -1, 0))]>,
- Requires<[PTX86, SM90]>;
+multiclass CP_ASYNC_BULK_G2S_CTA_INTR<bit has_ch, bit is_relaxed,
+ list<Predicate> preds = [hasCpAsyncBulkCTASupport]> {
+ defvar Intr =
+ !cast<Intrinsic>("int_nvvm_cp_async_bulk_global_to_shared_cta" #
+ !if(is_relaxed, "_relaxed", ""));
+
+ defvar asm = CpAsyncBulkStr<0, has_ch, 0, is_relaxed, 0>.G2S_CTA;
+ defvar asm_oob = CpAsyncBulkStr<0, has_ch, 0, is_relaxed, 1>.G2S_CTA;
+
+ // Relaxed form carries the extra scope operand.
+ defvar ins_tail = !if(is_relaxed,
+ (ins MemScopeFlags:$scope, TMAValidateDataFlags:$validate),
+ (ins TMAValidateDataFlags:$validate));
+
+ defvar ins_dag = !con((ins ADDR:$dst, ADDR:$mbar, ADDR:$src, B32:$size,
+ B64:$ch), ins_tail);
+ defvar ins_oob_dag = !con((ins ADDR:$dst, ADDR:$mbar, ADDR:$src, B32:$size,
+ B32:$ibl, B32:$ibr, B64:$ch), ins_tail);
+
+ // flag_oob/flag_ch are matched with bare constants (0 = false, -1 = true);
+ // validate_pattern must be forwarded, hence the bindable TImmLeaf.
+ defvar p_tail = !if(is_relaxed,
+ (Intr async_bulk_scope_imm:$scope, tma_validate_data_imm:$validate),
+ (Intr tma_validate_data_imm:$validate));
+ defvar p = !con((Intr addr:$dst, addr:$mbar, addr:$src, i32:$size,
+ (i32 srcvalue), (i32 srcvalue), i64:$ch, 0, !if(has_ch, -1, 0)), p_tail);
+ defvar p_oob = !con((Intr addr:$dst, addr:$mbar, addr:$src, i32:$size,
+ i32:$ibl, i32:$ibr, i64:$ch, -1, !if(has_ch, -1, 0)), p_tail);
+
+ defvar args_ch = !if(has_ch, ", $ch", "");
+
+ def "" : NVPTXInst<(outs), ins_dag,
+ asm # " [$dst], [$src], $size, [$mbar]" # args_ch # ";",
+ [p]>, Requires<preds>;
+
+ def _OOB : NVPTXInst<(outs), ins_oob_dag,
+ asm_oob # " [$dst], [$src], $size, $ibl, $ibr, [$mbar]" # args_ch # ";",
+ [p_oob]>, Requires<!listconcat([hasCpAsyncBulkOOBSupport], preds)>;
}
-defm CP_ASYNC_BULK_G2S_CTA : CP_ASYNC_BULK_G2S_CTA_INTR<has_ch = 0>;
-defm CP_ASYNC_BULK_G2S_CTA_CH : CP_ASYNC_BULK_G2S_CTA_INTR<has_ch = 1>;
+
+defm CP_ASYNC_BULK_G2S_CTA : CP_ASYNC_BULK_G2S_CTA_INTR</*has_ch=*/0, /*is_relaxed=*/0>;
+defm CP_ASYNC_BULK_G2S_CTA_CH : CP_ASYNC_BULK_G2S_CTA_INTR</*has_ch=*/1, /*is_relaxed=*/0>;
+defm CP_ASYNC_BULK_G2S_CTA_RELAXED : CP_ASYNC_BULK_G2S_CTA_INTR</*has_ch=*/0, /*is_relaxed=*/1, [hasCpAsyncBulkRelaxedSupport]>;
+defm CP_ASYNC_BULK_G2S_CTA_RELAXED_CH : CP_ASYNC_BULK_G2S_CTA_INTR</*has_ch=*/1, /*is_relaxed=*/1, [hasCpAsyncBulkRelaxedSupport]>;
def CP_ASYNC_BULK_CTA_TO_CLUSTER : NVPTXInst<(outs),
(ins ADDR:$dst, ADDR:$mbar, ADDR:$src, B32:$size),
@@ -677,17 +775,6 @@ def CTAGroupFlags : Operand<i32> {
def tma_cta_group_imm0 : TImmLeaf<i32, [{return Imm == 0;}]>;
def tma_cta_group_imm_any : TImmLeaf<i32, [{return Imm >= 0;}]>;
-def TMAValidateDataFlags : Operand<i32> {
- let PrintMethod = "printTMAValidateDataFlags";
-}
-
-// Matches validate pattern values 0-5. The 1-5 values are Rubin-only
-// validate patterns. 0 denotes disabled validate pattern (default).
-def tma_validate_data_imm : TImmLeaf<i32, [{
- return Imm == 0 || (Imm >= 1 && Imm <= 5 &&
- Subtarget->hasFeature(NVPTX::SM107f));
-}]>;
-
multiclass TMA_TENSOR_G2S_INTR<int dim, string mode, list<Predicate> pred,
TImmLeaf cta_group_type = tma_cta_group_imm_any> {
defvar dims_dag = TMA_DIMS_DAG_UTIL<dim>.ins_dag;
diff --git a/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll b/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
index f95215834320e..af312836da06f 100644
--- a/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
+++ b/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
@@ -113,6 +113,7 @@ declare double @llvm.nvvm.atomic.add.gen.f.sys.f64.p0(ptr, double)
declare ptr addrspace(3) @llvm.nvvm.mapa.shared.cluster(ptr addrspace(3), i32)
declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster(ptr addrspace(3), ptr addrspace(3), ptr addrspace(1), i32, i16, i64, i1, i1)
+declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cta(ptr addrspace(3), ptr addrspace(3), ptr addrspace(1), i32, i64, i1)
declare void @llvm.nvvm.cp.async.bulk.shared.cta.to.cluster(ptr addrspace(3), ptr addrspace(3), ptr addrspace(3), i32)
declare void @llvm.nvvm.tcgen05.commit.cg1(ptr)
@@ -489,10 +490,14 @@ define void @nvvm_shared_cluster_intrinsics(ptr addrspace(3) %p0, i32 %offset) {
}
; CHECK-LABEL: @nvvm_cp_async_bulk_intrinsics
-define void @nvvm_cp_async_bulk_intrinsics(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, ptr addrspace(3) %src_shared, i32 %size) {
-; CHECK: call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster(ptr addrspace(7) %1, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 0, i64 0, i1 false, i1 false)
+define void @nvvm_cp_async_bulk_intrinsics(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, ptr addrspace(3) %src_shared, ptr addrspace(7) %dst_cluster, i32 %size, i16 %mc, i64 %ch) {
+; CHECK: call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.i16(ptr addrspace(7) %1, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 0, i64 0, i1 false, i1 false, /* validate_pattern=disabled */ i32 0)
+; CHECK: call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.i16(ptr addrspace(7) %dst_cluster, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 true, i1 true, /* validate_pattern=disabled */ i32 0)
+; CHECK: call void @llvm.nvvm.cp.async.bulk.global.to.shared.cta(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 0, i32 0, i64 %ch, i1 false, i1 true, /* validate_pattern=disabled */ i32 0)
; CHECK: call void @llvm.nvvm.cp.async.bulk.shared.cta.to.cluster(ptr addrspace(7) %2, ptr addrspace(3) %bar, ptr addrspace(3) %src_shared, i32 %size)
call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 0, i64 0, i1 false, i1 false)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster(ptr addrspace(7) %dst_cluster, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 true, i1 true)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cta(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i64 %ch, i1 true)
call void @llvm.nvvm.cp.async.bulk.shared.cta.to.cluster(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr addrspace(3) %src_shared, i32 %size)
ret void
}
diff --git a/llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-cta-invalid.ll b/llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-cta-invalid.ll
new file mode 100644
index 0000000000000..c17ed65fdb908
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-cta-invalid.ll
@@ -0,0 +1,30 @@
+; RUN: not llvm-as -disable-output < %s 2>&1 | FileCheck %s
+
+declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cta(ptr addrspace(3) writeonly, ptr addrspace(3), ptr addrspace(1) readonly, i32, i32, i32, i64, i1 immarg, i1 immarg, i32 immarg range(i32 0, 6))
+declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cta.relaxed(ptr addrspace(3) writeonly, ptr addrspace(3), ptr addrspace(1) readonly, i32, i32, i32, i64, i1 immarg, i1 immarg, i32 immarg range(i32 0, 4), i32 immarg range(i32 0, 6))
+
+define void @test_cp_async_bulk_g2s_cta(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i64 %ch) {
+ ; CHECK: immarg value 6 for arg 9 out of range [0,6)
+ tail call void @llvm.nvvm.cp.async.bulk.global.to.shared.cta(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 0, i32 0, i64 %ch, i1 0, i1 0, i32 6)
+
+ ; CHECK: immarg value -1 for arg 9 out of range [0,6)
+ tail call void @llvm.nvvm.cp.async.bulk.global.to.shared.cta(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 0, i32 0, i64 %ch, i1 0, i1 0, i32 -1)
+
+ ret void
+}
+
+define void @test_cp_async_bulk_g2s_cta_relaxed(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i64 %ch) {
+ ; CHECK: immarg value 4 for arg 9 out of range [0,4)
+ tail call void @llvm.nvvm.cp.async.bulk.global.to.shared.cta.relaxed(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 0, i32 0, i64 %ch, i1 0, i1 0, i32 4, i32 0)
+
+ ; CHECK: immarg value -1 for arg 9 out of range [0,4)
+ tail call void @llvm.nvvm.cp.async.bulk.global.to.shared.cta.relaxed(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 0, i32 0, i64 %ch, i1 0, i1 0, i32 -1, i32 0)
+
+ ; CHECK: immarg value 6 for arg 10 out of range [0,6)
+ tail call void @llvm.nvvm.cp.async.bulk.global.to.shared.cta.relaxed(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 0, i32 0, i64 %ch, i1 0, i1 0, i32 0, i32 6)
+
+ ; CHECK: immarg value -1 for arg 10 out of range [0,6)
+ tail call void @llvm.nvvm.cp.async.bulk.global.to.shared.cta.relaxed(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 0, i32 0, i64 %ch, i1 0, i1 0, i32 0, i32 -1)
+
+ ret void
+}
diff --git a/llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-cta-oob.ll b/llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-cta-oob.ll
new file mode 100644
index 0000000000000..6356984b3d944
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-cta-oob.ll
@@ -0,0 +1,27 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx92 | FileCheck %s
+; RUN: %if ptxas-sm_90 && ptxas-isa-9.2 %{ llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx92 | %ptxas-verify -arch=sm_90 %}
+
+declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cta(ptr addrspace(3), ptr addrspace(3), ptr addrspace(1), i32, i32, i32, i64, i1, i1, i32)
+
+define void @cp_async_bulk_g2s_cta_oob(ptr addrspace(1) %src, ptr addrspace(3) %bar, ptr addrspace(3) %dst, i32 %size, i32 %ibl, i32 %ibr, i64 %ch) {
+; CHECK-LABEL: cp_async_bulk_g2s_cta_oob(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-NEXT: .reg .b64 %rd<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b64 %rd1, [cp_async_bulk_g2s_cta_oob_param_0];
+; CHECK-NEXT: ld.param::func.b64 %rd2, [cp_async_bulk_g2s_cta_oob_param_1];
+; CHECK-NEXT: ld.param::func.b64 %rd3, [cp_async_bulk_g2s_cta_oob_param_2];
+; CHECK-NEXT: ld.param::func.b32 %r1, [cp_async_bulk_g2s_cta_oob_param_3];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cp_async_bulk_g2s_cta_oob_param_4];
+; CHECK-NEXT: ld.param::func.b32 %r3, [cp_async_bulk_g2s_cta_oob_param_5];
+; CHECK-NEXT: ld.param::func.b64 %rd4, [cp_async_bulk_g2s_cta_oob_param_6];
+; CHECK-NEXT: cp.async.bulk.shared::cta.global.mbarrier::complete_tx::bytes.ignore_oob [%rd3], [%rd1], %r1, %r2, %r3, [%rd2];
+; CHECK-NEXT: cp.async.bulk.shared::cta.global.mbarrier::complete_tx::bytes.L2::cache_hint.ignore_oob [%rd3], [%rd1], %r1, %r2, %r3, [%rd2], %rd4;
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cta(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 %ibl, i32 %ibr, i64 %ch, i1 true, i1 false, i32 0)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cta(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 %ibl, i32 %ibr, i64 %ch, i1 true, i1 true, i32 0)
+ ret void
+}
diff --git a/llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-cta-relaxed.ll b/llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-cta-relaxed.ll
new file mode 100644
index 0000000000000..e2f83a496f2f2
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-cta-relaxed.ll
@@ -0,0 +1,75 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90a -mattr=+ptx93 | FileCheck %s
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100f -mattr=+ptx93 | FileCheck %s
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_110f -mattr=+ptx93 | FileCheck %s
+; RUN: %if ptxas-sm_90a && ptxas-isa-9.3 %{ llc < %s -march=nvptx64 -mcpu=sm_90a -mattr=+ptx93 | %ptxas-verify -arch=sm_90a %}
+; RUN: %if ptxas-sm_100f && ptxas-isa-9.3 %{ llc < %s -march=nvptx64 -mcpu=sm_100f -mattr=+ptx93 | %ptxas-verify -arch=sm_100f %}
+; RUN: %if ptxas-sm_110f && ptxas-isa-9.3 %{ llc < %s -march=nvptx64 -mcpu=sm_110f -mattr=+ptx93 | %ptxas-verify -arch=sm_110f %}
+
+declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cta.relaxed(ptr addrspace(3), ptr addrspace(3), ptr addrspace(1), i32, i32, i32, i64, i1, i1, i32, i32)
+
+define void @cp_async_bulk_g2s_cta_relaxed_scopes(ptr addrspace(1) %src, ptr addrspace(3) %bar, ptr addrspace(3) %dst, i32 %size, i64 %ch) {
+; CHECK-LABEL: cp_async_bulk_g2s_cta_relaxed_scopes(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b64 %rd1, [cp_async_bulk_g2s_cta_relaxed_scopes_param_0];
+; CHECK-NEXT: ld.param::func.b64 %rd2, [cp_async_bulk_g2s_cta_relaxed_scopes_param_1];
+; CHECK-NEXT: ld.param::func.b64 %rd3, [cp_async_bulk_g2s_cta_relaxed_scopes_param_2];
+; CHECK-NEXT: ld.param::func.b32 %r1, [cp_async_bulk_g2s_cta_relaxed_scopes_param_3];
+; CHECK-NEXT: ld.param::func.b64 %rd4, [cp_async_bulk_g2s_cta_relaxed_scopes_param_4];
+; CHECK-NEXT: cp.async.bulk.relaxed.cta.shared::cta.global.mbarrier::complete_tx::bytes.b128 [%rd3], [%rd1], %r1, [%rd2];
+; CHECK-NEXT: cp.async.bulk.relaxed.cluster.shared::cta.global.mbarrier::complete_tx::bytes.b128 [%rd3], [%rd1], %r1, [%rd2];
+; CHECK-NEXT: cp.async.bulk.relaxed.gpu.shared::cta.global.mbarrier::complete_tx::bytes.b128 [%rd3], [%rd1], %r1, [%rd2];
+; CHECK-NEXT: cp.async.bulk.relaxed.sys.shared::cta.global.mbarrier::complete_tx::bytes.b128 [%rd3], [%rd1], %r1, [%rd2];
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cta.relaxed(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 0, i32 0, i64 %ch, i1 false, i1 false, i32 0, i32 0)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cta.relaxed(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 0, i32 0, i64 %ch, i1 false, i1 false, i32 1, i32 0)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cta.relaxed(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 0, i32 0, i64 %ch, i1 false, i1 false, i32 2, i32 0)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cta.relaxed(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 0, i32 0, i64 %ch, i1 false, i1 false, i32 3, i32 0)
+ ret void
+}
+
+define void @cp_async_bulk_g2s_cta_relaxed_ch(ptr addrspace(1) %src, ptr addrspace(3) %bar, ptr addrspace(3) %dst, i32 %size, i64 %ch) {
+; CHECK-LABEL: cp_async_bulk_g2s_cta_relaxed_ch(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b64 %rd1, [cp_async_bulk_g2s_cta_relaxed_ch_param_0];
+; CHECK-NEXT: ld.param::func.b64 %rd2, [cp_async_bulk_g2s_cta_relaxed_ch_param_1];
+; CHECK-NEXT: ld.param::func.b64 %rd3, [cp_async_bulk_g2s_cta_relaxed_ch_param_2];
+; CHECK-NEXT: ld.param::func.b32 %r1, [cp_async_bulk_g2s_cta_relaxed_ch_param_3];
+; CHECK-NEXT: ld.param::func.b64 %rd4, [cp_async_bulk_g2s_cta_relaxed_ch_param_4];
+; CHECK-NEXT: cp.async.bulk.relaxed.cta.shared::cta.global.mbarrier::complete_tx::bytes.L2::cache_hint.b128 [%rd3], [%rd1], %r1, [%rd2], %rd4;
+; CHECK-NEXT: cp.async.bulk.relaxed.sys.shared::cta.global.mbarrier::complete_tx::bytes.L2::cache_hint.b128 [%rd3], [%rd1], %r1, [%rd2], %rd4;
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cta.relaxed(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 0, i32 0, i64 %ch, i1 false, i1 true, i32 0, i32 0)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cta.relaxed(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 0, i32 0, i64 %ch, i1 false, i1 true, i32 3, i32 0)
+ ret void
+}
+
+define void @cp_async_bulk_g2s_cta_relaxed_oob(ptr addrspace(1) %src, ptr addrspace(3) %bar, ptr addrspace(3) %dst, i32 %size, i32 %ibl, i32 %ibr, i64 %ch) {
+; CHECK-LABEL: cp_async_bulk_g2s_cta_relaxed_oob(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-NEXT: .reg .b64 %rd<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b64 %rd1, [cp_async_bulk_g2s_cta_relaxed_oob_param_0];
+; CHECK-NEXT: ld.param::func.b64 %rd2, [cp_async_bulk_g2s_cta_relaxed_oob_param_1];
+; CHECK-NEXT: ld.param::func.b64 %rd3, [cp_async_bulk_g2s_cta_relaxed_oob_param_2];
+; CHECK-NEXT: ld.param::func.b32 %r1, [cp_async_bulk_g2s_cta_relaxed_oob_param_3];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cp_async_bulk_g2s_cta_relaxed_oob_param_4];
+; CHECK-NEXT: ld.param::func.b32 %r3, [cp_async_bulk_g2s_cta_relaxed_oob_param_5];
+; CHECK-NEXT: ld.param::func.b64 %rd4, [cp_async_bulk_g2s_cta_relaxed_oob_param_6];
+; CHECK-NEXT: cp.async.bulk.relaxed.gpu.shared::cta.global.mbarrier::complete_tx::bytes.ignore_oob.b128 [%rd3], [%rd1], %r1, %r2, %r3, [%rd2];
+; CHECK-NEXT: cp.async.bulk.relaxed.gpu.shared::cta.global.mbarrier::complete_tx::bytes.L2::cache_hint.ignore_oob.b128 [%rd3], [%rd1], %r1, %r2, %r3, [%rd2], %rd4;
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cta.relaxed(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 %ibl, i32 %ibr, i64 %ch, i1 true, i1 false, i32 2, i32 0)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cta.relaxed(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 %ibl, i32 %ibr, i64 %ch, i1 true, i1 true, i32 2, i32 0)
+ ret void
+}
diff --git a/llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-cta-validate-data.ll b/llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-cta-validate-data.ll
new file mode 100644
index 0000000000000..33031e4e40b91
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-cta-validate-data.ll
@@ -0,0 +1,56 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_107f -mattr=+ptx94 | FileCheck %s
+; RUN: %if ptxas-sm_107f && ptxas-isa-9.4 %{ llc < %s -march=nvptx64 -mcpu=sm_107f -mattr=+ptx94 | %ptxas-verify -arch=sm_107f %}
+
+declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cta(ptr addrspace(3), ptr addrspace(3), ptr addrspace(1), i32, i32, i32, i64, i1, i1, i32)
+declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cta.relaxed(ptr addrspace(3), ptr addrspace(3), ptr addrspace(1), i32, i32, i32, i64, i1, i1, i32, i32)
+
+define void @cp_async_bulk_g2s_cta_validate_data(ptr addrspace(1) %src, ptr addrspace(3) %bar, ptr addrspace(3) %dst, i32 %size, i64 %ch) {
+; CHECK-LABEL: cp_async_bulk_g2s_cta_validate_data(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b64 %rd1, [cp_async_bulk_g2s_cta_validate_data_param_0];
+; CHECK-NEXT: ld.param::func.b64 %rd2, [cp_async_bulk_g2s_cta_validate_data_param_1];
+; CHECK-NEXT: ld.param::func.b64 %rd3, [cp_async_bulk_g2s_cta_validate_data_param_2];
+; CHECK-NEXT: ld.param::func.b32 %r1, [cp_async_bulk_g2s_cta_validate_data_param_3];
+; CHECK-NEXT: ld.param::func.b64 %rd4, [cp_async_bulk_g2s_cta_validate_data_param_4];
+; CHECK-NEXT: cp.async.bulk.shared::cta.global.mbarrier::complete_tx::bytes [%rd3], [%rd1], %r1, [%rd2];
+; CHECK-NEXT: cp.async.bulk.shared::cta.global.mbarrier::complete_tx::bytes.mbarrier::report::validity::per_16bytes::80000000 [%rd3], [%rd1], %r1, [%rd2];
+; CHECK-NEXT: cp.async.bulk.shared::cta.global.mbarrier::complete_tx::bytes.mbarrier::report::validity::per_16bytes::8000 [%rd3], [%rd1], %r1, [%rd2];
+; CHECK-NEXT: cp.async.bulk.shared::cta.global.mbarrier::complete_tx::bytes.mbarrier::report::validity::per_16bytes::80 [%rd3], [%rd1], %r1, [%rd2];
+; CHECK-NEXT: cp.async.bulk.shared::cta.global.mbarrier::complete_tx::bytes.mbarrier::report::validity::per_16bytes::8 [%rd3], [%rd1], %r1, [%rd2];
+; CHECK-NEXT: cp.async.bulk.shared::cta.global.mbarrier::complete_tx::bytes.mbarrier::report::validity::per_element::ff [%rd3], [%rd1], %r1, [%rd2];
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cta(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 0, i32 0, i64 %ch, i1 false, i1 false, i32 0)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cta(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 0, i32 0, i64 %ch, i1 false, i1 false, i32 1)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cta(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 0, i32 0, i64 %ch, i1 false, i1 false, i32 2)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cta(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 0, i32 0, i64 %ch, i1 false, i1 false, i32 3)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cta(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 0, i32 0, i64 %ch, i1 false, i1 false, i32 4)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cta(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 0, i32 0, i64 %ch, i1 false, i1 false, i32 5)
+ ret void
+}
+
+define void @cp_async_bulk_g2s_cta_validate_data_ch(ptr addrspace(1) %src, ptr addrspace(3) %bar, ptr addrspace(3) %dst, i32 %size, i64 %ch) {
+; CHECK-LABEL: cp_async_bulk_g2s_cta_validate_data_ch(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b64 %rd1, [cp_async_bulk_g2s_cta_validate_data_ch_param_0];
+; CHECK-NEXT: ld.param::func.b64 %rd2, [cp_async_bulk_g2s_cta_validate_data_ch_param_1];
+; CHECK-NEXT: ld.param::func.b64 %rd3, [cp_async_bulk_g2s_cta_validate_data_ch_param_2];
+; CHECK-NEXT: ld.param::func.b32 %r1, [cp_async_bulk_g2s_cta_validate_data_ch_param_3];
+; CHECK-NEXT: ld.param::func.b64 %rd4, [cp_async_bulk_g2s_cta_validate_data_ch_param_4];
+; CHECK-NEXT: cp.async.bulk.shared::cta.global.mbarrier::complete_tx::bytes.mbarrier::report::validity::per_16bytes::8000.L2::cache_hint [%rd3], [%rd1], %r1, [%rd2], %rd4;
+; CHECK-NEXT: cp.async.bulk.relaxed.cta.shared::cta.global.mbarrier::complete_tx::bytes.mbarrier::report::validity::per_16bytes::8000.b128 [%rd3], [%rd1], %r1, [%rd2];
+; CHECK-NEXT: cp.async.bulk.relaxed.sys.shared::cta.global.mbarrier::complete_tx::bytes.mbarrier::report::validity::per_16bytes::8000.L2::cache_hint.b128 [%rd3], [%rd1], %r1, [%rd2], %rd4;
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cta(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 0, i32 0, i64 %ch, i1 false, i1 true, i32 2)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cta.relaxed(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 0, i32 0, i64 %ch, i1 false, i1 false, i32 0, i32 2)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cta.relaxed(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 0, i32 0, i64 %ch, i1 false, i1 true, i32 3, i32 2)
+ ret void
+}
diff --git a/llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-invalid.ll b/llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-invalid.ll
new file mode 100644
index 0000000000000..0b5082479ab5c
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-invalid.ll
@@ -0,0 +1,31 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: not llvm-as -disable-output < %s 2>&1 | FileCheck %s
+
+declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.relaxed.i16(ptr addrspace(7) writeonly, ptr addrspace(3), ptr addrspace(1) readonly, i32, i16, i64, i1 immarg, i1 immarg, i32 immarg range(i32 0, 4), i32 immarg range(i32 0, 6))
+declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.i16(ptr addrspace(7) writeonly, ptr addrspace(3), ptr addrspace(1) readonly, i32, i16, i64, i1 immarg, i1 immarg, i32 immarg range(i32 0, 6))
+
+define void @test_cp_async_bulk_g2s_cluster(ptr addrspace(7) %d, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch) {
+ ; CHECK: immarg value 4 for arg 8 out of range [0,4)
+ tail call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.relaxed.i16(ptr addrspace(7) %d, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 0, i1 0, i32 4, i32 0)
+
+ ; CHECK: immarg value -1 for arg 8 out of range [0,4)
+ tail call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.relaxed.i16(ptr addrspace(7) %d, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 0, i1 0, i32 -1, i32 0)
+
+ ; CHECK: immarg value 6 for arg 9 out of range [0,6)
+ tail call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.relaxed.i16(ptr addrspace(7) %d, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 0, i1 0, i32 0, i32 6)
+
+ ; CHECK: immarg value -1 for arg 9 out of range [0,6)
+ tail call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.relaxed.i16(ptr addrspace(7) %d, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 0, i1 0, i32 0, i32 -1)
+
+ ret void
+}
+
+define void @test_cp_async_bulk_g2s_cluster_weak_validate(ptr addrspace(7) %d, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch) {
+ ; CHECK: immarg value 6 for arg 8 out of range [0,6)
+ tail call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.i16(ptr addrspace(7) %d, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 0, i1 0, i32 6)
+
+ ; CHECK: immarg value -1 for arg 8 out of range [0,6)
+ tail call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.i16(ptr addrspace(7) %d, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 0, i1 0, i32 -1)
+
+ ret void
+}
diff --git a/llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-mc32.ll b/llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-mc32.ll
new file mode 100644
index 0000000000000..edd4275d2f663
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-mc32.ll
@@ -0,0 +1,52 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_107f -mattr=+ptx94 | FileCheck %s
+; RUN: %if ptxas-sm_107f && ptxas-isa-9.4 %{ llc < %s -march=nvptx64 -mcpu=sm_107f -mattr=+ptx94 | %ptxas-verify -arch=sm_107f %}
+
+declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.i32(ptr addrspace(7), ptr addrspace(3), ptr addrspace(1), i32, i32, i64, i1, i1, i32)
+declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.relaxed.i32(ptr addrspace(7), ptr addrspace(3), ptr addrspace(1), i32, i32, i64, i1, i1, i32, i32)
+
+define void @cp_async_bulk_g2s_mc32(ptr addrspace(1) %src, ptr addrspace(3) %bar, ptr addrspace(7) %dst, i32 %size, i32 %mc, i64 %ch) {
+; CHECK-LABEL: cp_async_bulk_g2s_mc32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b64 %rd1, [cp_async_bulk_g2s_mc32_param_0];
+; CHECK-NEXT: ld.param::func.b64 %rd2, [cp_async_bulk_g2s_mc32_param_1];
+; CHECK-NEXT: ld.param::func.b64 %rd3, [cp_async_bulk_g2s_mc32_param_2];
+; CHECK-NEXT: ld.param::func.b32 %r1, [cp_async_bulk_g2s_mc32_param_3];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cp_async_bulk_g2s_mc32_param_4];
+; CHECK-NEXT: ld.param::func.b64 %rd4, [cp_async_bulk_g2s_mc32_param_5];
+; CHECK-NEXT: cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%rd3], [%rd1], %r1, [%rd2];
+; CHECK-NEXT: cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.multicast::cluster::32b [%rd3], [%rd1], %r1, [%rd2], %r2;
+; CHECK-NEXT: cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.multicast::cluster::32b.L2::cache_hint [%rd3], [%rd1], %r1, [%rd2], %r2, %rd4;
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.i32(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 %mc, i64 %ch, i1 false, i1 false, i32 0)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.i32(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 %mc, i64 %ch, i1 true, i1 false, i32 0)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.i32(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 %mc, i64 %ch, i1 true, i1 true, i32 0)
+ ret void
+}
+
+define void @cp_async_bulk_g2s_relaxed_mc32(ptr addrspace(1) %src, ptr addrspace(3) %bar, ptr addrspace(7) %dst, i32 %size, i32 %mc, i64 %ch) {
+; CHECK-LABEL: cp_async_bulk_g2s_relaxed_mc32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b64 %rd1, [cp_async_bulk_g2s_relaxed_mc32_param_0];
+; CHECK-NEXT: ld.param::func.b64 %rd2, [cp_async_bulk_g2s_relaxed_mc32_param_1];
+; CHECK-NEXT: ld.param::func.b64 %rd3, [cp_async_bulk_g2s_relaxed_mc32_param_2];
+; CHECK-NEXT: ld.param::func.b32 %r1, [cp_async_bulk_g2s_relaxed_mc32_param_3];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cp_async_bulk_g2s_relaxed_mc32_param_4];
+; CHECK-NEXT: ld.param::func.b64 %rd4, [cp_async_bulk_g2s_relaxed_mc32_param_5];
+; CHECK-NEXT: cp.async.bulk.relaxed.cluster.shared::cluster.global.mbarrier::complete_tx::bytes.multicast::cluster::32b.b128 [%rd3], [%rd1], %r1, [%rd2], %r2;
+; CHECK-NEXT: cp.async.bulk.relaxed.cluster.shared::cluster.global.mbarrier::complete_tx::bytes.multicast::cluster::32b.L2::cache_hint.b128 [%rd3], [%rd1], %r1, [%rd2], %r2, %rd4;
+; CHECK-NEXT: cp.async.bulk.relaxed.cluster.shared::cluster.global.mbarrier::complete_tx::bytes.mbarrier::report::validity::per_16bytes::8.multicast::cluster::32b.L2::cache_hint.b128 [%rd3], [%rd1], %r1, [%rd2], %r2, %rd4;
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.relaxed.i32(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 %mc, i64 %ch, i1 true, i1 false, i32 1, i32 0)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.relaxed.i32(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 %mc, i64 %ch, i1 true, i1 true, i32 1, i32 0)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.relaxed.i32(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 %mc, i64 %ch, i1 true, i1 true, i32 1, i32 4)
+ ret void
+}
diff --git a/llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-relaxed.ll b/llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-relaxed.ll
new file mode 100644
index 0000000000000..db904af1756c0
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-relaxed.ll
@@ -0,0 +1,59 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90a -mattr=+ptx93 | FileCheck %s
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100f -mattr=+ptx93 | FileCheck %s
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_110f -mattr=+ptx93 | FileCheck %s
+; RUN: %if ptxas-sm_90a && ptxas-isa-9.3 %{ llc < %s -march=nvptx64 -mcpu=sm_90a -mattr=+ptx93 | %ptxas-verify -arch=sm_90a %}
+; RUN: %if ptxas-sm_100f && ptxas-isa-9.3 %{ llc < %s -march=nvptx64 -mcpu=sm_100f -mattr=+ptx93 | %ptxas-verify -arch=sm_100f %}
+; RUN: %if ptxas-sm_110f && ptxas-isa-9.3 %{ llc < %s -march=nvptx64 -mcpu=sm_110f -mattr=+ptx93 | %ptxas-verify -arch=sm_110f %}
+
+declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.relaxed.i16(ptr addrspace(7), ptr addrspace(3), ptr addrspace(1), i32, i16, i64, i1, i1, i32, i32)
+
+define void @cp_async_bulk_g2s_relaxed_scopes(ptr addrspace(1) %src, ptr addrspace(3) %bar, ptr addrspace(7) %dst, i32 %size, i16 %mc, i64 %ch) {
+; CHECK-LABEL: cp_async_bulk_g2s_relaxed_scopes(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b64 %rd1, [cp_async_bulk_g2s_relaxed_scopes_param_0];
+; CHECK-NEXT: ld.param::func.b64 %rd2, [cp_async_bulk_g2s_relaxed_scopes_param_1];
+; CHECK-NEXT: ld.param::func.b64 %rd3, [cp_async_bulk_g2s_relaxed_scopes_param_2];
+; CHECK-NEXT: ld.param::func.b32 %r1, [cp_async_bulk_g2s_relaxed_scopes_param_3];
+; CHECK-NEXT: ld.param::func.b16 %rs1, [cp_async_bulk_g2s_relaxed_scopes_param_4];
+; CHECK-NEXT: ld.param::func.b64 %rd4, [cp_async_bulk_g2s_relaxed_scopes_param_5];
+; CHECK-NEXT: cp.async.bulk.relaxed.cta.shared::cluster.global.mbarrier::complete_tx::bytes.b128 [%rd3], [%rd1], %r1, [%rd2];
+; CHECK-NEXT: cp.async.bulk.relaxed.cluster.shared::cluster.global.mbarrier::complete_tx::bytes.b128 [%rd3], [%rd1], %r1, [%rd2];
+; CHECK-NEXT: cp.async.bulk.relaxed.gpu.shared::cluster.global.mbarrier::complete_tx::bytes.b128 [%rd3], [%rd1], %r1, [%rd2];
+; CHECK-NEXT: cp.async.bulk.relaxed.sys.shared::cluster.global.mbarrier::complete_tx::bytes.b128 [%rd3], [%rd1], %r1, [%rd2];
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.relaxed.i16(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 false, i1 false, i32 0, i32 0)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.relaxed.i16(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 false, i1 false, i32 1, i32 0)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.relaxed.i16(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 false, i1 false, i32 2, i32 0)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.relaxed.i16(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 false, i1 false, i32 3, i32 0)
+ ret void
+}
+
+define void @cp_async_bulk_g2s_relaxed_mc_ch(ptr addrspace(1) %src, ptr addrspace(3) %bar, ptr addrspace(7) %dst, i32 %size, i16 %mc, i64 %ch) {
+; CHECK-LABEL: cp_async_bulk_g2s_relaxed_mc_ch(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b64 %rd1, [cp_async_bulk_g2s_relaxed_mc_ch_param_0];
+; CHECK-NEXT: ld.param::func.b64 %rd2, [cp_async_bulk_g2s_relaxed_mc_ch_param_1];
+; CHECK-NEXT: ld.param::func.b64 %rd3, [cp_async_bulk_g2s_relaxed_mc_ch_param_2];
+; CHECK-NEXT: ld.param::func.b32 %r1, [cp_async_bulk_g2s_relaxed_mc_ch_param_3];
+; CHECK-NEXT: ld.param::func.b16 %rs1, [cp_async_bulk_g2s_relaxed_mc_ch_param_4];
+; CHECK-NEXT: ld.param::func.b64 %rd4, [cp_async_bulk_g2s_relaxed_mc_ch_param_5];
+; CHECK-NEXT: cp.async.bulk.relaxed.cta.shared::cluster.global.mbarrier::complete_tx::bytes.multicast::cluster.b128 [%rd3], [%rd1], %r1, [%rd2], %rs1;
+; CHECK-NEXT: cp.async.bulk.relaxed.cta.shared::cluster.global.mbarrier::complete_tx::bytes.L2::cache_hint.b128 [%rd3], [%rd1], %r1, [%rd2], %rd4;
+; CHECK-NEXT: cp.async.bulk.relaxed.cta.shared::cluster.global.mbarrier::complete_tx::bytes.multicast::cluster.L2::cache_hint.b128 [%rd3], [%rd1], %r1, [%rd2], %rs1, %rd4;
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.relaxed.i16(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 true, i1 false, i32 0, i32 0)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.relaxed.i16(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 false, i1 true, i32 0, i32 0)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.relaxed.i16(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 true, i1 true, i32 0, i32 0)
+ ret void
+}
diff --git a/llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-validate-data.ll b/llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-validate-data.ll
new file mode 100644
index 0000000000000..fd79a37310d9a
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/cp-async-bulk-g2s-validate-data.ll
@@ -0,0 +1,58 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_107f -mattr=+ptx94 | FileCheck %s
+; RUN: %if ptxas-sm_107f && ptxas-isa-9.4 %{ llc < %s -march=nvptx64 -mcpu=sm_107f -mattr=+ptx94 | %ptxas-verify -arch=sm_107f %}
+
+declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.i16(ptr addrspace(7), ptr addrspace(3), ptr addrspace(1), i32, i16, i64, i1, i1, i32)
+declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.relaxed.i16(ptr addrspace(7), ptr addrspace(3), ptr addrspace(1), i32, i16, i64, i1, i1, i32, i32)
+
+define void @cp_async_bulk_g2s_validate_data(ptr addrspace(1) %src, ptr addrspace(3) %bar, ptr addrspace(7) %dst, i32 %size, i16 %mc, i64 %ch) {
+; CHECK-LABEL: cp_async_bulk_g2s_validate_data(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b64 %rd1, [cp_async_bulk_g2s_validate_data_param_0];
+; CHECK-NEXT: ld.param::func.b64 %rd2, [cp_async_bulk_g2s_validate_data_param_1];
+; CHECK-NEXT: ld.param::func.b64 %rd3, [cp_async_bulk_g2s_validate_data_param_2];
+; CHECK-NEXT: ld.param::func.b32 %r1, [cp_async_bulk_g2s_validate_data_param_3];
+; CHECK-NEXT: ld.param::func.b16 %rs1, [cp_async_bulk_g2s_validate_data_param_4];
+; CHECK-NEXT: ld.param::func.b64 %rd4, [cp_async_bulk_g2s_validate_data_param_5];
+; CHECK-NEXT: cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%rd3], [%rd1], %r1, [%rd2];
+; CHECK-NEXT: cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.mbarrier::report::validity::per_16bytes::80000000 [%rd3], [%rd1], %r1, [%rd2];
+; CHECK-NEXT: cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.mbarrier::report::validity::per_16bytes::8000.multicast::cluster.L2::cache_hint [%rd3], [%rd1], %r1, [%rd2], %rs1, %rd4;
+; CHECK-NEXT: cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.mbarrier::report::validity::per_16bytes::80.multicast::cluster [%rd3], [%rd1], %r1, [%rd2], %rs1;
+; CHECK-NEXT: cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.mbarrier::report::validity::per_16bytes::8.L2::cache_hint [%rd3], [%rd1], %r1, [%rd2], %rd4;
+; CHECK-NEXT: cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.mbarrier::report::validity::per_element::ff [%rd3], [%rd1], %r1, [%rd2];
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.i16(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 false, i1 false, i32 0)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.i16(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 false, i1 false, i32 1)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.i16(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 true, i1 true, i32 2)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.i16(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 true, i1 false, i32 3)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.i16(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 false, i1 true, i32 4)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.i16(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 false, i1 false, i32 5)
+ ret void
+}
+
+define void @cp_async_bulk_g2s_relaxed_validate_data(ptr addrspace(1) %src, ptr addrspace(3) %bar, ptr addrspace(7) %dst, i32 %size, i16 %mc, i64 %ch) {
+; CHECK-LABEL: cp_async_bulk_g2s_relaxed_validate_data(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b64 %rd1, [cp_async_bulk_g2s_relaxed_validate_data_param_0];
+; CHECK-NEXT: ld.param::func.b64 %rd2, [cp_async_bulk_g2s_relaxed_validate_data_param_1];
+; CHECK-NEXT: ld.param::func.b64 %rd3, [cp_async_bulk_g2s_relaxed_validate_data_param_2];
+; CHECK-NEXT: ld.param::func.b32 %r1, [cp_async_bulk_g2s_relaxed_validate_data_param_3];
+; CHECK-NEXT: ld.param::func.b16 %rs1, [cp_async_bulk_g2s_relaxed_validate_data_param_4];
+; CHECK-NEXT: ld.param::func.b64 %rd4, [cp_async_bulk_g2s_relaxed_validate_data_param_5];
+; CHECK-NEXT: cp.async.bulk.relaxed.cta.shared::cluster.global.mbarrier::complete_tx::bytes.mbarrier::report::validity::per_16bytes::8.b128 [%rd3], [%rd1], %r1, [%rd2];
+; CHECK-NEXT: cp.async.bulk.relaxed.cta.shared::cluster.global.mbarrier::complete_tx::bytes.mbarrier::report::validity::per_16bytes::8.L2::cache_hint.b128 [%rd3], [%rd1], %r1, [%rd2], %rd4;
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.relaxed.i16(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 false, i1 false, i32 0, i32 4)
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.relaxed.i16(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 false, i1 true, i32 0, i32 4)
+ ret void
+}
diff --git a/llvm/test/Verifier/NVPTX/cp-async-bulk-g2s-cta.ll b/llvm/test/Verifier/NVPTX/cp-async-bulk-g2s-cta.ll
new file mode 100644
index 0000000000000..20287cea61cab
--- /dev/null
+++ b/llvm/test/Verifier/NVPTX/cp-async-bulk-g2s-cta.ll
@@ -0,0 +1,16 @@
+; RUN: not llvm-as %s -o /dev/null 2>&1 | FileCheck %s
+
+declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cta(ptr addrspace(3) writeonly, ptr addrspace(3), ptr addrspace(1) readonly, i32, i32, i32, i64, i1 immarg, i1 immarg, i32 immarg range(i32 0, 6))
+declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cta.relaxed(ptr addrspace(3) writeonly, ptr addrspace(3), ptr addrspace(1) readonly, i32, i32, i32, i64, i1 immarg, i1 immarg, i32 immarg range(i32 0, 4), i32 immarg range(i32 0, 6))
+
+define void @test_cp_async_bulk_g2s_cta(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i64 %ch) {
+ ; CHECK: validate_pattern must be 0 (disabled) when ignore_oob is enabled
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cta(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 0, i32 0, i64 %ch, i1 1, i1 0, i32 2)
+ ret void
+}
+
+define void @test_cp_async_bulk_g2s_cta_relaxed(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i64 %ch) {
+ ; CHECK: validate_pattern must be 0 (disabled) when ignore_oob is enabled
+ call void @llvm.nvvm.cp.async.bulk.global.to.shared.cta.relaxed(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i32 0, i32 0, i64 %ch, i1 1, i1 1, i32 0, i32 4)
+ ret void
+}
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 80a0094190383..c04487aaaa210 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -4982,7 +4982,7 @@ def NVVM_CpAsyncBulkGlobalToSharedClusterOp :
string llvmBuilder = [{
auto [id, args] = NVVM::CpAsyncBulkGlobalToSharedClusterOp::getIntrinsicIDAndArgs(
*op, moduleTranslation, builder);
- createIntrinsicCall(builder, id, args);
+ createIntrinsicCall(builder, id, builder.getVoidTy(), args);
}];
}
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index 68f2a79fa7e36..6b60b6d0b8d31 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -4537,6 +4537,9 @@ mlir::NVVM::IDArgPair CpAsyncBulkGlobalToSharedClusterOp::getIntrinsicIDAndArgs(
llvm::Value *i16Unused = llvm::ConstantInt::get(builder.getInt16Ty(), 0);
args.push_back(hasMulticastMask ? mt.lookupValue(multicastMask)
: i16Unused);
+ } else {
+ args.push_back(builder.getInt32(0)); // ignore_bytes_left
+ args.push_back(builder.getInt32(0)); // ignore_bytes_right
}
// Cache hint, if available.
@@ -4545,10 +4548,15 @@ mlir::NVVM::IDArgPair CpAsyncBulkGlobalToSharedClusterOp::getIntrinsicIDAndArgs(
llvm::Value *i64Unused = llvm::ConstantInt::get(builder.getInt64Ty(), 0);
args.push_back(hasCacheHint ? mt.lookupValue(cacheHint) : i64Unused);
- // Flag arguments for multicast and cachehint.
+ // Flag arguments for multicast/ignore_oob and cachehint.
if (!isSharedCTA)
- args.push_back(builder.getInt1(hasMulticastMask));
- args.push_back(builder.getInt1(hasCacheHint));
+ args.push_back(builder.getInt1(hasMulticastMask)); // flag_mc
+ else
+ args.push_back(builder.getInt1(false)); // flag_oob
+ args.push_back(builder.getInt1(hasCacheHint)); // flag_ch
+
+ // validate_pattern = disabled.
+ args.push_back(builder.getInt32(0));
llvm::Intrinsic::ID id =
isSharedCTA
diff --git a/mlir/test/Target/LLVMIR/nvvm/tma_bulk_copy.mlir b/mlir/test/Target/LLVMIR/nvvm/tma_bulk_copy.mlir
index 240fab5b63908..441ae4348fb7e 100644
--- a/mlir/test/Target/LLVMIR/nvvm/tma_bulk_copy.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/tma_bulk_copy.mlir
@@ -2,10 +2,10 @@
// CHECK-LABEL: @llvm_nvvm_cp_async_bulk_global_to_shared_cluster
llvm.func @llvm_nvvm_cp_async_bulk_global_to_shared_cluster(%dst : !llvm.ptr<7>, %src : !llvm.ptr<1>, %mbar : !llvm.ptr<3>, %size : i32, %mc : i16, %ch : i64) {
- // CHECK: call 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 0, i64 0, i1 false, i1 false)
- // CHECK: call 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 0, i64 %[[CH:.*]], i1 false, i1 true)
- // CHECK: call 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 0, i1 true, i1 false)
- // CHECK: call 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 true, i1 true)
+ // CHECK: call 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 0, i64 0, i1 false, i1 false, /* validate_pattern=disabled */ i32 0)
+ // CHECK: call 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 0, i64 %[[CH:.*]], i1 false, i1 true, /* validate_pattern=disabled */ i32 0)
+ // CHECK: call 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 0, i1 true, i1 false, /* validate_pattern=disabled */ i32 0)
+ // CHECK: call 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 true, i1 true, /* validate_pattern=disabled */ i32 0)
nvvm.cp.async.bulk.shared.cluster.global %dst, %src, %mbar, %size : !llvm.ptr<7>, !llvm.ptr<1>
nvvm.cp.async.bulk.shared.cluster.global %dst, %src, %mbar, %size l2_cache_hint = %ch : !llvm.ptr<7>, !llvm.ptr<1>
@@ -18,8 +18,8 @@ llvm.func @llvm_nvvm_cp_async_bulk_global_to_shared_cluster(%dst : !llvm.ptr<7>,
// CHECK-LABEL: @llvm_nvvm_cp_async_bulk_global_to_shared_cta
llvm.func @llvm_nvvm_cp_async_bulk_global_to_shared_cta(%dst : !llvm.ptr<3>, %src : !llvm.ptr<1>, %mbar : !llvm.ptr<3>, %size : i32, %ch : i64) {
- // CHECK: call 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 0, i1 false)
- // CHECK: call 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 true)
+ // CHECK: call 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 0, i32 0, i64 0, i1 false, i1 false, /* validate_pattern=disabled */ i32 0)
+ // CHECK: call 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 0, i32 0, i64 %[[CH:.*]], i1 false, i1 true, /* validate_pattern=disabled */ i32 0)
nvvm.cp.async.bulk.shared.cluster.global %dst, %src, %mbar, %size : !llvm.ptr<3>, !llvm.ptr<1>
nvvm.cp.async.bulk.shared.cluster.global %dst, %src, %mbar, %size l2_cache_hint = %ch : !llvm.ptr<3>, !llvm.ptr<1>
More information about the llvm-commits
mailing list