Skip to content

Commit 9a5a8c9

Browse files
authored
[NVPTX] Add intrinsics for st.bulk instruction (#128856)
Adds NVVM intrinsics and NVPTX codegen for the `st.bulk` instruction introduced in ptx8.6 for sm_100. Tests added in `CodeGen/NVPTX/st_bulk.ll` and verified through ptxas 12.8.0. PTX Spec Reference: https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-st-bulk
1 parent 494fe0b commit 9a5a8c9

File tree

4 files changed

+106
-0
lines changed

4 files changed

+106
-0
lines changed

llvm/docs/NVPTXUsage.rst

Lines changed: 32 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1446,6 +1446,38 @@ The last argument `i1 %unpack` is a compile-time constant which when set, indica
14461446
For more information, refer to the
14471447
`PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-instructions-tcgen05-st>`__.
14481448

1449+
Store Intrinsics
1450+
----------------
1451+
1452+
'``llvm.nvvm.st.bulk.*``'
1453+
^^^^^^^^^^^^^^^^^^^^^^^^^
1454+
1455+
Syntax:
1456+
"""""""
1457+
1458+
.. code-block:: llvm
1459+
1460+
declare void @llvm.nvvm.st.bulk(ptr addrspace(1) %dst, i64 %size, i64 immarg %initval)
1461+
declare void @llvm.nvvm.st.bulk.shared.cta(ptr addrspace(3) %dst, i64 %size, i64 immarg %initval)
1462+
1463+
Overview:
1464+
"""""""""
1465+
1466+
The '``@llvm.nvvm.st.bulk.*``' intrinsics initialize a region of shared memory
1467+
starting from the location specified by the destination address operand `%dst`.
1468+
1469+
The integer operand `%size` specifies the amount of memory to be initialized in
1470+
terms of number of bytes and must be a multiple of 8. Otherwise, the behavior
1471+
is undefined.
1472+
1473+
The integer immediate operand `%initval` specifies the initialization value for
1474+
the memory locations. The only numeric value allowed is 0.
1475+
1476+
The ``@llvm.nvvm.st.bulk.shared.cta`` and ``@llvm.nvvm.st.bulk`` intrinsics are
1477+
similar but the latter uses generic addressing (see `Generic Addressing <https://docs.nvidia.com/cuda/parallel-thread-execution/#generic-addressing>`__).
1478+
1479+
For more information, refer `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-st-bulk>`__.
1480+
14491481
Other Intrinsics
14501482
----------------
14511483

llvm/include/llvm/IR/IntrinsicsNVVM.td

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -5269,4 +5269,18 @@ foreach shape = ["16x64b", "16x128b", "16x256b", "32x32b", "16x32bx2"] in {
52695269
}
52705270
}
52715271

5272+
//
5273+
// Bulk store intrinsics
5274+
//
5275+
5276+
def int_nvvm_st_bulk: DefaultAttrsIntrinsic<[],
5277+
[llvm_global_ptr_ty, llvm_i64_ty, llvm_i64_ty],
5278+
[IntrArgMemOnly, IntrWriteMem,
5279+
WriteOnly<ArgIndex<0>>, NoCapture<ArgIndex<0>>, ImmArg<ArgIndex<2>>]>;
5280+
5281+
def int_nvvm_st_bulk_shared_cta : DefaultAttrsIntrinsic<[],
5282+
[llvm_shared_ptr_ty, llvm_i64_ty, llvm_i64_ty],
5283+
[IntrArgMemOnly, IntrWriteMem,
5284+
WriteOnly<ArgIndex<0>>, NoCapture<ArgIndex<0>>, ImmArg<ArgIndex<2>>]>;
5285+
52725286
} // let TargetPrefix = "nvvm"

llvm/lib/Target/NVPTX/NVPTXIntrinsics.td

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -7732,3 +7732,17 @@ foreach shape = ["16x64b", "16x128b", "16x256b", "32x32b", "16x32bx2"] in {
77327732
}
77337733

77347734
} // isConvergent
7735+
7736+
// Bulk store instructions
7737+
7738+
def INT_NVVM_ST_BULK_GENERIC :
7739+
NVPTXInst<(outs), (ins ADDR:$dest_addr, Int64Regs:$size),
7740+
"st.bulk [$dest_addr], $size, 0;",
7741+
[(int_nvvm_st_bulk addr:$dest_addr, i64:$size, (i64 0))]>,
7742+
Requires<[hasSM<100>, hasPTX<86>]>;
7743+
7744+
def INT_NVVM_ST_BULK_SHARED_CTA:
7745+
NVPTXInst<(outs), (ins ADDR:$dest_addr, Int64Regs:$size),
7746+
"st.bulk.shared::cta [$dest_addr], $size, 0;",
7747+
[(int_nvvm_st_bulk_shared_cta addr:$dest_addr, i64:$size, (i64 0))]>,
7748+
Requires<[hasSM<100>, hasPTX<86>]>;

llvm/test/CodeGen/NVPTX/st_bulk.ll

Lines changed: 46 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,46 @@
1+
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
2+
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx86 | FileCheck --check-prefixes=CHECK,CHECK-PTX64 %s
3+
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx86 --nvptx-short-ptr | FileCheck --check-prefixes=CHECK,CHECK-PTX-SHARED32 %s
4+
; RUN: %if ptxas-12.8 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx86 | %ptxas-verify -arch=sm_100 %}
5+
; RUN: %if ptxas-12.8 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx86 --nvptx-short-ptr | %ptxas-verify -arch=sm_100 %}
6+
7+
declare void @llvm.nvvm.st.bulk(ptr addrspace(1), i64, i64)
8+
define void @st_bulk(ptr addrspace(1) %dest_addr, i64 %size) {
9+
; CHECK-LABEL: st_bulk(
10+
; CHECK: {
11+
; CHECK-NEXT: .reg .b64 %rd<3>;
12+
; CHECK-EMPTY:
13+
; CHECK-NEXT: // %bb.0:
14+
; CHECK-NEXT: ld.param.u64 %rd1, [st_bulk_param_0];
15+
; CHECK-NEXT: ld.param.u64 %rd2, [st_bulk_param_1];
16+
; CHECK-NEXT: st.bulk [%rd1], %rd2, 0;
17+
; CHECK-NEXT: ret;
18+
call void @llvm.nvvm.st.bulk(ptr addrspace(1) %dest_addr, i64 %size, i64 0)
19+
ret void
20+
}
21+
22+
declare void @llvm.nvvm.st.bulk.shared.cta(ptr addrspace(3), i64, i64)
23+
define void @st_bulk_shared_cta(ptr addrspace(3) %dest_addr, i64 %size) {
24+
; CHECK-PTX64-LABEL: st_bulk_shared_cta(
25+
; CHECK-PTX64: {
26+
; CHECK-PTX64-NEXT: .reg .b64 %rd<3>;
27+
; CHECK-PTX64-EMPTY:
28+
; CHECK-PTX64-NEXT: // %bb.0:
29+
; CHECK-PTX64-NEXT: ld.param.u64 %rd1, [st_bulk_shared_cta_param_0];
30+
; CHECK-PTX64-NEXT: ld.param.u64 %rd2, [st_bulk_shared_cta_param_1];
31+
; CHECK-PTX64-NEXT: st.bulk.shared::cta [%rd1], %rd2, 0;
32+
; CHECK-PTX64-NEXT: ret;
33+
;
34+
; CHECK-PTX-SHARED32-LABEL: st_bulk_shared_cta(
35+
; CHECK-PTX-SHARED32: {
36+
; CHECK-PTX-SHARED32-NEXT: .reg .b32 %r<2>;
37+
; CHECK-PTX-SHARED32-NEXT: .reg .b64 %rd<2>;
38+
; CHECK-PTX-SHARED32-EMPTY:
39+
; CHECK-PTX-SHARED32-NEXT: // %bb.0:
40+
; CHECK-PTX-SHARED32-NEXT: ld.param.u32 %r1, [st_bulk_shared_cta_param_0];
41+
; CHECK-PTX-SHARED32-NEXT: ld.param.u64 %rd1, [st_bulk_shared_cta_param_1];
42+
; CHECK-PTX-SHARED32-NEXT: st.bulk.shared::cta [%r1], %rd1, 0;
43+
; CHECK-PTX-SHARED32-NEXT: ret;
44+
call void @llvm.nvvm.st.bulk.shared.cta(ptr addrspace(3) %dest_addr, i64 %size, i64 0)
45+
ret void
46+
}

0 commit comments

Comments
 (0)