llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-s2g.ll

   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
   7 target triple = "nvptx64-nvidia-cuda"
   8
   9 declare void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.1d(ptr addrspace(3) %s, ptr %tm, i32 %d0, i64 %ch, i1 %flag);
  10 declare void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.2d(ptr addrspace(3) %s, ptr %tm, i32 %d0, i32 %d1, i64 %ch, i1 %flag);
  11 declare void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.3d(ptr addrspace(3) %s, ptr %tm, i32 %d0, i32 %d1, i32 %d2, i64 %ch, i1 %flag);
  12 declare void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.4d(ptr addrspace(3) %s, ptr %tm, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i64 %ch, i1 %flag);
  13 declare void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.5d(ptr addrspace(3) %s, ptr %tm, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, i64 %ch, i1 %flag);
  14
  15 declare void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.3d(ptr addrspace(3) %s, ptr %tm, i32 %d0, i32 %d1, i32 %d2, i64 %ch, i1 %flag);
  16 declare void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.4d(ptr addrspace(3) %s, ptr %tm, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i64 %ch, i1 %flag);
  17 declare void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.5d(ptr addrspace(3) %s, ptr %tm, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, i64 %ch, i1 %flag);
  18
  19 ; CHECK-LABEL: cp_async_bulk_tensor_s2g_tile_1d
  20 define void @cp_async_bulk_tensor_s2g_tile_1d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i64 %ch) {
  21 ; CHECK-PTX64-LABEL: cp_async_bulk_tensor_s2g_tile_1d(
  22 ; CHECK-PTX64:       {
  23 ; CHECK-PTX64-NEXT:    .reg .b32 %r<2>;
  24 ; CHECK-PTX64-NEXT:    .reg .b64 %rd<4>;
  25 ; CHECK-PTX64-EMPTY:
  26 ; CHECK-PTX64-NEXT:  // %bb.0:
  27 ; CHECK-PTX64-NEXT:    ld.param.u64 %rd1, [cp_async_bulk_tensor_s2g_tile_1d_param_0];
  28 ; CHECK-PTX64-NEXT:    ld.param.u64 %rd2, [cp_async_bulk_tensor_s2g_tile_1d_param_1];
  29 ; CHECK-PTX64-NEXT:    ld.param.u32 %r1, [cp_async_bulk_tensor_s2g_tile_1d_param_2];
  30 ; CHECK-PTX64-NEXT:    cp.async.bulk.tensor.1d.global.shared::cta.tile.bulk_group [%rd2, {%r1}], [%rd1];
  31 ; CHECK-PTX64-NEXT:    ld.param.u64 %rd3, [cp_async_bulk_tensor_s2g_tile_1d_param_3];
  32 ; CHECK-PTX64-NEXT:    cp.async.bulk.tensor.1d.global.shared::cta.tile.bulk_group.L2::cache_hint [%rd2, {%r1}], [%rd1], %rd3;
  33 ; CHECK-PTX64-NEXT:    ret;
  34 ;
  35 ; CHECK-PTX-SHARED32-LABEL: cp_async_bulk_tensor_s2g_tile_1d(
  36 ; CHECK-PTX-SHARED32:       {
  37 ; CHECK-PTX-SHARED32-NEXT:    .reg .b32 %r<3>;
  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.u32 %r1, [cp_async_bulk_tensor_s2g_tile_1d_param_0];
  42 ; CHECK-PTX-SHARED32-NEXT:    ld.param.u64 %rd1, [cp_async_bulk_tensor_s2g_tile_1d_param_1];
  43 ; CHECK-PTX-SHARED32-NEXT:    ld.param.u32 %r2, [cp_async_bulk_tensor_s2g_tile_1d_param_2];
  44 ; CHECK-PTX-SHARED32-NEXT:    cp.async.bulk.tensor.1d.global.shared::cta.tile.bulk_group [%rd1, {%r2}], [%r1];
  45 ; CHECK-PTX-SHARED32-NEXT:    ld.param.u64 %rd2, [cp_async_bulk_tensor_s2g_tile_1d_param_3];
  46 ; CHECK-PTX-SHARED32-NEXT:    cp.async.bulk.tensor.1d.global.shared::cta.tile.bulk_group.L2::cache_hint [%rd1, {%r2}], [%r1], %rd2;
  47 ; CHECK-PTX-SHARED32-NEXT:    ret;
  48   tail call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.1d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i64 undef, i1 0)
  49   tail call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.1d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i64 %ch, i1 1)
  50   ret void
  51 }
  52
  53 ; CHECK-LABEL: cp_async_bulk_tensor_s2g_tile_2d
  54 define void @cp_async_bulk_tensor_s2g_tile_2d(i32 %flag, ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i64 %ch) {
  55 ; CHECK-PTX64-LABEL: cp_async_bulk_tensor_s2g_tile_2d(
  56 ; CHECK-PTX64:       {
  57 ; CHECK-PTX64-NEXT:    .reg .b32 %r<3>;
  58 ; CHECK-PTX64-NEXT:    .reg .b64 %rd<4>;
  59 ; CHECK-PTX64-EMPTY:
  60 ; CHECK-PTX64-NEXT:  // %bb.0:
  61 ; CHECK-PTX64-NEXT:    ld.param.u64 %rd1, [cp_async_bulk_tensor_s2g_tile_2d_param_1];
  62 ; CHECK-PTX64-NEXT:    ld.param.u64 %rd2, [cp_async_bulk_tensor_s2g_tile_2d_param_2];
  63 ; CHECK-PTX64-NEXT:    ld.param.u32 %r1, [cp_async_bulk_tensor_s2g_tile_2d_param_3];
  64 ; CHECK-PTX64-NEXT:    ld.param.u32 %r2, [cp_async_bulk_tensor_s2g_tile_2d_param_4];
  65 ; CHECK-PTX64-NEXT:    cp.async.bulk.tensor.2d.global.shared::cta.tile.bulk_group [%rd2, {%r1, %r2}], [%rd1];
  66 ; CHECK-PTX64-NEXT:    ld.param.u64 %rd3, [cp_async_bulk_tensor_s2g_tile_2d_param_5];
  67 ; CHECK-PTX64-NEXT:    cp.async.bulk.tensor.2d.global.shared::cta.tile.bulk_group.L2::cache_hint [%rd2, {%r1, %r2}], [%rd1], %rd3;
  68 ; CHECK-PTX64-NEXT:    ret;
  69 ;
  70 ; CHECK-PTX-SHARED32-LABEL: cp_async_bulk_tensor_s2g_tile_2d(
  71 ; CHECK-PTX-SHARED32:       {
  72 ; CHECK-PTX-SHARED32-NEXT:    .reg .b32 %r<4>;
  73 ; CHECK-PTX-SHARED32-NEXT:    .reg .b64 %rd<3>;
  74 ; CHECK-PTX-SHARED32-EMPTY:
  75 ; CHECK-PTX-SHARED32-NEXT:  // %bb.0:
  76 ; CHECK-PTX-SHARED32-NEXT:    ld.param.u32 %r1, [cp_async_bulk_tensor_s2g_tile_2d_param_1];
  77 ; CHECK-PTX-SHARED32-NEXT:    ld.param.u64 %rd1, [cp_async_bulk_tensor_s2g_tile_2d_param_2];
  78 ; CHECK-PTX-SHARED32-NEXT:    ld.param.u32 %r2, [cp_async_bulk_tensor_s2g_tile_2d_param_3];
  79 ; CHECK-PTX-SHARED32-NEXT:    ld.param.u32 %r3, [cp_async_bulk_tensor_s2g_tile_2d_param_4];
  80 ; CHECK-PTX-SHARED32-NEXT:    cp.async.bulk.tensor.2d.global.shared::cta.tile.bulk_group [%rd1, {%r2, %r3}], [%r1];
  81 ; CHECK-PTX-SHARED32-NEXT:    ld.param.u64 %rd2, [cp_async_bulk_tensor_s2g_tile_2d_param_5];
  82 ; CHECK-PTX-SHARED32-NEXT:    cp.async.bulk.tensor.2d.global.shared::cta.tile.bulk_group.L2::cache_hint [%rd1, {%r2, %r3}], [%r1], %rd2;
  83 ; CHECK-PTX-SHARED32-NEXT:    ret;
  84   tail call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.2d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i64 undef, i1 0)
  85   tail call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.2d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i64 %ch, i1 1)
  86   ret void
  87 }
  88
  89 ; CHECK-LABEL: cp_async_bulk_tensor_s2g_3d
  90 define void @cp_async_bulk_tensor_s2g_3d(i32 %flag, ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i64 %ch) {
  91 ; CHECK-PTX64-LABEL: cp_async_bulk_tensor_s2g_3d(
  92 ; CHECK-PTX64:       {
  93 ; CHECK-PTX64-NEXT:    .reg .b32 %r<4>;
  94 ; CHECK-PTX64-NEXT:    .reg .b64 %rd<4>;
  95 ; CHECK-PTX64-EMPTY:
  96 ; CHECK-PTX64-NEXT:  // %bb.0:
  97 ; CHECK-PTX64-NEXT:    ld.param.u64 %rd1, [cp_async_bulk_tensor_s2g_3d_param_1];
  98 ; CHECK-PTX64-NEXT:    ld.param.u64 %rd2, [cp_async_bulk_tensor_s2g_3d_param_2];
  99 ; CHECK-PTX64-NEXT:    ld.param.u32 %r1, [cp_async_bulk_tensor_s2g_3d_param_3];
 100 ; CHECK-PTX64-NEXT:    ld.param.u32 %r2, [cp_async_bulk_tensor_s2g_3d_param_4];
 101 ; CHECK-PTX64-NEXT:    ld.param.u32 %r3, [cp_async_bulk_tensor_s2g_3d_param_5];
 102 ; CHECK-PTX64-NEXT:    cp.async.bulk.tensor.3d.global.shared::cta.tile.bulk_group [%rd2, {%r1, %r2, %r3}], [%rd1];
 103 ; CHECK-PTX64-NEXT:    ld.param.u64 %rd3, [cp_async_bulk_tensor_s2g_3d_param_6];
 104 ; CHECK-PTX64-NEXT:    cp.async.bulk.tensor.3d.global.shared::cta.tile.bulk_group.L2::cache_hint [%rd2, {%r1, %r2, %r3}], [%rd1], %rd3;
 105 ; CHECK-PTX64-NEXT:    cp.async.bulk.tensor.3d.global.shared::cta.im2col_no_offs.bulk_group [%rd2, {%r1, %r2, %r3}], [%rd1];
 106 ; CHECK-PTX64-NEXT:    cp.async.bulk.tensor.3d.global.shared::cta.im2col_no_offs.bulk_group.L2::cache_hint [%rd2, {%r1, %r2, %r3}], [%rd1], %rd3;
 107 ; CHECK-PTX64-NEXT:    ret;
 108 ;
 109 ; CHECK-PTX-SHARED32-LABEL: cp_async_bulk_tensor_s2g_3d(
 110 ; CHECK-PTX-SHARED32:       {
 111 ; CHECK-PTX-SHARED32-NEXT:    .reg .b32 %r<5>;
 112 ; CHECK-PTX-SHARED32-NEXT:    .reg .b64 %rd<3>;
 113 ; CHECK-PTX-SHARED32-EMPTY:
 114 ; CHECK-PTX-SHARED32-NEXT:  // %bb.0:
 115 ; CHECK-PTX-SHARED32-NEXT:    ld.param.u32 %r1, [cp_async_bulk_tensor_s2g_3d_param_1];
 116 ; CHECK-PTX-SHARED32-NEXT:    ld.param.u64 %rd1, [cp_async_bulk_tensor_s2g_3d_param_2];
 117 ; CHECK-PTX-SHARED32-NEXT:    ld.param.u32 %r2, [cp_async_bulk_tensor_s2g_3d_param_3];
 118 ; CHECK-PTX-SHARED32-NEXT:    ld.param.u32 %r3, [cp_async_bulk_tensor_s2g_3d_param_4];
 119 ; CHECK-PTX-SHARED32-NEXT:    ld.param.u32 %r4, [cp_async_bulk_tensor_s2g_3d_param_5];
 120 ; CHECK-PTX-SHARED32-NEXT:    cp.async.bulk.tensor.3d.global.shared::cta.tile.bulk_group [%rd1, {%r2, %r3, %r4}], [%r1];
 121 ; CHECK-PTX-SHARED32-NEXT:    ld.param.u64 %rd2, [cp_async_bulk_tensor_s2g_3d_param_6];
 122 ; CHECK-PTX-SHARED32-NEXT:    cp.async.bulk.tensor.3d.global.shared::cta.tile.bulk_group.L2::cache_hint [%rd1, {%r2, %r3, %r4}], [%r1], %rd2;
 123 ; CHECK-PTX-SHARED32-NEXT:    cp.async.bulk.tensor.3d.global.shared::cta.im2col_no_offs.bulk_group [%rd1, {%r2, %r3, %r4}], [%r1];
 124 ; CHECK-PTX-SHARED32-NEXT:    cp.async.bulk.tensor.3d.global.shared::cta.im2col_no_offs.bulk_group.L2::cache_hint [%rd1, {%r2, %r3, %r4}], [%r1], %rd2;
 125 ; CHECK-PTX-SHARED32-NEXT:    ret;
 126   tail call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.3d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i64 undef, i1 0)
 127   tail call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.3d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i64 %ch, i1 1)
 128
 129   tail call void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.3d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i64 undef, i1 0)
 130   tail call void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.3d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i64 %ch, i1 1)
 131   ret void
 132 }
 133
 134 ; CHECK-LABEL: cp_async_bulk_tensor_s2g_4d
 135 define void @cp_async_bulk_tensor_s2g_4d(i32 %flag, ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i64 %ch) {
 136 ; CHECK-PTX64-LABEL: cp_async_bulk_tensor_s2g_4d(
 137 ; CHECK-PTX64:       {
 138 ; CHECK-PTX64-NEXT:    .reg .b32 %r<5>;
 139 ; CHECK-PTX64-NEXT:    .reg .b64 %rd<4>;
 140 ; CHECK-PTX64-EMPTY:
 141 ; CHECK-PTX64-NEXT:  // %bb.0:
 142 ; CHECK-PTX64-NEXT:    ld.param.u64 %rd1, [cp_async_bulk_tensor_s2g_4d_param_1];
 143 ; CHECK-PTX64-NEXT:    ld.param.u64 %rd2, [cp_async_bulk_tensor_s2g_4d_param_2];
 144 ; CHECK-PTX64-NEXT:    ld.param.u32 %r1, [cp_async_bulk_tensor_s2g_4d_param_3];
 145 ; CHECK-PTX64-NEXT:    ld.param.u32 %r2, [cp_async_bulk_tensor_s2g_4d_param_4];
 146 ; CHECK-PTX64-NEXT:    ld.param.u32 %r3, [cp_async_bulk_tensor_s2g_4d_param_5];
 147 ; CHECK-PTX64-NEXT:    ld.param.u32 %r4, [cp_async_bulk_tensor_s2g_4d_param_6];
 148 ; CHECK-PTX64-NEXT:    cp.async.bulk.tensor.4d.global.shared::cta.tile.bulk_group [%rd2, {%r1, %r2, %r3, %r4}], [%rd1];
 149 ; CHECK-PTX64-NEXT:    ld.param.u64 %rd3, [cp_async_bulk_tensor_s2g_4d_param_7];
 150 ; CHECK-PTX64-NEXT:    cp.async.bulk.tensor.4d.global.shared::cta.tile.bulk_group.L2::cache_hint [%rd2, {%r1, %r2, %r3, %r4}], [%rd1], %rd3;
 151 ; CHECK-PTX64-NEXT:    cp.async.bulk.tensor.4d.global.shared::cta.im2col_no_offs.bulk_group [%rd2, {%r1, %r2, %r3, %r4}], [%rd1];
 152 ; CHECK-PTX64-NEXT:    cp.async.bulk.tensor.4d.global.shared::cta.im2col_no_offs.bulk_group.L2::cache_hint [%rd2, {%r1, %r2, %r3, %r4}], [%rd1], %rd3;
 153 ; CHECK-PTX64-NEXT:    ret;
 154 ;
 155 ; CHECK-PTX-SHARED32-LABEL: cp_async_bulk_tensor_s2g_4d(
 156 ; CHECK-PTX-SHARED32:       {
 157 ; CHECK-PTX-SHARED32-NEXT:    .reg .b32 %r<6>;
 158 ; CHECK-PTX-SHARED32-NEXT:    .reg .b64 %rd<3>;
 159 ; CHECK-PTX-SHARED32-EMPTY:
 160 ; CHECK-PTX-SHARED32-NEXT:  // %bb.0:
 161 ; CHECK-PTX-SHARED32-NEXT:    ld.param.u32 %r1, [cp_async_bulk_tensor_s2g_4d_param_1];
 162 ; CHECK-PTX-SHARED32-NEXT:    ld.param.u64 %rd1, [cp_async_bulk_tensor_s2g_4d_param_2];
 163 ; CHECK-PTX-SHARED32-NEXT:    ld.param.u32 %r2, [cp_async_bulk_tensor_s2g_4d_param_3];
 164 ; CHECK-PTX-SHARED32-NEXT:    ld.param.u32 %r3, [cp_async_bulk_tensor_s2g_4d_param_4];
 165 ; CHECK-PTX-SHARED32-NEXT:    ld.param.u32 %r4, [cp_async_bulk_tensor_s2g_4d_param_5];
 166 ; CHECK-PTX-SHARED32-NEXT:    ld.param.u32 %r5, [cp_async_bulk_tensor_s2g_4d_param_6];
 167 ; CHECK-PTX-SHARED32-NEXT:    cp.async.bulk.tensor.4d.global.shared::cta.tile.bulk_group [%rd1, {%r2, %r3, %r4, %r5}], [%r1];
 168 ; CHECK-PTX-SHARED32-NEXT:    ld.param.u64 %rd2, [cp_async_bulk_tensor_s2g_4d_param_7];
 169 ; CHECK-PTX-SHARED32-NEXT:    cp.async.bulk.tensor.4d.global.shared::cta.tile.bulk_group.L2::cache_hint [%rd1, {%r2, %r3, %r4, %r5}], [%r1], %rd2;
 170 ; CHECK-PTX-SHARED32-NEXT:    cp.async.bulk.tensor.4d.global.shared::cta.im2col_no_offs.bulk_group [%rd1, {%r2, %r3, %r4, %r5}], [%r1];
 171 ; CHECK-PTX-SHARED32-NEXT:    cp.async.bulk.tensor.4d.global.shared::cta.im2col_no_offs.bulk_group.L2::cache_hint [%rd1, {%r2, %r3, %r4, %r5}], [%r1], %rd2;
 172 ; CHECK-PTX-SHARED32-NEXT:    ret;
 173   tail call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.4d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i64 undef, i1 0)
 174   tail call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.4d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i64 %ch, i1 1)
 175
 176   tail call void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.4d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i64 undef, i1 0)
 177   tail call void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.4d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i64 %ch, i1 1)
 178   ret void
 179 }
 180
 181 ; CHECK-LABEL: cp_async_bulk_tensor_s2g_5d
 182 define void @cp_async_bulk_tensor_s2g_5d(i32 %flag, ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, i64 %ch) {
 183 ; CHECK-PTX64-LABEL: cp_async_bulk_tensor_s2g_5d(
 184 ; CHECK-PTX64:       {
 185 ; CHECK-PTX64-NEXT:    .reg .b32 %r<6>;
 186 ; CHECK-PTX64-NEXT:    .reg .b64 %rd<4>;
 187 ; CHECK-PTX64-EMPTY:
 188 ; CHECK-PTX64-NEXT:  // %bb.0:
 189 ; CHECK-PTX64-NEXT:    ld.param.u64 %rd1, [cp_async_bulk_tensor_s2g_5d_param_1];
 190 ; CHECK-PTX64-NEXT:    ld.param.u64 %rd2, [cp_async_bulk_tensor_s2g_5d_param_2];
 191 ; CHECK-PTX64-NEXT:    ld.param.u32 %r1, [cp_async_bulk_tensor_s2g_5d_param_3];
 192 ; CHECK-PTX64-NEXT:    ld.param.u32 %r2, [cp_async_bulk_tensor_s2g_5d_param_4];
 193 ; CHECK-PTX64-NEXT:    ld.param.u32 %r3, [cp_async_bulk_tensor_s2g_5d_param_5];
 194 ; CHECK-PTX64-NEXT:    ld.param.u32 %r4, [cp_async_bulk_tensor_s2g_5d_param_6];
 195 ; CHECK-PTX64-NEXT:    ld.param.u32 %r5, [cp_async_bulk_tensor_s2g_5d_param_7];
 196 ; CHECK-PTX64-NEXT:    cp.async.bulk.tensor.5d.global.shared::cta.tile.bulk_group [%rd2, {%r1, %r2, %r3, %r4, %r5}], [%rd1];
 197 ; CHECK-PTX64-NEXT:    ld.param.u64 %rd3, [cp_async_bulk_tensor_s2g_5d_param_8];
 198 ; CHECK-PTX64-NEXT:    cp.async.bulk.tensor.5d.global.shared::cta.tile.bulk_group.L2::cache_hint [%rd2, {%r1, %r2, %r3, %r4, %r5}], [%rd1], %rd3;
 199 ; CHECK-PTX64-NEXT:    cp.async.bulk.tensor.5d.global.shared::cta.im2col_no_offs.bulk_group [%rd2, {%r1, %r2, %r3, %r4, %r5}], [%rd1];
 200 ; CHECK-PTX64-NEXT:    cp.async.bulk.tensor.5d.global.shared::cta.im2col_no_offs.bulk_group.L2::cache_hint [%rd2, {%r1, %r2, %r3, %r4, %r5}], [%rd1], %rd3;
 201 ; CHECK-PTX64-NEXT:    ret;
 202 ;
 203 ; CHECK-PTX-SHARED32-LABEL: cp_async_bulk_tensor_s2g_5d(
 204 ; CHECK-PTX-SHARED32:       {
 205 ; CHECK-PTX-SHARED32-NEXT:    .reg .b32 %r<7>;
 206 ; CHECK-PTX-SHARED32-NEXT:    .reg .b64 %rd<3>;
 207 ; CHECK-PTX-SHARED32-EMPTY:
 208 ; CHECK-PTX-SHARED32-NEXT:  // %bb.0:
 209 ; CHECK-PTX-SHARED32-NEXT:    ld.param.u32 %r1, [cp_async_bulk_tensor_s2g_5d_param_1];
 210 ; CHECK-PTX-SHARED32-NEXT:    ld.param.u64 %rd1, [cp_async_bulk_tensor_s2g_5d_param_2];
 211 ; CHECK-PTX-SHARED32-NEXT:    ld.param.u32 %r2, [cp_async_bulk_tensor_s2g_5d_param_3];
 212 ; CHECK-PTX-SHARED32-NEXT:    ld.param.u32 %r3, [cp_async_bulk_tensor_s2g_5d_param_4];
 213 ; CHECK-PTX-SHARED32-NEXT:    ld.param.u32 %r4, [cp_async_bulk_tensor_s2g_5d_param_5];
 214 ; CHECK-PTX-SHARED32-NEXT:    ld.param.u32 %r5, [cp_async_bulk_tensor_s2g_5d_param_6];
 215 ; CHECK-PTX-SHARED32-NEXT:    ld.param.u32 %r6, [cp_async_bulk_tensor_s2g_5d_param_7];
 216 ; CHECK-PTX-SHARED32-NEXT:    cp.async.bulk.tensor.5d.global.shared::cta.tile.bulk_group [%rd1, {%r2, %r3, %r4, %r5, %r6}], [%r1];
 217 ; CHECK-PTX-SHARED32-NEXT:    ld.param.u64 %rd2, [cp_async_bulk_tensor_s2g_5d_param_8];
 218 ; CHECK-PTX-SHARED32-NEXT:    cp.async.bulk.tensor.5d.global.shared::cta.tile.bulk_group.L2::cache_hint [%rd1, {%r2, %r3, %r4, %r5, %r6}], [%r1], %rd2;
 219 ; CHECK-PTX-SHARED32-NEXT:    cp.async.bulk.tensor.5d.global.shared::cta.im2col_no_offs.bulk_group [%rd1, {%r2, %r3, %r4, %r5, %r6}], [%r1];
 220 ; CHECK-PTX-SHARED32-NEXT:    cp.async.bulk.tensor.5d.global.shared::cta.im2col_no_offs.bulk_group.L2::cache_hint [%rd1, {%r2, %r3, %r4, %r5, %r6}], [%r1], %rd2;
 221 ; CHECK-PTX-SHARED32-NEXT:    ret;
 222   tail call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.5d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, i64 undef, i1 0)
 223   tail call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.5d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, i64 %ch, i1 1)
 224
 225   tail call void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.5d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, i64 undef, i1 0)
 226   tail call void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.5d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, i64 %ch, i1 1)
 227   ret void
 228 }