[llvm] 6702faa - [NVPTX] Support 32-bit size operand for st.bulk (#217703)
via llvm-commits
llvm-commits at lists.llvm.org
Thu Sep 24 09:53:28 PDT 2026
Author: robertvirany
Date: 2026-09-24T10:53:20-06:00
New Revision: 6702faa9c52cedeaad7d8f8fbc82be44cbc8f637
URL: https://github.com/llvm/llvm-project/commit/6702faa9c52cedeaad7d8f8fbc82be44cbc8f637
DIFF: https://github.com/llvm/llvm-project/commit/6702faa9c52cedeaad7d8f8fbc82be44cbc8f637.diff
LOG: [NVPTX] Support 32-bit size operand for st.bulk (#217703)
PTX ISA 9.0 extends the `st.bulk` instruction to accept a 32-bit `size`
operand.
Add `i32` variants of the generic and `.shared::cta` NVVM intrinsics and
select them to `st.bulk` instructions using 32-bit registers. The new
patterns require PTX ISA 9.0 and `sm_100`.
The existing intrinsics continue to represent the original 64-bit forms.
Add code-generation coverage for both the default and short-pointer
ABIs.
Added:
llvm/test/CodeGen/NVPTX/st_bulk_i32.ll
Modified:
llvm/docs/NVPTXUsage.md
llvm/include/llvm/IR/IntrinsicsNVVM.td
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
llvm/test/CodeGen/NVPTX/st_bulk.ll
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
mlir/lib/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.cpp
mlir/test/Dialect/LLVMIR/nvvm.mlir
mlir/test/Target/LLVMIR/nvvmir-invalid.mlir
mlir/test/Target/LLVMIR/nvvmir.mlir
Removed:
################################################################################
diff --git a/llvm/docs/NVPTXUsage.md b/llvm/docs/NVPTXUsage.md
index 136b4cd954312..088035ee83f1f 100644
--- a/llvm/docs/NVPTXUsage.md
+++ b/llvm/docs/NVPTXUsage.md
@@ -4377,8 +4377,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:
@@ -4386,15 +4388,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 95e1def383428..002a8864d3959 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -3850,10 +3850,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 a3fc7407e595f..5327d39c34d03 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -4750,6 +4750,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 7b0848cef1186..d5108e1f32582 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -6961,17 +6961,26 @@ let isConvergent = true in {
// Bulk store instructions
def st_bulk_imm : TImmLeaf<i64, [{ return Imm == 0; }]>;
-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)]>;
-
- 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)]>;
-}
+class StBulkInst<NVPTXAddressSpace AS, string InstSuffix,
+ RegTyInfo Size, list<Predicate> Preds>
+ : BasicNVPTXInst<
+ (outs),
+ (ins ADDR:$dest_addr, Size.RC:$size, i64imm:$value),
+ "st.bulk" # InstSuffix,
+ [(IntrinsicInAS<int_nvvm_st_bulk, AS>
+ addr:$dest_addr, Size.Ty:$size, st_bulk_imm:$value)]> {
+ let Predicates = Preds;
+}
+
+multiclass StBulkInsts<RegTyInfo Size, list<Predicate> Preds> {
+ def _GENERIC :
+ StBulkInst<AddrSpaceGeneric, "", Size, Preds>;
+ def _SHARED_CTA :
+ StBulkInst<AddrSpaceShared, ".shared::cta", Size, Preds>;
+}
+
+defm INT_NVVM_ST_BULK_I64 : StBulkInsts<I64RT, [SM100]>;
+defm INT_NVVM_ST_BULK_I32 : StBulkInsts<I32RT, [PTX90, SM100]>;
//
// Asynchronous store instructions
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 28846502267d9..ada95402d5d70 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -5446,7 +5446,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 = [{
@@ -5460,13 +5460,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 612425754e620..f7f2503eee428 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 e619f6e3092a3..87a1ed369e747 100644
--- a/mlir/test/Dialect/LLVMIR/nvvm.mlir
+++ b/mlir/test/Dialect/LLVMIR/nvvm.mlir
@@ -596,11 +596,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 llvm-commits
mailing list