[llvm] [NVPTX] Support 32-bit size operand for st.bulk (PR #217703)
via llvm-commits
llvm-commits at lists.llvm.org
Fri Aug 21 10:27:40 PDT 2026
https://github.com/robertvirany updated https://github.com/llvm/llvm-project/pull/217703
>From 15ee3cc478df3dcb34ec95703b6dfc633c400773 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 +++++++++++++++++++++
6 files changed, 92 insertions(+), 18 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
+}
More information about the llvm-commits
mailing list