PTX: Add intrinsics to list of built-in intrinsics, which allows them to be

     used by Clang.  To help Clang integration, the PTX target has been split
     into two targets: ptx32 and ptx64, depending on the desired pointer size.

- Add GCCBuiltin class to all intrinsics
- Split PTX target into ptx32 and ptx64

llvm-svn: 129851

19 files changed, 24 insertions, 24 deletions

diff --git a/llvm/test/CodeGen/PTX/add.ll b/llvm/test/CodeGen/PTX/add.ll
index 598591c0fcb..235b00e8782 100644
--- a/llvm/test/CodeGen/PTX/add.ll
+++ b/llvm/test/CodeGen/PTX/add.ll
@@ -1,4 +1,4 @@
-; RUN: llc < %s -march=ptx | FileCheck %s
+; RUN: llc < %s -march=ptx32 | FileCheck %s
 
 define ptx_device i16 @t1_u16(i16 %x, i16 %y) {
 ; CHECK: add.u16 rh0, rh1, rh2;
diff --git a/llvm/test/CodeGen/PTX/bra.ll b/llvm/test/CodeGen/PTX/bra.ll
index 0506a990668..49383eb3cf9 100644
--- a/llvm/test/CodeGen/PTX/bra.ll
+++ b/llvm/test/CodeGen/PTX/bra.ll
@@ -1,4 +1,4 @@
-; RUN: llc < %s -march=ptx | FileCheck %s
+; RUN: llc < %s -march=ptx32 | FileCheck %s
 
 define ptx_device void @test_bra_direct() {
 ; CHECK: bra $L__BB0_1;
diff --git a/llvm/test/CodeGen/PTX/exit.ll b/llvm/test/CodeGen/PTX/exit.ll
index 4071babb80c..7816c801728 100644
--- a/llvm/test/CodeGen/PTX/exit.ll
+++ b/llvm/test/CodeGen/PTX/exit.ll
@@ -1,4 +1,4 @@
-; RUN: llc < %s -march=ptx | FileCheck %s
+; RUN: llc < %s -march=ptx32 | FileCheck %s
 
 define ptx_kernel void @t1() {
 ; CHECK: exit;
diff --git a/llvm/test/CodeGen/PTX/fdiv-sm10.ll b/llvm/test/CodeGen/PTX/fdiv-sm10.ll
index 42f615d0c8d..121360ce9be 100644
--- a/llvm/test/CodeGen/PTX/fdiv-sm10.ll
+++ b/llvm/test/CodeGen/PTX/fdiv-sm10.ll
@@ -1,4 +1,4 @@
-; RUN: llc < %s -march=ptx -mattr=+sm10 | FileCheck %s
+; RUN: llc < %s -march=ptx32 -mattr=+sm10 | FileCheck %s
 
 define ptx_device float @t1_f32(float %x, float %y) {
 ; CHECK: div.approx.f32 f0, f1, f2;
diff --git a/llvm/test/CodeGen/PTX/fdiv-sm13.ll b/llvm/test/CodeGen/PTX/fdiv-sm13.ll
index eb20f787639..0ec7bae8030 100644
--- a/llvm/test/CodeGen/PTX/fdiv-sm13.ll
+++ b/llvm/test/CodeGen/PTX/fdiv-sm13.ll
@@ -1,4 +1,4 @@
-; RUN: llc < %s -march=ptx -mattr=+sm13 | FileCheck %s
+; RUN: llc < %s -march=ptx32 -mattr=+sm13 | FileCheck %s
 
 define ptx_device float @t1_f32(float %x, float %y) {
 ; CHECK: div.approx.f32 f0, f1, f2;
diff --git a/llvm/test/CodeGen/PTX/intrinsic.ll b/llvm/test/CodeGen/PTX/intrinsic.ll
index 7405dd6f5e5..cea41827ca4 100644
--- a/llvm/test/CodeGen/PTX/intrinsic.ll
+++ b/llvm/test/CodeGen/PTX/intrinsic.ll
@@ -1,4 +1,4 @@
-; RUN: llc < %s -march=ptx -mattr=+ptx20,+sm20 | FileCheck %s
+; RUN: llc < %s -march=ptx32 -mattr=+ptx20,+sm20 | FileCheck %s
 
 define ptx_device i32 @test_tid_x() {
 ; CHECK: mov.u32 r0, %tid.x;
diff --git a/llvm/test/CodeGen/PTX/ld.ll b/llvm/test/CodeGen/PTX/ld.ll
index 1119aa46944..58e16a20a45 100644
--- a/llvm/test/CodeGen/PTX/ld.ll
+++ b/llvm/test/CodeGen/PTX/ld.ll
@@ -1,4 +1,4 @@
-; RUN: llc < %s -march=ptx | FileCheck %s
+; RUN: llc < %s -march=ptx32 | FileCheck %s
 
 ;CHECK: .extern .global .b8 array_i16[20];
 @array_i16 = external global [10 x i16]
diff --git a/llvm/test/CodeGen/PTX/llvm-intrinsic.ll b/llvm/test/CodeGen/PTX/llvm-intrinsic.ll
index 3ce4c29f9f5..1e265f5b7b3 100644
--- a/llvm/test/CodeGen/PTX/llvm-intrinsic.ll
+++ b/llvm/test/CodeGen/PTX/llvm-intrinsic.ll
@@ -1,4 +1,4 @@
-; RUN: llc < %s -march=ptx -mattr=+ptx20,+sm20 | FileCheck %s
+; RUN: llc < %s -march=ptx32 -mattr=+ptx20,+sm20 | FileCheck %s
 
 define ptx_device float @test_sqrt_f32(float %x) {
 entry:
diff --git a/llvm/test/CodeGen/PTX/mad.ll b/llvm/test/CodeGen/PTX/mad.ll
index 786345b2913..0c25f2c0030 100644
--- a/llvm/test/CodeGen/PTX/mad.ll
+++ b/llvm/test/CodeGen/PTX/mad.ll
@@ -1,4 +1,4 @@
-; RUN: llc < %s -march=ptx -mattr=+sm13 | FileCheck %s
+; RUN: llc < %s -march=ptx32 -mattr=+sm13 | FileCheck %s
 
 define ptx_device float @t1_f32(float %x, float %y, float %z) {
 ; CHECK: mad.rn.f32 f0, f1, f2, f3;
diff --git a/llvm/test/CodeGen/PTX/mov.ll b/llvm/test/CodeGen/PTX/mov.ll
index 00dcf19f1da..120572a0e86 100644
--- a/llvm/test/CodeGen/PTX/mov.ll
+++ b/llvm/test/CodeGen/PTX/mov.ll
@@ -1,4 +1,4 @@
-; RUN: llc < %s -march=ptx | FileCheck %s
+; RUN: llc < %s -march=ptx32 | FileCheck %s
 
 define ptx_device i16 @t1_u16() {
 ; CHECK: mov.u16 rh0, 0;
diff --git a/llvm/test/CodeGen/PTX/mul.ll b/llvm/test/CodeGen/PTX/mul.ll
index fd0788fce66..5ce042675dc 100644
--- a/llvm/test/CodeGen/PTX/mul.ll
+++ b/llvm/test/CodeGen/PTX/mul.ll
@@ -1,4 +1,4 @@
-; RUN: llc < %s -march=ptx | FileCheck %s
+; RUN: llc < %s -march=ptx32 | FileCheck %s
 
 ;define ptx_device i32 @t1(i32 %x, i32 %y) {
 ;	%z = mul i32 %x, %y
diff --git a/llvm/test/CodeGen/PTX/options.ll b/llvm/test/CodeGen/PTX/options.ll
index 6576a6d8bbb..ac33fef0d6e 100644
--- a/llvm/test/CodeGen/PTX/options.ll
+++ b/llvm/test/CodeGen/PTX/options.ll
@@ -1,9 +1,9 @@
-; RUN: llc < %s -march=ptx -mattr=ptx20 | grep ".version 2.0"
-; RUN: llc < %s -march=ptx -mattr=ptx21 | grep ".version 2.1"
-; RUN: llc < %s -march=ptx -mattr=ptx22 | grep ".version 2.2"
-; RUN: llc < %s -march=ptx -mattr=sm10 | grep ".target sm_10"
-; RUN: llc < %s -march=ptx -mattr=sm13 | grep ".target sm_13"
-; RUN: llc < %s -march=ptx -mattr=sm20 | grep ".target sm_20"
+; RUN: llc < %s -march=ptx32 -mattr=ptx20 | grep ".version 2.0"
+; RUN: llc < %s -march=ptx32 -mattr=ptx21 | grep ".version 2.1"
+; RUN: llc < %s -march=ptx32 -mattr=ptx22 | grep ".version 2.2"
+; RUN: llc < %s -march=ptx32 -mattr=sm10 | grep ".target sm_10"
+; RUN: llc < %s -march=ptx32 -mattr=sm13 | grep ".target sm_13"
+; RUN: llc < %s -march=ptx32 -mattr=sm20 | grep ".target sm_20"
 
 define ptx_device void @t1() {
 	ret void
diff --git a/llvm/test/CodeGen/PTX/parameter-order.ll b/llvm/test/CodeGen/PTX/parameter-order.ll
index dbbbb67a140..8131f13a6e8 100644
--- a/llvm/test/CodeGen/PTX/parameter-order.ll
+++ b/llvm/test/CodeGen/PTX/parameter-order.ll
@@ -1,4 +1,4 @@
-; RUN: llc < %s -march=ptx | FileCheck %s
+; RUN: llc < %s -march=ptx32 | FileCheck %s
 
 ; CHECK: .func (.reg .u32 r0) test_parameter_order (.reg .u32 r1, .reg .u32 r2)
 define ptx_device i32 @test_parameter_order(i32 %x, i32 %y) {
diff --git a/llvm/test/CodeGen/PTX/ret.ll b/llvm/test/CodeGen/PTX/ret.ll
index d5037f25fd3..ba0523f6424 100644
--- a/llvm/test/CodeGen/PTX/ret.ll
+++ b/llvm/test/CodeGen/PTX/ret.ll
@@ -1,4 +1,4 @@
-; RUN: llc < %s -march=ptx | FileCheck %s
+; RUN: llc < %s -march=ptx32 | FileCheck %s
 
 define ptx_device void @t1() {
 ; CHECK: ret;
diff --git a/llvm/test/CodeGen/PTX/setp.ll b/llvm/test/CodeGen/PTX/setp.ll
index 5348482e093..5836122049e 100644
--- a/llvm/test/CodeGen/PTX/setp.ll
+++ b/llvm/test/CodeGen/PTX/setp.ll
@@ -1,4 +1,4 @@
-; RUN: llc < %s -march=ptx | FileCheck %s
+; RUN: llc < %s -march=ptx32 | FileCheck %s
 
 define ptx_device i32 @test_setp_eq_u32_rr(i32 %x, i32 %y) {
 ; CHECK: setp.eq.u32 p0, r1, r2;
diff --git a/llvm/test/CodeGen/PTX/shl.ll b/llvm/test/CodeGen/PTX/shl.ll
index b564b43ab93..6e72c922132 100644
--- a/llvm/test/CodeGen/PTX/shl.ll
+++ b/llvm/test/CodeGen/PTX/shl.ll
@@ -1,4 +1,4 @@
-; RUN: llc < %s -march=ptx | FileCheck %s
+; RUN: llc < %s -march=ptx32 | FileCheck %s
 
 define ptx_device i32 @t1(i32 %x, i32 %y) {
 ; CHECK: shl.b32 r0, r1, r2
diff --git a/llvm/test/CodeGen/PTX/shr.ll b/llvm/test/CodeGen/PTX/shr.ll
index 3f8ade862b7..8693e0ecf49 100644
--- a/llvm/test/CodeGen/PTX/shr.ll
+++ b/llvm/test/CodeGen/PTX/shr.ll
@@ -1,4 +1,4 @@
-; RUN: llc < %s -march=ptx | FileCheck %s
+; RUN: llc < %s -march=ptx32 | FileCheck %s
 
 define ptx_device i32 @t1(i32 %x, i32 %y) {
 ; CHECK: shr.u32 r0, r1, r2
diff --git a/llvm/test/CodeGen/PTX/st.ll b/llvm/test/CodeGen/PTX/st.ll
index 4e9b08a33a2..dee5c61abe6 100644
--- a/llvm/test/CodeGen/PTX/st.ll
+++ b/llvm/test/CodeGen/PTX/st.ll
@@ -1,4 +1,4 @@
-; RUN: llc < %s -march=ptx | FileCheck %s
+; RUN: llc < %s -march=ptx32 | FileCheck %s
 
 ;CHECK: .extern .global .b8 array_i16[20];
 @array_i16 = external global [10 x i16]
diff --git a/llvm/test/CodeGen/PTX/sub.ll b/llvm/test/CodeGen/PTX/sub.ll
index 4810e4fc055..7dd2c6f6ac7 100644
--- a/llvm/test/CodeGen/PTX/sub.ll
+++ b/llvm/test/CodeGen/PTX/sub.ll
@@ -1,4 +1,4 @@
-; RUN: llc < %s -march=ptx | FileCheck %s
+; RUN: llc < %s -march=ptx32 | FileCheck %s
 
 define ptx_device i16 @t1_u16(i16 %x, i16 %y) {
 ; CHECK: sub.u16 rh0, rh1, rh2;