[Mlir-commits] [llvm] [mlir] [NVPTX] Support 32-bit size operand for st.bulk (PR #217703)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Fri Aug 21 10:30:33 PDT 2026
https://github.com/robertvirany updated https://github.com/llvm/llvm-project/pull/217703
>From 576a1adeb45906eb824eebe9d18a63158aaea414 Mon Sep 17 00:00:00 2001
From: Robert Virany <robertvirany at gmail.com>
Date: Thu, 20 Aug 2026 11:12:57 -0600
Subject: [PATCH] [NVPTX] Support 32-bit size operand for st.bulk
---
llvm/docs/NVPTXUsage.md | 16 ++++---
llvm/include/llvm/IR/IntrinsicsNVVM.td | 5 +-
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 14 ++++++
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 20 ++++++--
llvm/test/CodeGen/NVPTX/st_bulk.ll | 8 ++--
llvm/test/CodeGen/NVPTX/st_bulk_i32.ll | 47 +++++++++++++++++++
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 11 ++---
.../Dialect/NVVM/NVVMToLLVMIRTranslation.cpp | 9 ----
mlir/test/Dialect/LLVMIR/nvvm.mlir | 14 ++++--
mlir/test/Target/LLVMIR/nvvmir-invalid.mlir | 2 +-
mlir/test/Target/LLVMIR/nvvmir.mlir | 23 +++++----
11 files changed, 121 insertions(+), 48 deletions(-)
create mode 100644 llvm/test/CodeGen/NVPTX/st_bulk_i32.ll
diff --git a/llvm/docs/NVPTXUsage.md b/llvm/docs/NVPTXUsage.md
index efad04b638494..e7459b015a2d7 100644
--- a/llvm/docs/NVPTXUsage.md
+++ b/llvm/docs/NVPTXUsage.md
@@ -3607,8 +3607,10 @@ The following tables describes the possible values of the flag arguments
##### Syntax:
```llvm
-declare void @llvm.nvvm.st.bulk(ptr addrspace(1) %dst, i64 %size, i64 immarg %initval)
-declare void @llvm.nvvm.st.bulk.shared.cta(ptr addrspace(3) %dst, i64 %size, i64 immarg %initval)
+declare void @llvm.nvvm.st.bulk.p0.i32(ptr %dst, i32 %size, i64 immarg %initval)
+declare void @llvm.nvvm.st.bulk.p0.i64(ptr %dst, i64 %size, i64 immarg %initval)
+declare void @llvm.nvvm.st.bulk.p3.i32(ptr addrspace(3) %dst, i32 %size, i64 immarg %initval)
+declare void @llvm.nvvm.st.bulk.p3.i64(ptr addrspace(3) %dst, i64 %size, i64 immarg %initval)
```
##### Overview:
@@ -3616,15 +3618,15 @@ declare void @llvm.nvvm.st.bulk.shared.cta(ptr addrspace(3) %dst, i64 %size, i64
The '`@llvm.nvvm.st.bulk.*`' intrinsics initialize a region of shared memory
starting from the location specified by the destination address operand `%dst`.
-The integer operand `%size` specifies the amount of memory to be initialized in
-terms of number of bytes and must be a multiple of 8. Otherwise, the behavior
-is undefined.
+The integer operand `%size`, which may have type `i32` or `i64`, specifies the
+amount of memory to be initialized in terms of number of bytes and must be a
+multiple of 8. Otherwise, the behavior is undefined.
The integer immediate operand `%initval` specifies the initialization value for
the memory locations. The only numeric value allowed is 0.
-The `@llvm.nvvm.st.bulk.shared.cta` and `@llvm.nvvm.st.bulk` intrinsics are
-similar but the latter uses generic addressing (see [Generic Addressing](https://docs.nvidia.com/cuda/parallel-thread-execution/#generic-addressing)).
+The `@llvm.nvvm.st.bulk.p3.*` and `@llvm.nvvm.st.bulk.p0.*` intrinsics are
+similar, but the latter uses generic addressing (see [Generic Addressing](https://docs.nvidia.com/cuda/parallel-thread-execution/#generic-addressing)).
For more information, refer [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-st-bulk).
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 6d6671e79dd24..3061fa1ef7845 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -3345,10 +3345,7 @@ foreach shape = ["32x32b", "16x32bx2"] in {
let IntrProperties = [IntrArgMemOnly, IntrWriteMem, WriteOnly<ArgIndex<0>>,
NoCapture<ArgIndex<0>>, ImmArg<ArgIndex<2>>] in {
def int_nvvm_st_bulk :
- DefaultAttrsIntrinsic<[], [llvm_ptr_ty, llvm_i64_ty, llvm_i64_ty]>;
-
- def int_nvvm_st_bulk_shared_cta :
- DefaultAttrsIntrinsic<[], [llvm_shared_ptr_ty, llvm_i64_ty, llvm_i64_ty]>;
+ DefaultAttrsIntrinsic<[], [llvm_anyptr_ty, llvm_anyint_ty, llvm_i64_ty]>;
}
//
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 2b01aab98d54a..f85955f39da96 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -4716,6 +4716,20 @@ void NVPTXTargetLowering::getTgtMemIntrinsic(
return;
}
+ case Intrinsic::nvvm_st_bulk: {
+ Value *Dst = I.getArgOperand(0);
+ Value *Val = I.getArgOperand(2);
+ Info.opc = ISD::INTRINSIC_VOID;
+ Info.memVT = MVT::getVT(Val->getType());
+ Info.ptrVal = Dst;
+ Info.offset = 0;
+ Info.flags = MachineMemOperand::MOStore;
+ Info.align.reset();
+ Info.size = MemoryLocation::UnknownSize;
+ Infos.push_back(Info);
+ return;
+ }
+
case Intrinsic::nvvm_prefetch_tensormap: {
auto &DL = I.getDataLayout();
Info.opc = ISD::INTRINSIC_VOID;
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index df5248a5bf7b5..be1cef47617b3 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -6180,12 +6180,26 @@ let Predicates = [SM100] in {
def INT_NVVM_ST_BULK_GENERIC :
BasicNVPTXInst<(outs), (ins ADDR:$dest_addr, B64:$size, i64imm:$value),
"st.bulk",
- [(int_nvvm_st_bulk addr:$dest_addr, i64:$size, st_bulk_imm:$value)]>;
-
+ [(IntrinsicInAS<int_nvvm_st_bulk, AddrSpaceGeneric>
+ addr:$dest_addr, i64:$size, st_bulk_imm:$value)]>;
def INT_NVVM_ST_BULK_SHARED_CTA:
BasicNVPTXInst<(outs), (ins ADDR:$dest_addr, B64:$size, i64imm:$value),
"st.bulk.shared::cta",
- [(int_nvvm_st_bulk_shared_cta addr:$dest_addr, i64:$size, st_bulk_imm:$value)]>;
+ [(IntrinsicInAS<int_nvvm_st_bulk, AddrSpaceShared>
+ addr:$dest_addr, i64:$size, st_bulk_imm:$value)]>;
+}
+
+let Predicates = [PTX90, SM100] in {
+ def INT_NVVM_ST_BULK_GENERIC_I32 :
+ BasicNVPTXInst<(outs), (ins ADDR:$dest_addr, B32:$size, i64imm:$value),
+ "st.bulk",
+ [(IntrinsicInAS<int_nvvm_st_bulk, AddrSpaceGeneric>
+ addr:$dest_addr, i32:$size, st_bulk_imm:$value)]>;
+ def INT_NVVM_ST_BULK_SHARED_CTA_I32:
+ BasicNVPTXInst<(outs), (ins ADDR:$dest_addr, B32:$size, i64imm:$value),
+ "st.bulk.shared::cta",
+ [(IntrinsicInAS<int_nvvm_st_bulk, AddrSpaceShared>
+ addr:$dest_addr, i32:$size, st_bulk_imm:$value)]>;
}
//
diff --git a/llvm/test/CodeGen/NVPTX/st_bulk.ll b/llvm/test/CodeGen/NVPTX/st_bulk.ll
index 9ac0da5620991..ff963b2db8c15 100644
--- a/llvm/test/CodeGen/NVPTX/st_bulk.ll
+++ b/llvm/test/CodeGen/NVPTX/st_bulk.ll
@@ -4,7 +4,7 @@
; RUN: %if ptxas-sm_100 && ptxas-isa-8.6 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx86 | %ptxas-verify -arch=sm_100 %}
; RUN: %if ptxas-sm_100 && ptxas-isa-8.6 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx86 -target-abi=shortptr | %ptxas-verify -arch=sm_100 %}
-declare void @llvm.nvvm.st.bulk(ptr, i64, i64)
+declare void @llvm.nvvm.st.bulk.p0.i64(ptr, i64, i64)
define void @st_bulk(ptr %dest_addr, i64 %size) {
; CHECK-LABEL: st_bulk(
; CHECK: {
@@ -15,11 +15,11 @@ define void @st_bulk(ptr %dest_addr, i64 %size) {
; CHECK-NEXT: ld.param::func.b64 %rd2, [st_bulk_param_1];
; CHECK-NEXT: st.bulk [%rd1], %rd2, 0;
; CHECK-NEXT: ret;
- call void @llvm.nvvm.st.bulk(ptr %dest_addr, i64 %size, i64 0)
+ call void @llvm.nvvm.st.bulk.p0.i64(ptr %dest_addr, i64 %size, i64 0)
ret void
}
-declare void @llvm.nvvm.st.bulk.shared.cta(ptr addrspace(3), i64, i64)
+declare void @llvm.nvvm.st.bulk.p3.i64(ptr addrspace(3), i64, i64)
define void @st_bulk_shared_cta(ptr addrspace(3) %dest_addr, i64 %size) {
; CHECK-PTX64-LABEL: st_bulk_shared_cta(
; CHECK-PTX64: {
@@ -41,6 +41,6 @@ define void @st_bulk_shared_cta(ptr addrspace(3) %dest_addr, i64 %size) {
; CHECK-PTX-SHARED32-NEXT: ld.param::func.b64 %rd1, [st_bulk_shared_cta_param_1];
; CHECK-PTX-SHARED32-NEXT: st.bulk.shared::cta [%r1], %rd1, 0;
; CHECK-PTX-SHARED32-NEXT: ret;
- call void @llvm.nvvm.st.bulk.shared.cta(ptr addrspace(3) %dest_addr, i64 %size, i64 0)
+ call void @llvm.nvvm.st.bulk.p3.i64(ptr addrspace(3) %dest_addr, i64 %size, i64 0)
ret void
}
diff --git a/llvm/test/CodeGen/NVPTX/st_bulk_i32.ll b/llvm/test/CodeGen/NVPTX/st_bulk_i32.ll
new file mode 100644
index 0000000000000..b7b11050246a3
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/st_bulk_i32.ll
@@ -0,0 +1,47 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx90 | FileCheck --check-prefixes=CHECK,CHECK-PTX64 %s
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx90 -target-abi=shortptr | FileCheck --check-prefixes=CHECK,CHECK-PTX-SHARED32 %s
+; RUN: %if ptxas-sm_100 && ptxas-isa-9.0 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx90 | %ptxas-verify -arch=sm_100 %}
+; RUN: %if ptxas-sm_100 && ptxas-isa-9.0 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx90 -target-abi=shortptr | %ptxas-verify -arch=sm_100 %}
+
+declare void @llvm.nvvm.st.bulk.p0.i32(ptr, i32, i64)
+define void @st_bulk_i32(ptr %dest_addr, i32 %size) {
+; CHECK-LABEL: st_bulk_i32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b64 %rd1, [st_bulk_i32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r1, [st_bulk_i32_param_1];
+; CHECK-NEXT: st.bulk [%rd1], %r1, 0;
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.st.bulk.p0.i32(ptr %dest_addr, i32 %size, i64 0)
+ ret void
+}
+
+declare void @llvm.nvvm.st.bulk.p3.i32(ptr addrspace(3), i32, i64)
+define void @st_bulk_shared_cta_i32(ptr addrspace(3) %dest_addr, i32 %size) {
+; CHECK-PTX64-LABEL: st_bulk_shared_cta_i32(
+; CHECK-PTX64: {
+; CHECK-PTX64-NEXT: .reg .b32 %r<2>;
+; CHECK-PTX64-NEXT: .reg .b64 %rd<2>;
+; CHECK-PTX64-EMPTY:
+; CHECK-PTX64-NEXT: // %bb.0:
+; CHECK-PTX64-NEXT: ld.param::func.b64 %rd1, [st_bulk_shared_cta_i32_param_0];
+; CHECK-PTX64-NEXT: ld.param::func.b32 %r1, [st_bulk_shared_cta_i32_param_1];
+; CHECK-PTX64-NEXT: st.bulk.shared::cta [%rd1], %r1, 0;
+; CHECK-PTX64-NEXT: ret;
+;
+; CHECK-PTX-SHARED32-LABEL: st_bulk_shared_cta_i32(
+; CHECK-PTX-SHARED32: {
+; CHECK-PTX-SHARED32-NEXT: .reg .b32 %r<3>;
+; CHECK-PTX-SHARED32-EMPTY:
+; CHECK-PTX-SHARED32-NEXT: // %bb.0:
+; CHECK-PTX-SHARED32-NEXT: ld.param::func.b32 %r1, [st_bulk_shared_cta_i32_param_0];
+; CHECK-PTX-SHARED32-NEXT: ld.param::func.b32 %r2, [st_bulk_shared_cta_i32_param_1];
+; CHECK-PTX-SHARED32-NEXT: st.bulk.shared::cta [%r1], %r2, 0;
+; CHECK-PTX-SHARED32-NEXT: ret;
+ call void @llvm.nvvm.st.bulk.p3.i32(ptr addrspace(3) %dest_addr, i32 %size, i64 0)
+ ret void
+}
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 6ac489f9ad2bc..bea92b6062825 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -5239,7 +5239,7 @@ def NVVM_MatchSyncOp : NVVM_Op<"match.sync", [InferTypeOpAdaptor]>,
//===----------------------------------------------------------------------===//
def NVVM_BulkStoreOp: NVVM_Op<"st.bulk"> {
- let arguments = (ins AnyTypeOf<[LLVM_PointerGeneric, LLVM_PointerShared]>:$addr, I64:$size, DefaultValuedAttr<I64Attr, "0">:$initVal);
+ let arguments = (ins AnyTypeOf<[LLVM_PointerGeneric, LLVM_PointerShared]>:$addr, AnyTypeOf<[I32, I64]>:$size, DefaultValuedAttr<I64Attr, "0">:$initVal);
let summary = "Bulk Store Op";
let description = [{
@@ -5253,13 +5253,12 @@ def NVVM_BulkStoreOp: NVVM_Op<"st.bulk"> {
}];
string llvmBuilder = [{
- auto intId = getStBulkIntrinsicId(
- llvm::cast<LLVM::LLVMPointerType>(op.getAddr().getType()));
- createIntrinsicCall(builder, intId,
- {$addr, $size, builder.getInt64($initVal)});
+ createIntrinsicCall(builder, llvm::Intrinsic::nvvm_st_bulk,
+ builder.getVoidTy(),
+ {$addr, $size, builder.getInt64($initVal)});
}];
- let assemblyFormat = "$addr `,` `size` `=` $size (`,` `init` `=` $initVal^)? attr-dict `:` type($addr)";
+ let assemblyFormat = "$addr `,` `size` `=` $size (`,` `init` `=` $initVal^)? attr-dict `:` type($addr) `,` type($size)";
let hasVerifier = 1;
}
diff --git a/mlir/lib/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.cpp b/mlir/lib/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.cpp
index 4201aabc02c3d..00aded8a9b782 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.cpp
+++ b/mlir/lib/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.cpp
@@ -249,15 +249,6 @@ getStMatrixIntrinsicId(NVVM::MMALayout layout, int32_t num,
llvm_unreachable("unknown stmatrix kind");
}
-/// Return the intrinsic ID associated with st.bulk for the given address type.
-static llvm::Intrinsic::ID
-getStBulkIntrinsicId(LLVM::LLVMPointerType addrType) {
- bool isSharedMemory = addrType.getAddressSpace() ==
- static_cast<unsigned>(NVVM::NVVMMemorySpace::Shared);
- return isSharedMemory ? llvm::Intrinsic::nvvm_st_bulk_shared_cta
- : llvm::Intrinsic::nvvm_st_bulk;
-}
-
static unsigned getUnidirectionalFenceProxyID(NVVM::ProxyKind fromProxy,
NVVM::ProxyKind toProxy,
NVVM::MemScopeKind scope,
diff --git a/mlir/test/Dialect/LLVMIR/nvvm.mlir b/mlir/test/Dialect/LLVMIR/nvvm.mlir
index 70e7816c6c1ec..5e147c3dfa880 100644
--- a/mlir/test/Dialect/LLVMIR/nvvm.mlir
+++ b/mlir/test/Dialect/LLVMIR/nvvm.mlir
@@ -556,11 +556,15 @@ func.func @match_sync(%val32: i32, %val64: i64, %thread_mask: i32) {
}
// CHECK-LABEL: @st_bulk
-func.func @st_bulk(%addr_gen: !llvm.ptr, %addr_shared: !llvm.ptr<3>, %size: i64) {
- // CHECK: nvvm.st.bulk %{{.*}}, size = %{{.*}} : !llvm.ptr
- nvvm.st.bulk %addr_gen, size = %size, init = 0 : !llvm.ptr
- // CHECK: nvvm.st.bulk %{{.*}}, size = %{{.*}} : !llvm.ptr<3>
- nvvm.st.bulk %addr_shared, size = %size, init = 0 : !llvm.ptr<3>
+func.func @st_bulk(%addr_gen: !llvm.ptr, %addr_shared: !llvm.ptr<3>, %size32: i32, %size: i64) {
+ // CHECK: nvvm.st.bulk %{{.*}}, size = %{{.*}} : !llvm.ptr, i32
+ nvvm.st.bulk %addr_gen, size = %size32, init = 0 : !llvm.ptr, i32
+ // CHECK: nvvm.st.bulk %{{.*}}, size = %{{.*}} : !llvm.ptr<3>, i32
+ nvvm.st.bulk %addr_shared, size = %size32, init = 0 : !llvm.ptr<3>, i32
+ // CHECK: nvvm.st.bulk %{{.*}}, size = %{{.*}} : !llvm.ptr, i64
+ nvvm.st.bulk %addr_gen, size = %size, init = 0 : !llvm.ptr, i64
+ // CHECK: nvvm.st.bulk %{{.*}}, size = %{{.*}} : !llvm.ptr<3>, i64
+ nvvm.st.bulk %addr_shared, size = %size, init = 0 : !llvm.ptr<3>, i64
return
}
diff --git a/mlir/test/Target/LLVMIR/nvvmir-invalid.mlir b/mlir/test/Target/LLVMIR/nvvmir-invalid.mlir
index f48b7f48dad90..9027f760d3483 100644
--- a/mlir/test/Target/LLVMIR/nvvmir-invalid.mlir
+++ b/mlir/test/Target/LLVMIR/nvvmir-invalid.mlir
@@ -94,7 +94,7 @@ llvm.func @convert_float_to_tf32_no_rnd_mode(%src : f32) -> i32 {
llvm.func @nvvm_st_bulk_initval_nonzero(%addr : !llvm.ptr, %size : i64) {
// expected-error @below {{only 0 is supported for initVal, got 1}}
- nvvm.st.bulk %addr, size = %size, init = 1 : !llvm.ptr
+ nvvm.st.bulk %addr, size = %size, init = 1 : !llvm.ptr, i64
llvm.return
}
diff --git a/mlir/test/Target/LLVMIR/nvvmir.mlir b/mlir/test/Target/LLVMIR/nvvmir.mlir
index c3b9fb301c3f9..c10b55a074dd5 100644
--- a/mlir/test/Target/LLVMIR/nvvmir.mlir
+++ b/mlir/test/Target/LLVMIR/nvvmir.mlir
@@ -865,15 +865,20 @@ llvm.func @nvvm_match_sync(%mask: i32, %val32: i32, %val64: i64) {
// -----
// CHECK-LABEL: @nvvm_st_bulk
-llvm.func @nvvm_st_bulk(%addr_gen: !llvm.ptr, %addr_shared: !llvm.ptr<3>, %size: i64) {
- // CHECK: call void @llvm.nvvm.st.bulk(ptr %{{.*}}, i64 %{{.*}}, i64 0)
- nvvm.st.bulk %addr_gen, size = %size : !llvm.ptr
- // CHECK: call void @llvm.nvvm.st.bulk.shared.cta(ptr addrspace(3) %{{.*}}, i64 %{{.*}}, i64 0)
- nvvm.st.bulk %addr_shared, size = %size: !llvm.ptr<3>
- // CHECK: call void @llvm.nvvm.st.bulk(ptr %{{.*}}, i64 %{{.*}}, i64 0)
- nvvm.st.bulk %addr_gen, size = %size, init = 0 : !llvm.ptr
- // CHECK: call void @llvm.nvvm.st.bulk.shared.cta(ptr addrspace(3) %{{.*}}, i64 %{{.*}}, i64 0)
- nvvm.st.bulk %addr_shared, size = %size, init = 0: !llvm.ptr<3>
+llvm.func @nvvm_st_bulk(%addr_gen: !llvm.ptr, %addr_shared: !llvm.ptr<3>, %size32: i32, %size: i64) {
+ // CHECK: call void @llvm.nvvm.st.bulk.p0.i32(ptr %{{.*}}, i32 %{{.*}}, i64 0)
+ nvvm.st.bulk %addr_gen, size = %size32 : !llvm.ptr, i32
+ // CHECK: call void @llvm.nvvm.st.bulk.p3.i32(ptr addrspace(3) %{{.*}}, i32 %{{.*}}, i64 0)
+ nvvm.st.bulk %addr_shared, size = %size32 : !llvm.ptr<3>, i32
+
+ // CHECK: call void @llvm.nvvm.st.bulk.p0.i64(ptr %{{.*}}, i64 %{{.*}}, i64 0)
+ nvvm.st.bulk %addr_gen, size = %size : !llvm.ptr, i64
+ // CHECK: call void @llvm.nvvm.st.bulk.p3.i64(ptr addrspace(3) %{{.*}}, i64 %{{.*}}, i64 0)
+ nvvm.st.bulk %addr_shared, size = %size: !llvm.ptr<3>, i64
+ // CHECK: call void @llvm.nvvm.st.bulk.p0.i64(ptr %{{.*}}, i64 %{{.*}}, i64 0)
+ nvvm.st.bulk %addr_gen, size = %size, init = 0 : !llvm.ptr, i64
+ // CHECK: call void @llvm.nvvm.st.bulk.p3.i64(ptr addrspace(3) %{{.*}}, i64 %{{.*}}, i64 0)
+ nvvm.st.bulk %addr_shared, size = %size, init = 0 : !llvm.ptr<3>, i64
llvm.return
}
More information about the Mlir-commits
mailing list