1// RUN: mlir-translate -mlir-to-nvvmir %s | FileCheck %s
2
3llvm.func @nvvm_special_regs() -> !llvm.i32 {
4  // CHECK: %1 = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
5  %1 = nvvm.read.ptx.sreg.tid.x : !llvm.i32
6  // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.y()
7  %2 = nvvm.read.ptx.sreg.tid.y : !llvm.i32
8  // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.z()
9  %3 = nvvm.read.ptx.sreg.tid.z : !llvm.i32
10  // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
11  %4 = nvvm.read.ptx.sreg.ntid.x : !llvm.i32
12  // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ntid.y()
13  %5 = nvvm.read.ptx.sreg.ntid.y : !llvm.i32
14  // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ntid.z()
15  %6 = nvvm.read.ptx.sreg.ntid.z : !llvm.i32
16  // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
17  %7 = nvvm.read.ptx.sreg.ctaid.x : !llvm.i32
18  // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ctaid.y()
19  %8 = nvvm.read.ptx.sreg.ctaid.y : !llvm.i32
20  // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ctaid.z()
21  %9 = nvvm.read.ptx.sreg.ctaid.z : !llvm.i32
22  // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nctaid.x()
23  %10 = nvvm.read.ptx.sreg.nctaid.x : !llvm.i32
24  // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nctaid.y()
25  %11 = nvvm.read.ptx.sreg.nctaid.y : !llvm.i32
26  // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nctaid.z()
27  %12 = nvvm.read.ptx.sreg.nctaid.z : !llvm.i32
28  // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
29  %13 = nvvm.read.ptx.sreg.warpsize : !llvm.i32
30  // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.laneid()
31  %14 = nvvm.read.ptx.sreg.laneid : !llvm.i32
32  llvm.return %1 : !llvm.i32
33}
34
35llvm.func @llvm.nvvm.barrier0() {
36  // CHECK: call void @llvm.nvvm.barrier0()
37  nvvm.barrier0
38  llvm.return
39}
40
41llvm.func @nvvm_shfl(
42    %0 : !llvm.i32, %1 : !llvm.i32, %2 : !llvm.i32,
43    %3 : !llvm.i32, %4 : !llvm.float) -> !llvm.i32 {
44  // CHECK: call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}})
45  %6 = nvvm.shfl.sync.bfly %0, %3, %1, %2 : !llvm.i32
46  // CHECK: call float @llvm.nvvm.shfl.sync.bfly.f32(i32 %{{.*}}, float %{{.*}}, i32 %{{.*}}, i32 %{{.*}})
47  %7 = nvvm.shfl.sync.bfly %0, %4, %1, %2 : !llvm.float
48  llvm.return %6 : !llvm.i32
49}
50
51llvm.func @nvvm_shfl_pred(
52    %0 : !llvm.i32, %1 : !llvm.i32, %2 : !llvm.i32,
53    %3 : !llvm.i32, %4 : !llvm.float) -> !llvm.struct<(i32, i1)> {
54  // CHECK: call { i32, i1 } @llvm.nvvm.shfl.sync.bfly.i32p(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}})
55  %6 = nvvm.shfl.sync.bfly %0, %3, %1, %2 {return_value_and_is_valid} : !llvm.struct<(i32, i1)>
56  // CHECK: call { float, i1 } @llvm.nvvm.shfl.sync.bfly.f32p(i32 %{{.*}}, float %{{.*}}, i32 %{{.*}}, i32 %{{.*}})
57  %7 = nvvm.shfl.sync.bfly %0, %4, %1, %2 {return_value_and_is_valid} : !llvm.struct<(float, i1)>
58  llvm.return %6 : !llvm.struct<(i32, i1)>
59}
60
61llvm.func @nvvm_vote(%0 : !llvm.i32, %1 : !llvm.i1) -> !llvm.i32 {
62  // CHECK: call i32 @llvm.nvvm.vote.ballot.sync(i32 %{{.*}}, i1 %{{.*}})
63  %3 = nvvm.vote.ballot.sync %0, %1 : !llvm.i32
64  llvm.return %3 : !llvm.i32
65}
66
67llvm.func @nvvm_mma(%a0 : !llvm.vec<2 x half>, %a1 : !llvm.vec<2 x half>,
68                    %b0 : !llvm.vec<2 x half>, %b1 : !llvm.vec<2 x half>,
69                    %c0 : !llvm.float, %c1 : !llvm.float, %c2 : !llvm.float, %c3 : !llvm.float,
70                    %c4 : !llvm.float, %c5 : !llvm.float, %c6 : !llvm.float, %c7 : !llvm.float) {
71  // CHECK: call { float, float, float, float, float, float, float, float } @llvm.nvvm.mma.m8n8k4.row.col.f32.f32
72  %0 = nvvm.mma.sync %a0, %a1, %b0, %b1, %c0, %c1, %c2, %c3, %c4, %c5, %c6, %c7 {alayout="row", blayout="col"} : (!llvm.vec<2 x half>, !llvm.vec<2 x half>, !llvm.vec<2 x half>, !llvm.vec<2 x half>, !llvm.float, !llvm.float, !llvm.float, !llvm.float, !llvm.float, !llvm.float, !llvm.float, !llvm.float) -> !llvm.struct<(float, float, float, float, float, float, float, float)>
73  llvm.return %0 : !llvm.struct<(float, float, float, float, float, float, float, float)>
74}
75
76// This function has the "kernel" attribute attached and should appear in the
77// NVVM annotations after conversion.
78llvm.func @kernel_func() attributes {gpu.kernel} {
79  llvm.return
80}
81
82// CHECK:     !nvvm.annotations =
83// CHECK-NOT: {i32 ()* @nvvm_special_regs, !"kernel", i32 1}
84// CHECK:     {void ()* @kernel_func, !"kernel", i32 1}
85