198 lines · plain
1; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80| FileCheck --check-prefixes=CHECK %s2; RUN: %if ptxas-sm_90 && ptxas-isa-8.0 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80| %ptxas-verify -arch=sm_90 %}3 4; CHECK-LABEL: test_isspacep5define i1 @test_isspacep_shared_cluster(ptr %p) {6; CHECK: isspacep.shared::cluster7 %a = tail call i1 @llvm.nvvm.isspacep.shared.cluster(ptr %p)8; CHECK: ret9 ret i1 %a10}11 12; CHECK-LABEL: test_mapa(13define ptr @test_mapa(ptr %p, i32 %r) {14; CHECK64: mapa.u6415 %a = call ptr @llvm.nvvm.mapa(ptr %p, i32 %r)16 ret ptr %a17}18 19; CHECK-LABEL: test_mapa_shared_cluster(20define ptr addrspace(3) @test_mapa_shared_cluster(ptr addrspace(3) %p, i32 %r) {21; CHECK: mapa.shared::cluster.u6422 %a = call ptr addrspace(3) @llvm.nvvm.mapa.shared.cluster(ptr addrspace(3) %p, i32 %r)23 ret ptr addrspace(3) %a24}25 26; CHECK-LABEL: test_getctarank(27define i32 @test_getctarank(ptr %p) {28; CHECK: getctarank.u6429 %a = call i32 @llvm.nvvm.getctarank(ptr %p)30 ret i32 %a31}32 33; CHECK-LABEL: test_getctarank_shared_cluster(34define i32 @test_getctarank_shared_cluster(ptr addrspace(3) %p) {35; CHECK64: getctarank.shared::cluster.u6436; CHECK32: getctarank.shared::cluster.u3237 %a = call i32 @llvm.nvvm.getctarank.shared.cluster(ptr addrspace(3) %p)38 ret i32 %a39}40 41; CHECK-LABEL: test_clusterid_x(42define 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 %x47}48; CHECK-LABEL: test_clusterid_y(49define 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 %x54}55; CHECK-LABEL: test_clusterid_z(56define 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 %x61}62; CHECK-LABEL: test_clusterid_w(63define 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 %x68}69 70; CHECK-LABEL: test_nclusterid_x(71define 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 %x76}77; CHECK-LABEL: test_nclusterid_y(78define 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 %x83}84; CHECK-LABEL: test_nclusterid_z(85define 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 %x90}91; CHECK-LABEL: test_nclusterid_w(92define 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 %x97}98 99; CHECK-LABEL: test_cluster_ctarank(100define 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 %x105}106 107; CHECK-LABEL: test_cluster_nctarank(108define 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 %x113}114 115; CHECK-LABEL: test_is_explicit_cluster(116define 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 %x121}122 123; CHECK-LABEL: test_barrier_cluster(124define 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.cluster132 call void @llvm.nvvm.fence.sc.cluster()133 ret void134}135 136; CHECK-LABEL: test_barrier_cluster_aligned(137define 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 void145}146 147; CHECK-LABEL: test_cp_async_bulk_commit_group(148define 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 void152}153 154; CHECK-LABEL: test_cp_async_bulk_wait_group(155define 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 void161}162 163; CHECK-LABEL: test_cp_async_bulk_wait_group_read(164define 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 void170}171 172declare i1 @llvm.nvvm.isspacep.shared.cluster(ptr %p);173declare ptr @llvm.nvvm.mapa(ptr %p, i32 %r);174declare ptr addrspace(3) @llvm.nvvm.mapa.shared.cluster(ptr addrspace(3) %p, i32 %r);175declare i32 @llvm.nvvm.getctarank(ptr %p);176declare i32 @llvm.nvvm.getctarank.shared.cluster(ptr addrspace(3) %p);177declare i32 @llvm.nvvm.read.ptx.sreg.clusterid.x()178declare i32 @llvm.nvvm.read.ptx.sreg.clusterid.y()179declare i32 @llvm.nvvm.read.ptx.sreg.clusterid.z()180declare i32 @llvm.nvvm.read.ptx.sreg.clusterid.w()181declare i32 @llvm.nvvm.read.ptx.sreg.nclusterid.x()182declare i32 @llvm.nvvm.read.ptx.sreg.nclusterid.y()183declare i32 @llvm.nvvm.read.ptx.sreg.nclusterid.z()184declare i32 @llvm.nvvm.read.ptx.sreg.nclusterid.w()185declare i32 @llvm.nvvm.read.ptx.sreg.cluster.ctarank()186declare i32 @llvm.nvvm.read.ptx.sreg.cluster.nctarank()187declare i1 @llvm.nvvm.is_explicit_cluster()188declare void @llvm.nvvm.barrier.cluster.arrive()189declare void @llvm.nvvm.barrier.cluster.arrive.relaxed()190declare void @llvm.nvvm.barrier.cluster.wait()191declare void @llvm.nvvm.barrier.cluster.arrive.aligned()192declare void @llvm.nvvm.barrier.cluster.arrive.relaxed.aligned()193declare void @llvm.nvvm.barrier.cluster.wait.aligned()194declare void @llvm.nvvm.fence.sc.cluster()195declare void @llvm.nvvm.cp.async.bulk.commit.group()196declare void @llvm.nvvm.cp.async.bulk.wait.group(i32)197declare void @llvm.nvvm.cp.async.bulk.wait.group.read(i32)198