llvm/test/CodeGen/NVPTX/intrinsics-sm90.ll

   1 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx80| FileCheck --check-prefixes=CHECK %s
   2 ; RUN: %if ptxas-12.0 %{ llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx80| %ptxas-verify -arch=sm_90 %}
   3
   4 ; CHECK-LABEL: test_isspacep
   5 define i1 @test_isspacep_shared_cluster(ptr %p) {
   6 ; CHECK: isspacep.shared::cluster
   7   %a = tail call i1 @llvm.nvvm.isspacep.shared.cluster(ptr %p)
   8 ; CHECK: ret
   9   ret i1 %a
  10 }
  11
  12 ; CHECK-LABEL: test_mapa(
  13 define ptr @test_mapa(ptr %p, i32 %r) {
  14 ; CHECK64: mapa.u64
  15   %a = call ptr @llvm.nvvm.mapa(ptr %p, i32 %r)
  16   ret ptr %a
  17 }
  18
  19 ; CHECK-LABEL: test_mapa_shared_cluster(
  20 define ptr addrspace(3) @test_mapa_shared_cluster(ptr addrspace(3) %p, i32 %r) {
  21 ; CHECK: mapa.shared::cluster.u64
  22   %a = call ptr addrspace(3) @llvm.nvvm.mapa.shared.cluster(ptr addrspace(3) %p, i32 %r)
  23   ret ptr addrspace(3) %a
  24 }
  25
  26 ; CHECK-LABEL: test_getctarank(
  27 define i32 @test_getctarank(ptr %p) {
  28 ; CHECK: getctarank.u64
  29   %a = call i32 @llvm.nvvm.getctarank(ptr %p)
  30   ret i32 %a
  31 }
  32
  33 ; CHECK-LABEL: test_getctarank_shared_cluster(
  34 define i32 @test_getctarank_shared_cluster(ptr addrspace(3) %p) {
  35 ; CHECK64: getctarank.shared::cluster.u64
  36 ; CHECK32: getctarank.shared::cluster.u32
  37   %a = call i32 @llvm.nvvm.getctarank.shared.cluster(ptr addrspace(3) %p)
  38   ret i32 %a
  39 }
  40
  41 ; CHECK-LABEL: test_clusterid_x(
  42 define i32 @test_clusterid_x() {
  43 ; CHECK: mov.u32 %r{{[0-9]+}}, %clusterid.x;
  44 ; CHECK: ret;
  45         %x = call i32 @llvm.nvvm.read.ptx.sreg.clusterid.x()
  46         ret i32 %x
  47 }
  48 ; CHECK-LABEL: test_clusterid_y(
  49 define i32 @test_clusterid_y() {
  50 ; CHECK: mov.u32 %r{{[0-9]+}}, %clusterid.y;
  51 ; CHECK: ret;
  52         %x = call i32 @llvm.nvvm.read.ptx.sreg.clusterid.y()
  53         ret i32 %x
  54 }
  55 ; CHECK-LABEL: test_clusterid_z(
  56 define i32 @test_clusterid_z() {
  57 ; CHECK: mov.u32 %r{{[0-9]+}}, %clusterid.z;
  58 ; CHECK: ret;
  59         %x = call i32 @llvm.nvvm.read.ptx.sreg.clusterid.z()
  60         ret i32 %x
  61 }
  62 ; CHECK-LABEL: test_clusterid_w(
  63 define i32 @test_clusterid_w() {
  64 ; CHECK: mov.u32 %r{{[0-9]+}}, %clusterid.w;
  65 ; CHECK: ret;
  66         %x = call i32 @llvm.nvvm.read.ptx.sreg.clusterid.w()
  67         ret i32 %x
  68 }
  69
  70 ; CHECK-LABEL: test_nclusterid_x(
  71 define i32 @test_nclusterid_x() {
  72 ; CHECK: mov.u32 %r{{[0-9]+}}, %nclusterid.x;
  73 ; CHECK: ret;
  74         %x = call i32 @llvm.nvvm.read.ptx.sreg.nclusterid.x()
  75         ret i32 %x
  76 }
  77 ; CHECK-LABEL: test_nclusterid_y(
  78 define i32 @test_nclusterid_y() {
  79 ; CHECK: mov.u32 %r{{[0-9]+}}, %nclusterid.y;
  80 ; CHECK: ret;
  81         %x = call i32 @llvm.nvvm.read.ptx.sreg.nclusterid.y()
  82         ret i32 %x
  83 }
  84 ; CHECK-LABEL: test_nclusterid_z(
  85 define i32 @test_nclusterid_z() {
  86 ; CHECK: mov.u32 %r{{[0-9]+}}, %nclusterid.z;
  87 ; CHECK: ret;
  88         %x = call i32 @llvm.nvvm.read.ptx.sreg.nclusterid.z()
  89         ret i32 %x
  90 }
  91 ; CHECK-LABEL: test_nclusterid_w(
  92 define i32 @test_nclusterid_w() {
  93 ; CHECK: mov.u32 %r{{[0-9]+}}, %nclusterid.w;
  94 ; CHECK: ret;
  95         %x = call i32 @llvm.nvvm.read.ptx.sreg.nclusterid.w()
  96         ret i32 %x
  97 }
  98
  99 ; CHECK-LABEL: test_cluster_ctarank(
 100 define i32 @test_cluster_ctarank() {
 101 ; CHECK: mov.u32 %r{{[0-9]+}}, %cluster_ctarank;
 102 ; CHECK: ret;
 103         %x = call i32 @llvm.nvvm.read.ptx.sreg.cluster.ctarank()
 104         ret i32 %x
 105 }
 106
 107 ; CHECK-LABEL: test_cluster_nctarank(
 108 define i32 @test_cluster_nctarank() {
 109 ; CHECK: mov.u32 %r{{[0-9]+}}, %cluster_nctarank;
 110 ; CHECK: ret;
 111         %x = call i32 @llvm.nvvm.read.ptx.sreg.cluster.nctarank()
 112         ret i32 %x
 113 }
 114
 115 ; CHECK-LABEL: test_is_explicit_cluster(
 116 define i1 @test_is_explicit_cluster() {
 117 ; CHECK: mov.pred %p{{[0-9]+}}, %is_explicit_cluster;
 118 ; CHECK: ret;
 119         %x = call i1 @llvm.nvvm.is_explicit_cluster()
 120         ret i1 %x
 121 }
 122
 123 ; CHECK-LABEL: test_barrier_cluster(
 124 define void @test_barrier_cluster() {
 125 ; CHECK: barrier.cluster.arrive;
 126        call void @llvm.nvvm.barrier.cluster.arrive()
 127 ; CHECK: barrier.cluster.arrive.relaxed;
 128        call void @llvm.nvvm.barrier.cluster.arrive.relaxed()
 129 ; CHECK: barrier.cluster.wait;
 130        call void @llvm.nvvm.barrier.cluster.wait()
 131 ; CHECK: fence.sc.cluster
 132        call void @llvm.nvvm.fence.sc.cluster()
 133        ret void
 134 }
 135
 136 ; CHECK-LABEL: test_barrier_cluster_aligned(
 137 define void @test_barrier_cluster_aligned() {
 138 ; CHECK: barrier.cluster.arrive.aligned;
 139        call void @llvm.nvvm.barrier.cluster.arrive.aligned()
 140 ; CHECK: barrier.cluster.arrive.relaxed.aligned;
 141        call void @llvm.nvvm.barrier.cluster.arrive.relaxed.aligned()
 142 ; CHECK: barrier.cluster.wait.aligned;
 143        call void @llvm.nvvm.barrier.cluster.wait.aligned()
 144        ret void
 145 }
 146
 147 ; CHECK-LABEL: test_cp_async_bulk_commit_group(
 148 define void @test_cp_async_bulk_commit_group() {
 149 ; CHECK: cp.async.bulk.commit_group;
 150        call void @llvm.nvvm.cp.async.bulk.commit.group()
 151        ret void
 152 }
 153
 154 ; CHECK-LABEL: test_cp_async_bulk_wait_group(
 155 define void @test_cp_async_bulk_wait_group() {
 156 ; CHECK: cp.async.bulk.wait_group 8;
 157        call void @llvm.nvvm.cp.async.bulk.wait.group(i32 8)
 158 ; CHECK: cp.async.bulk.wait_group 0;
 159        call void @llvm.nvvm.cp.async.bulk.wait.group(i32 0)
 160        ret void
 161 }
 162
 163 ; CHECK-LABEL: test_cp_async_bulk_wait_group_read(
 164 define void @test_cp_async_bulk_wait_group_read() {
 165 ; CHECK: cp.async.bulk.wait_group.read 8;
 166        call void @llvm.nvvm.cp.async.bulk.wait.group.read(i32 8)
 167 ; CHECK: cp.async.bulk.wait_group.read 0;
 168        call void @llvm.nvvm.cp.async.bulk.wait.group.read(i32 0)
 169        ret void
 170 }
 171
 172 declare i1 @llvm.nvvm.isspacep.shared.cluster(ptr %p);
 173 declare ptr @llvm.nvvm.mapa(ptr %p, i32 %r);
 174 declare ptr addrspace(3) @llvm.nvvm.mapa.shared.cluster(ptr addrspace(3) %p, i32 %r);
 175 declare i32 @llvm.nvvm.getctarank(ptr %p);
 176 declare i32 @llvm.nvvm.getctarank.shared.cluster(ptr addrspace(3) %p);
 177 declare i32 @llvm.nvvm.read.ptx.sreg.clusterid.x()
 178 declare i32 @llvm.nvvm.read.ptx.sreg.clusterid.y()
 179 declare i32 @llvm.nvvm.read.ptx.sreg.clusterid.z()
 180 declare i32 @llvm.nvvm.read.ptx.sreg.clusterid.w()
 181 declare i32 @llvm.nvvm.read.ptx.sreg.nclusterid.x()
 182 declare i32 @llvm.nvvm.read.ptx.sreg.nclusterid.y()
 183 declare i32 @llvm.nvvm.read.ptx.sreg.nclusterid.z()
 184 declare i32 @llvm.nvvm.read.ptx.sreg.nclusterid.w()
 185 declare i32 @llvm.nvvm.read.ptx.sreg.cluster.ctarank()
 186 declare i32 @llvm.nvvm.read.ptx.sreg.cluster.nctarank()
 187 declare i1 @llvm.nvvm.is_explicit_cluster()
 188 declare void @llvm.nvvm.barrier.cluster.arrive()
 189 declare void @llvm.nvvm.barrier.cluster.arrive.relaxed()
 190 declare void @llvm.nvvm.barrier.cluster.wait()
 191 declare void @llvm.nvvm.barrier.cluster.arrive.aligned()
 192 declare void @llvm.nvvm.barrier.cluster.arrive.relaxed.aligned()
 193 declare void @llvm.nvvm.barrier.cluster.wait.aligned()
 194 declare void @llvm.nvvm.fence.sc.cluster()
 195 declare void @llvm.nvvm.cp.async.bulk.commit.group()
 196 declare void @llvm.nvvm.cp.async.bulk.wait.group(i32)
 197 declare void @llvm.nvvm.cp.async.bulk.wait.group.read(i32)