xref: /llvm-project/llvm/test/CodeGen/NVPTX/cp-async-bulk.ll (revision fa7f0e582bc25a91d89dab7c488a1619060f9bef) 
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_90 -mattr=+ptx80| FileCheck --check-prefixes=CHECK-PTX64 %s
3; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80 --nvptx-short-ptr| FileCheck --check-prefixes=CHECK-PTX-SHARED32 %s
4; RUN: %if ptxas-12.3 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80| %ptxas-verify -arch=sm_90 %}
5; RUN: %if ptxas-12.3 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80 --nvptx-short-ptr| %ptxas-verify -arch=sm_90 %}
6
7target triple = "nvptx64-nvidia-cuda"
8
9declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster(ptr addrspace(3), ptr addrspace(3), ptr addrspace(1), i32, i16, i64, i1, i1)
10declare void @llvm.nvvm.cp.async.bulk.shared.cta.to.global(ptr addrspace(1), ptr addrspace(3), i32, i64, i1)
11declare void @llvm.nvvm.cp.async.bulk.shared.cta.to.cluster(ptr addrspace(3), ptr addrspace(3), ptr addrspace(3), i32)
12declare void @llvm.nvvm.cp.async.bulk.prefetch.L2(ptr addrspace(1), i32, i64, i1)
13
14define void @cp_async_bulk_g2s(ptr addrspace(1) %src, ptr addrspace(3) %bar, ptr addrspace(3) %dst, i32 %size, i16 %mc, i64 %ch) {
15; CHECK-PTX64-LABEL: cp_async_bulk_g2s(
16; CHECK-PTX64:       {
17; CHECK-PTX64-NEXT:    .reg .b16 %rs<2>;
18; CHECK-PTX64-NEXT:    .reg .b32 %r<2>;
19; CHECK-PTX64-NEXT:    .reg .b64 %rd<5>;
20; CHECK-PTX64-EMPTY:
21; CHECK-PTX64-NEXT:  // %bb.0:
22; CHECK-PTX64-NEXT:    ld.param.u64 %rd1, [cp_async_bulk_g2s_param_0];
23; CHECK-PTX64-NEXT:    ld.param.u64 %rd2, [cp_async_bulk_g2s_param_1];
24; CHECK-PTX64-NEXT:    ld.param.u64 %rd3, [cp_async_bulk_g2s_param_2];
25; CHECK-PTX64-NEXT:    ld.param.u32 %r1, [cp_async_bulk_g2s_param_3];
26; CHECK-PTX64-NEXT:    cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%rd3], [%rd1], %r1, [%rd2];
27; CHECK-PTX64-NEXT:    ld.param.u64 %rd4, [cp_async_bulk_g2s_param_5];
28; CHECK-PTX64-NEXT:    cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.L2::cache_hint [%rd3], [%rd1], %r1, [%rd2], %rd4;
29; CHECK-PTX64-NEXT:    ld.param.u16 %rs1, [cp_async_bulk_g2s_param_4];
30; CHECK-PTX64-NEXT:    cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.multicast::cluster [%rd3], [%rd1], %r1, [%rd2], %rs1;
31; CHECK-PTX64-NEXT:    cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.multicast::cluster.L2::cache_hint [%rd3], [%rd1], %r1, [%rd2], %rs1, %rd4;
32; CHECK-PTX64-NEXT:    ret;
33;
34; CHECK-PTX-SHARED32-LABEL: cp_async_bulk_g2s(
35; CHECK-PTX-SHARED32:       {
36; CHECK-PTX-SHARED32-NEXT:    .reg .b16 %rs<2>;
37; CHECK-PTX-SHARED32-NEXT:    .reg .b32 %r<4>;
38; CHECK-PTX-SHARED32-NEXT:    .reg .b64 %rd<3>;
39; CHECK-PTX-SHARED32-EMPTY:
40; CHECK-PTX-SHARED32-NEXT:  // %bb.0:
41; CHECK-PTX-SHARED32-NEXT:    ld.param.u64 %rd1, [cp_async_bulk_g2s_param_0];
42; CHECK-PTX-SHARED32-NEXT:    ld.param.u32 %r1, [cp_async_bulk_g2s_param_1];
43; CHECK-PTX-SHARED32-NEXT:    ld.param.u32 %r2, [cp_async_bulk_g2s_param_2];
44; CHECK-PTX-SHARED32-NEXT:    ld.param.u32 %r3, [cp_async_bulk_g2s_param_3];
45; CHECK-PTX-SHARED32-NEXT:    cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%r2], [%rd1], %r3, [%r1];
46; CHECK-PTX-SHARED32-NEXT:    ld.param.u64 %rd2, [cp_async_bulk_g2s_param_5];
47; CHECK-PTX-SHARED32-NEXT:    cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.L2::cache_hint [%r2], [%rd1], %r3, [%r1], %rd2;
48; CHECK-PTX-SHARED32-NEXT:    ld.param.u16 %rs1, [cp_async_bulk_g2s_param_4];
49; CHECK-PTX-SHARED32-NEXT:    cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.multicast::cluster [%r2], [%rd1], %r3, [%r1], %rs1;
50; CHECK-PTX-SHARED32-NEXT:    cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.multicast::cluster.L2::cache_hint [%r2], [%rd1], %r3, [%r1], %rs1, %rd2;
51; CHECK-PTX-SHARED32-NEXT:    ret;
52  tail 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 0, i1 0)
53  tail 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 %ch, i1 0, i1 1)
54  tail 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 %mc, i64 0, i1 1, i1 0)
55  tail 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 %mc, i64 %ch, i1 1, i1 1)
56  ret void
57}
58
59define void @cp_async_bulk_s2g(ptr addrspace(3) %src, ptr addrspace(1) %dst, i32 %size, i64 %ch) {
60; CHECK-PTX64-LABEL: cp_async_bulk_s2g(
61; CHECK-PTX64:       {
62; CHECK-PTX64-NEXT:    .reg .b32 %r<2>;
63; CHECK-PTX64-NEXT:    .reg .b64 %rd<4>;
64; CHECK-PTX64-EMPTY:
65; CHECK-PTX64-NEXT:  // %bb.0:
66; CHECK-PTX64-NEXT:    ld.param.u64 %rd1, [cp_async_bulk_s2g_param_0];
67; CHECK-PTX64-NEXT:    ld.param.u64 %rd2, [cp_async_bulk_s2g_param_1];
68; CHECK-PTX64-NEXT:    ld.param.u32 %r1, [cp_async_bulk_s2g_param_2];
69; CHECK-PTX64-NEXT:    cp.async.bulk.global.shared::cta.bulk_group [%rd2], [%rd1], %r1;
70; CHECK-PTX64-NEXT:    ld.param.u64 %rd3, [cp_async_bulk_s2g_param_3];
71; CHECK-PTX64-NEXT:    cp.async.bulk.global.shared::cta.bulk_group.L2::cache_hint [%rd2], [%rd1], %r1, %rd3;
72; CHECK-PTX64-NEXT:    ret;
73;
74; CHECK-PTX-SHARED32-LABEL: cp_async_bulk_s2g(
75; CHECK-PTX-SHARED32:       {
76; CHECK-PTX-SHARED32-NEXT:    .reg .b32 %r<3>;
77; CHECK-PTX-SHARED32-NEXT:    .reg .b64 %rd<3>;
78; CHECK-PTX-SHARED32-EMPTY:
79; CHECK-PTX-SHARED32-NEXT:  // %bb.0:
80; CHECK-PTX-SHARED32-NEXT:    ld.param.u32 %r1, [cp_async_bulk_s2g_param_0];
81; CHECK-PTX-SHARED32-NEXT:    ld.param.u64 %rd1, [cp_async_bulk_s2g_param_1];
82; CHECK-PTX-SHARED32-NEXT:    ld.param.u32 %r2, [cp_async_bulk_s2g_param_2];
83; CHECK-PTX-SHARED32-NEXT:    cp.async.bulk.global.shared::cta.bulk_group [%rd1], [%r1], %r2;
84; CHECK-PTX-SHARED32-NEXT:    ld.param.u64 %rd2, [cp_async_bulk_s2g_param_3];
85; CHECK-PTX-SHARED32-NEXT:    cp.async.bulk.global.shared::cta.bulk_group.L2::cache_hint [%rd1], [%r1], %r2, %rd2;
86; CHECK-PTX-SHARED32-NEXT:    ret;
87  tail call void @llvm.nvvm.cp.async.bulk.shared.cta.to.global(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i64 0, i1 0)
88  tail call void @llvm.nvvm.cp.async.bulk.shared.cta.to.global(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i64 %ch, i1 1)
89  ret void
90}
91
92define void @cp_async_bulk_cta_to_cluster(ptr addrspace(3) %src, ptr addrspace(3) %bar, ptr addrspace(3) %dst, i32 %size) {
93; CHECK-PTX64-LABEL: cp_async_bulk_cta_to_cluster(
94; CHECK-PTX64:       {
95; CHECK-PTX64-NEXT:    .reg .b32 %r<2>;
96; CHECK-PTX64-NEXT:    .reg .b64 %rd<4>;
97; CHECK-PTX64-EMPTY:
98; CHECK-PTX64-NEXT:  // %bb.0:
99; CHECK-PTX64-NEXT:    ld.param.u64 %rd1, [cp_async_bulk_cta_to_cluster_param_0];
100; CHECK-PTX64-NEXT:    ld.param.u64 %rd2, [cp_async_bulk_cta_to_cluster_param_1];
101; CHECK-PTX64-NEXT:    ld.param.u64 %rd3, [cp_async_bulk_cta_to_cluster_param_2];
102; CHECK-PTX64-NEXT:    ld.param.u32 %r1, [cp_async_bulk_cta_to_cluster_param_3];
103; CHECK-PTX64-NEXT:    cp.async.bulk.shared::cluster.shared::cta.mbarrier::complete_tx::bytes [%rd3], [%rd1], %r1, [%rd2];
104; CHECK-PTX64-NEXT:    ret;
105;
106; CHECK-PTX-SHARED32-LABEL: cp_async_bulk_cta_to_cluster(
107; CHECK-PTX-SHARED32:       {
108; CHECK-PTX-SHARED32-NEXT:    .reg .b32 %r<5>;
109; CHECK-PTX-SHARED32-EMPTY:
110; CHECK-PTX-SHARED32-NEXT:  // %bb.0:
111; CHECK-PTX-SHARED32-NEXT:    ld.param.u32 %r1, [cp_async_bulk_cta_to_cluster_param_0];
112; CHECK-PTX-SHARED32-NEXT:    ld.param.u32 %r2, [cp_async_bulk_cta_to_cluster_param_1];
113; CHECK-PTX-SHARED32-NEXT:    ld.param.u32 %r3, [cp_async_bulk_cta_to_cluster_param_2];
114; CHECK-PTX-SHARED32-NEXT:    ld.param.u32 %r4, [cp_async_bulk_cta_to_cluster_param_3];
115; CHECK-PTX-SHARED32-NEXT:    cp.async.bulk.shared::cluster.shared::cta.mbarrier::complete_tx::bytes [%r3], [%r1], %r4, [%r2];
116; CHECK-PTX-SHARED32-NEXT:    ret;
117  tail call void @llvm.nvvm.cp.async.bulk.shared.cta.to.cluster(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr addrspace(3) %src, i32 %size)
118  ret void
119}
120
121define void @cp_async_bulk_prefetch(ptr addrspace(1) %src, i32 %size, i64 %ch) {
122; CHECK-PTX64-LABEL: cp_async_bulk_prefetch(
123; CHECK-PTX64:       {
124; CHECK-PTX64-NEXT:    .reg .b32 %r<2>;
125; CHECK-PTX64-NEXT:    .reg .b64 %rd<3>;
126; CHECK-PTX64-EMPTY:
127; CHECK-PTX64-NEXT:  // %bb.0:
128; CHECK-PTX64-NEXT:    ld.param.u64 %rd1, [cp_async_bulk_prefetch_param_0];
129; CHECK-PTX64-NEXT:    ld.param.u32 %r1, [cp_async_bulk_prefetch_param_1];
130; CHECK-PTX64-NEXT:    ld.param.u64 %rd2, [cp_async_bulk_prefetch_param_2];
131; CHECK-PTX64-NEXT:    cp.async.bulk.prefetch.L2.global.L2::cache_hint [%rd1], %r1, %rd2;
132; CHECK-PTX64-NEXT:    cp.async.bulk.prefetch.L2.global [%rd1], %r1;
133; CHECK-PTX64-NEXT:    ret;
134  tail call void @llvm.nvvm.cp.async.bulk.prefetch.L2(ptr addrspace(1) %src, i32 %size, i64 %ch, i1 1)
135  tail call void @llvm.nvvm.cp.async.bulk.prefetch.L2(ptr addrspace(1) %src, i32 %size, i64 0, i1 0)
136  ret void
137}
138