llvm.org GIT mirror llvm / 0df2c50
ptx: add basic intrinsic support git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@127084 91177308-0d34-0410-b5e6-96231b3b80d8 Che-Liang Chiou 9 years ago
6 changed file(s) with 117 addition(s) and 2 deletion(s). Raw diff Collapse all Expand all
489489 include "llvm/IntrinsicsCellSPU.td"
490490 include "llvm/IntrinsicsAlpha.td"
491491 include "llvm/IntrinsicsXCore.td"
492 include "llvm/IntrinsicsPTX.td"
0 //===- IntrinsicsARM.td - Defines ARM intrinsics -----------*- tablegen -*-===//
1 //
1 //
22 // The LLVM Compiler Infrastructure
33 //
44 // This file is distributed under the University of Illinois Open Source
55 // License. See LICENSE.TXT for details.
6 //
6 //
77 //===----------------------------------------------------------------------===//
88 //
99 // This file defines all of the ARM-specific intrinsics.
0 //===- IntrinsicsPTX.td - Defines PTX intrinsics -----------*- tablegen -*-===//
1 //
2 // The LLVM Compiler Infrastructure
3 //
4 // This file is distributed under the University of Illinois Open Source
5 // License. See LICENSE.TXT for details.
6 //
7 //===----------------------------------------------------------------------===//
8 //
9 // This file defines all of the PTX-specific intrinsics.
10 //
11 //===----------------------------------------------------------------------===//
12
13 let TargetPrefix = "ptx" in {
14 multiclass PTXReadSpecialRegisterIntrinsic {
15 def _r64 : Intrinsic<[llvm_i64_ty], [], [IntrNoMem]>;
16 def _v4i16 : Intrinsic<[llvm_v4i16_ty], [], [IntrNoMem]>;
17 }
18
19 multiclass PTXReadSpecialSubRegisterIntrinsic {
20 def _x : Intrinsic<[llvm_i16_ty], [], [IntrNoMem]>;
21 def _y : Intrinsic<[llvm_i16_ty], [], [IntrNoMem]>;
22 def _z : Intrinsic<[llvm_i16_ty], [], [IntrNoMem]>;
23 def _w : Intrinsic<[llvm_i16_ty], [], [IntrNoMem]>;
24 }
25 }
26
27 defm int_ptx_read_tid : PTXReadSpecialRegisterIntrinsic;
28 defm int_ptx_read_tid : PTXReadSpecialSubRegisterIntrinsic;
29
30 let TargetPrefix = "ptx" in
31 def int_ptx_bar_sync : Intrinsic<[], [llvm_i32_ty], []>;
389389 def EXIT : InstPTX<(outs), (ins), "exit", [(PTXexit)]>;
390390 def RET : InstPTX<(outs), (ins), "ret", [(PTXret)]>;
391391 }
392
393 ///===- Intrinsic Instructions --------------------------------------------===//
394
395 include "PTXIntrinsicInstrInfo.td"
0 //===- PTXIntrinsicInstrInfo.td - Defines PTX intrinsics ---*- tablegen -*-===//
1 //
2 // The LLVM Compiler Infrastructure
3 //
4 // This file is distributed under the University of Illinois Open Source
5 // License. See LICENSE.TXT for details.
6 //
7 //===----------------------------------------------------------------------===//
8 //
9 // This file defines all of the PTX-specific intrinsic instructions.
10 //
11 //===----------------------------------------------------------------------===//
12
13 // PTX Special Purpose Register Accessor Intrinsics
14
15 class PTX_READ_SPECIAL_REGISTER
16 : InstPTX<(outs RRegu64:$d), (ins),
17 !strconcat("mov.u64\t$d, ", regname),
18 [(set RRegu64:$d, (intop))]>;
19
20 class PTX_READ_SPECIAL_SUB_REGISTER
21 : InstPTX<(outs RRegu16:$d), (ins),
22 !strconcat("mov.u16\t$d, ", regname),
23 [(set RRegu16:$d, (intop))]>;
24
25 def PTX_READ_TID_R64 : PTX_READ_SPECIAL_REGISTER<"tid", int_ptx_read_tid_r64>;
26 def PTX_READ_TID_X : PTX_READ_SPECIAL_SUB_REGISTER<"tid.x", int_ptx_read_tid_x>;
27 def PTX_READ_TID_Y : PTX_READ_SPECIAL_SUB_REGISTER<"tid.y", int_ptx_read_tid_y>;
28 def PTX_READ_TID_Z : PTX_READ_SPECIAL_SUB_REGISTER<"tid.z", int_ptx_read_tid_z>;
29 def PTX_READ_TID_W : PTX_READ_SPECIAL_SUB_REGISTER<"tid.w", int_ptx_read_tid_w>;
30
31 // PTX Parallel Synchronization and Communication Intrinsics
32
33 def PTX_BAR_SYNC : InstPTX<(outs), (ins i32imm:$i), "bar.sync\t$i",
34 [(int_ptx_bar_sync imm:$i)]>;
0 ; RUN: llc < %s -march=ptx | FileCheck %s
1
2 define ptx_device i16 @tid_x() {
3 ; CHECK: mov.u16 rh0, tid.x;
4 ; CHECK-NEXT: ret;
5 %x = call i16 @llvm.ptx.read.tid.x()
6 ret i16 %x
7 }
8
9 define ptx_device i16 @tid_y() {
10 ; CHECK: mov.u16 rh0, tid.y;
11 ; CHECK-NEXT: ret;
12 %x = call i16 @llvm.ptx.read.tid.y()
13 ret i16 %x
14 }
15
16 define ptx_device i16 @tid_z() {
17 ; CHECK: mov.u16 rh0, tid.z;
18 ; CHECK-NEXT: ret;
19 %x = call i16 @llvm.ptx.read.tid.z()
20 ret i16 %x
21 }
22
23 define ptx_device i16 @tid_w() {
24 ; CHECK: mov.u16 rh0, tid.w;
25 ; CHECK-NEXT: ret;
26 %x = call i16 @llvm.ptx.read.tid.w()
27 ret i16 %x
28 }
29
30 define ptx_device void @bar_sync() {
31 ; CHECK: bar.sync 0
32 ; CHECK-NEXT: ret;
33 call void @llvm.ptx.bar.sync(i32 0)
34 ret void
35 }
36
37 declare i16 @llvm.ptx.read.tid.x()
38 declare i16 @llvm.ptx.read.tid.y()
39 declare i16 @llvm.ptx.read.tid.z()
40 declare i16 @llvm.ptx.read.tid.w()
41
42 declare void @llvm.ptx.bar.sync(i32 %i)