PTX: Reverting implementation of i8.

The .b8 operations in PTX are far more limiting than I first thought. The mov operation isn't even supported, so there's no way of converting a .pred value into a .b8 without going via .b16, which is
not sensible. An improved implementation needs to use the fact that loads and stores automatically extend and truncate to implement support for EXTLOAD and TRUNCSTORE in order to correctly support
boolean values.

llvm-svn: 133873

4 files changed, 0 insertions, 250 deletions

diff --git a/llvm/test/CodeGen/PTX/cvt.ll b/llvm/test/CodeGen/PTX/cvt.ll
index dbabbf8938e..18f7ef365b4 100644
--- a/llvm/test/CodeGen/PTX/cvt.ll
+++ b/llvm/test/CodeGen/PTX/cvt.ll
@@ -3,17 +3,6 @@
 ; preds 
 ; (note: we convert back to i32 to return)
 
-define ptx_device i32 @cvt_pred_i8(i8 %x, i1 %y) {
-; CHECK: setp.gt.b8 p[[P0:[0-9]+]], rq{{[0-9]+}}, 0
-; CHECK-NEXT: and.pred p0, p[[P0:[0-9]+]], p{{[0-9]+}};
-; CHECK-NEXT: selp.u32 r{{[0-9]+}}, 1, 0, p[[P0:[0-9]+]];
-; CHECK-NEXT: ret;
-	%a = trunc i8 %x to i1
-	%b = and i1 %a, %y
-	%c = zext i1 %b to i32
-	ret i32 %c
-}
-
 define ptx_device i32 @cvt_pred_i16(i16 %x, i1 %y) {
 ; CHECK: setp.gt.b16 p[[P0:[0-9]+]], rh{{[0-9]+}}, 0
 ; CHECK-NEXT: and.pred p0, p[[P0:[0-9]+]], p{{[0-9]+}};
@@ -69,43 +58,6 @@ define ptx_device i32 @cvt_pred_f64(double %x, i1 %y) {
 	ret i32 %c
 }
 
-; i8
-
-define ptx_device i8 @cvt_i8_preds(i1 %x) {
-; CHECK: selp.u8 rq{{[0-9]+}}, 1, 0, p{{[0-9]+}};
-; CHECK-NEXT: ret;
-	%a = zext i1 %x to i8
-	ret i8 %a
-}
-
-define ptx_device i8 @cvt_i8_i32(i32 %x) {
-; CHECK: cvt.u8.u32 rq{{[0-9]+}}, r{{[0-9]+}};
-; CHECK-NEXT: ret;
-	%a = trunc i32 %x to i8
-	ret i8 %a
-}
-
-define ptx_device i8 @cvt_i8_i64(i64 %x) {
-; CHECK: cvt.u8.u64 rq{{[0-9]+}}, rd{{[0-9]+}};
-; CHECK-NEXT: ret;
-	%a = trunc i64 %x to i8
-	ret i8 %a
-}
-
-define ptx_device i8 @cvt_i8_f32(float %x) {
-; CHECK: cvt.rzi.u8.f32 rq{{[0-9]+}}, r{{[0-9]+}};
-; CHECK-NEXT: ret;
-	%a = fptoui float %x to i8
-	ret i8 %a
-}
-
-define ptx_device i8 @cvt_i8_f64(double %x) {
-; CHECK: cvt.rzi.u8.f64 rq{{[0-9]+}}, rd{{[0-9]+}};
-; CHECK-NEXT: ret;
-	%a = fptoui double %x to i8
-	ret i8 %a
-}
-
 ; i16
 
 define ptx_device i16 @cvt_i16_preds(i1 %x) {
@@ -115,13 +67,6 @@ define ptx_device i16 @cvt_i16_preds(i1 %x) {
 	ret i16 %a
 }
 
-define ptx_device i16 @cvt_i16_i8(i8 %x) {
-; CHECK: cvt.u16.u8 rh{{[0-9]+}}, rq{{[0-9]+}};
-; CHECK-NEXT: ret;
-	%a = zext i8 %x to i16
-	ret i16 %a
-}
-
 define ptx_device i16 @cvt_i16_i32(i32 %x) {
 ; CHECK: cvt.u16.u32 rh{{[0-9]+}}, r{{[0-9]+}};
 ; CHECK-NEXT: ret;
@@ -159,13 +104,6 @@ define ptx_device i32 @cvt_i32_preds(i1 %x) {
 	ret i32 %a
 }
 
-define ptx_device i32 @cvt_i32_i8(i8 %x) {
-; CHECK: cvt.u32.u8 r{{[0-9]+}}, rq{{[0-9]+}};
-; CHECK-NEXT: ret;
-	%a = zext i8 %x to i32
-	ret i32 %a
-}
-
 define ptx_device i32 @cvt_i32_i16(i16 %x) {
 ; CHECK: cvt.u32.u16 r{{[0-9]+}}, rh{{[0-9]+}};
 ; CHECK-NEXT: ret;
@@ -203,13 +141,6 @@ define ptx_device i64 @cvt_i64_preds(i1 %x) {
 	ret i64 %a
 }
 
-define ptx_device i64 @cvt_i64_i8(i8 %x) {
-; CHECK: cvt.u64.u8 rd{{[0-9]+}}, rq{{[0-9]+}};
-; CHECK-NEXT: ret;
-	%a = zext i8 %x to i64
-	ret i64 %a
-}
-
 define ptx_device i64 @cvt_i64_i16(i16 %x) {
 ; CHECK: cvt.u64.u16 rd{{[0-9]+}}, rh{{[0-9]+}};
 ; CHECK-NEXT: ret;
@@ -247,13 +178,6 @@ define ptx_device float @cvt_f32_preds(i1 %x) {
 	ret float %a
 }
 
-define ptx_device float @cvt_f32_i8(i8 %x) {
-; CHECK: cvt.rn.f32.u8 r{{[0-9]+}}, rq{{[0-9]+}};
-; CHECK-NEXT: ret;
-	%a = uitofp i8 %x to float
-	ret float %a
-}
-
 define ptx_device float @cvt_f32_i16(i16 %x) {
 ; CHECK: cvt.rn.f32.u16 r{{[0-9]+}}, rh{{[0-9]+}};
 ; CHECK-NEXT: ret;
@@ -291,13 +215,6 @@ define ptx_device double @cvt_f64_preds(i1 %x) {
 	ret double %a
 }
 
-define ptx_device double @cvt_f64_i8(i8 %x) {
-; CHECK: cvt.rn.f64.u8 rd{{[0-9]+}}, rq{{[0-9]+}};
-; CHECK-NEXT: ret;
-	%a = uitofp i8 %x to double
-	ret double %a
-}
-
 define ptx_device double @cvt_f64_i16(i16 %x) {
 ; CHECK: cvt.rn.f64.u16 rd{{[0-9]+}}, rh{{[0-9]+}};
 ; CHECK-NEXT: ret;
diff --git a/llvm/test/CodeGen/PTX/ld.ll b/llvm/test/CodeGen/PTX/ld.ll
index 951b14b8644..d184d1243ab 100644
--- a/llvm/test/CodeGen/PTX/ld.ll
+++ b/llvm/test/CodeGen/PTX/ld.ll
@@ -1,17 +1,5 @@
 ; RUN: llc < %s -march=ptx32 | FileCheck %s
 
-;CHECK: .extern .global .b8 array_i8[10];
-@array_i8 = external global [10 x i8]
-
-;CHECK: .extern .const .b8 array_constant_i8[10];
-@array_constant_i8 = external addrspace(1) constant [10 x i8]
-
-;CHECK: .extern .local .b8 array_local_i8[10];
-@array_local_i8 = external addrspace(2) global [10 x i8]
-
-;CHECK: .extern .shared .b8 array_shared_i8[10];
-@array_shared_i8 = external addrspace(4) global [10 x i8]
-
 ;CHECK: .extern .global .b8 array_i16[20];
 @array_i16 = external global [10 x i16]
 
@@ -72,13 +60,6 @@
 ;CHECK: .extern .shared .b8 array_shared_double[80];
 @array_shared_double = external addrspace(4) global [10 x double]
 
-define ptx_device i8 @t1_u8(i8* %p) {
-entry:
-;CHECK: ld.global.u8 rq{{[0-9]+}}, [r{{[0-9]+}}];
-;CHECK-NEXT: ret;
-  %x = load i8* %p
-  ret i8 %x
-}
 
 define ptx_device i16 @t1_u16(i16* %p) {
 entry:
@@ -120,15 +101,6 @@ entry:
   ret double %x
 }
 
-define ptx_device i8 @t2_u8(i8* %p) {
-entry:
-;CHECK: ld.global.u8 rq{{[0-9]+}}, [r{{[0-9]+}}+1];
-;CHECK-NEXT: ret;
-  %i = getelementptr i8* %p, i32 1
-  %x = load i8* %i
-  ret i8 %x
-}
-
 define ptx_device i16 @t2_u16(i16* %p) {
 entry:
 ;CHECK: ld.global.u16 rh{{[0-9]+}}, [r{{[0-9]+}}+2];
@@ -174,15 +146,6 @@ entry:
   ret double %x
 }
 
-define ptx_device i8 @t3_u8(i8* %p, i32 %q) {
-entry:
-;CHECK: add.u32 r[[R0:[0-9]+]], r{{[0-9]+}}, r{{[0-9]+}};
-;CHECK-NEXT: ld.global.u8 rq{{[0-9]+}}, [r[[R0]]];
-  %i = getelementptr i8* %p, i32 %q
-  %x = load i8* %i
-  ret i8 %x
-}
-
 define ptx_device i16 @t3_u16(i16* %p, i32 %q) {
 entry:
 ;CHECK: shl.b32 r[[R0:[0-9]+]], r{{[0-9]+}}, 1;
@@ -233,16 +196,6 @@ entry:
   ret double %x
 }
 
-define ptx_device i8 @t4_global_u8() {
-entry:
-;CHECK: mov.u32 r[[R0:[0-9]+]], array_i8;
-;CHECK-NEXT: ld.global.u8 rq{{[0-9]+}}, [r[[R0]]];
-;CHECK-NEXT: ret;
-  %i = getelementptr [10 x i8]* @array_i8, i32 0, i32 0
-  %x = load i8* %i
-  ret i8 %x
-}
-
 define ptx_device i16 @t4_global_u16() {
 entry:
 ;CHECK: mov.u32 r[[R0:[0-9]+]], array_i16;
@@ -343,16 +296,6 @@ entry:
   ret double %x
 }
 
-define ptx_device i8 @t4_local_u8() {
-entry:
-;CHECK: mov.u32 r[[R0:[0-9]+]], array_local_i8;
-;CHECK-NEXT: ld.local.u8 rq{{[0-9]+}}, [r[[R0]]];
-;CHECK-NEXT: ret;
-  %i = getelementptr [10 x i8] addrspace(2)* @array_local_i8, i32 0, i32 0
-  %x = load i8 addrspace(2)* %i
-  ret i8 %x
-}
-
 define ptx_device i16 @t4_local_u16() {
 entry:
 ;CHECK: mov.u32 r[[R0:[0-9]+]], array_local_i16;
@@ -403,16 +346,6 @@ entry:
   ret double %x
 }
 
-define ptx_device i8 @t4_shared_u8() {
-entry:
-;CHECK: mov.u32 r[[R0:[0-9]+]], array_shared_i8;
-;CHECK-NEXT: ld.shared.u8 rq{{[0-9]+}}, [r[[R0]]];
-;CHECK-NEXT: ret;
-  %i = getelementptr [10 x i8] addrspace(4)* @array_shared_i8, i32 0, i32 0
-  %x = load i8 addrspace(4)* %i
-  ret i8 %x
-}
-
 define ptx_device i16 @t4_shared_u16() {
 entry:
 ;CHECK: mov.u32 r[[R0:[0-9]+]], array_shared_i16;
@@ -463,16 +396,6 @@ entry:
   ret double %x
 }
 
-define ptx_device i8 @t5_u8() {
-entry:
-;CHECK: mov.u32 r[[R0:[0-9]+]], array_i8;
-;CHECK-NEXT: ld.global.u8 rq{{[0-9]+}}, [r[[R0]]+1];
-;CHECK-NEXT: ret;
-  %i = getelementptr [10 x i8]* @array_i8, i32 0, i32 1
-  %x = load i8* %i
-  ret i8 %x
-}
-
 define ptx_device i16 @t5_u16() {
 entry:
 ;CHECK: mov.u32 r[[R0:[0-9]+]], array_i16;
diff --git a/llvm/test/CodeGen/PTX/mov.ll b/llvm/test/CodeGen/PTX/mov.ll
index b930b4caefb..cce6a5b8976 100644
--- a/llvm/test/CodeGen/PTX/mov.ll
+++ b/llvm/test/CodeGen/PTX/mov.ll
@@ -1,11 +1,5 @@
 ; RUN: llc < %s -march=ptx32 | FileCheck %s
 
-define ptx_device i8 @t1_u8() {
-; CHECK: mov.u8 rq{{[0-9]+}}, 0;
-; CHECK: ret;
-	ret i8 0
-}
-
 define ptx_device i16 @t1_u16() {
 ; CHECK: mov.u16 rh{{[0-9]+}}, 0;
 ; CHECK: ret;
@@ -36,12 +30,6 @@ define ptx_device double @t1_f64() {
 	ret double 0.0
 }
 
-define ptx_device i8 @t2_u8(i8 %x) {
-; CHECK: mov.u8 rq{{[0-9]+}}, rq{{[0-9]+}};
-; CHECK: ret;
-	ret i8 %x
-}
-
 define ptx_device i16 @t2_u16(i16 %x) {
 ; CHECK: mov.u16 rh{{[0-9]+}}, rh{{[0-9]+}};
 ; CHECK: ret;
diff --git a/llvm/test/CodeGen/PTX/st.ll b/llvm/test/CodeGen/PTX/st.ll
index 596d189e4b7..b08528e1c3c 100644
--- a/llvm/test/CodeGen/PTX/st.ll
+++ b/llvm/test/CodeGen/PTX/st.ll
@@ -1,17 +1,5 @@
 ; RUN: llc < %s -march=ptx32 | FileCheck %s
 
-;CHECK: .extern .global .b8 array_i8[10];
-@array_i8 = external global [10 x i8]
-
-;CHECK: .extern .const .b8 array_constant_i8[10];
-@array_constant_i8 = external addrspace(1) constant [10 x i8]
-
-;CHECK: .extern .local .b8 array_local_i8[10];
-@array_local_i8 = external addrspace(2) global [10 x i8]
-
-;CHECK: .extern .shared .b8 array_shared_i8[10];
-@array_shared_i8 = external addrspace(4) global [10 x i8]
-
 ;CHECK: .extern .global .b8 array_i16[20];
 @array_i16 = external global [10 x i16]
 
@@ -72,13 +60,6 @@
 ;CHECK: .extern .shared .b8 array_shared_double[80];
 @array_shared_double = external addrspace(4) global [10 x double]
 
-define ptx_device void @t1_u8(i8* %p, i8 %x) {
-entry:
-;CHECK: st.global.u8 [r{{[0-9]+}}], rq{{[0-9]+}};
-;CHECK-NEXT: ret;
-  store i8 %x, i8* %p
-  ret void
-}
 
 define ptx_device void @t1_u16(i16* %p, i16 %x) {
 entry:
@@ -120,15 +101,6 @@ entry:
   ret void
 }
 
-define ptx_device void @t2_u8(i8* %p, i8 %x) {
-entry:
-;CHECK: st.global.u8 [r{{[0-9]+}}+1], rq{{[0-9]+}};
-;CHECK-NEXT: ret;
-  %i = getelementptr i8* %p, i32 1
-  store i8 %x, i8* %i
-  ret void
-}
-
 define ptx_device void @t2_u16(i16* %p, i16 %x) {
 entry:
 ;CHECK: st.global.u16 [r{{[0-9]+}}+2], rh{{[0-9]+}};
@@ -174,16 +146,6 @@ entry:
   ret void
 }
 
-define ptx_device void @t3_u8(i8* %p, i32 %q, i8 %x) {
-entry:
-;CHECK: add.u32 r[[R0:[0-9]+]], r{{[0-9]+}}, r{{[0-9]+}};
-;CHECK-NEXT: st.global.u8 [r[[R0]]], rq{{[0-9]+}};
-;CHECK-NEXT: ret;
-  %i = getelementptr i8* %p, i32 %q
-  store i8 %x, i8* %i
-  ret void
-}
-
 define ptx_device void @t3_u16(i16* %p, i32 %q, i16 %x) {
 entry:
 ;CHECK: shl.b32 r[[R0:[0-9]+]], r{{[0-9]+}}, 1;
@@ -239,16 +201,6 @@ entry:
   ret void
 }
 
-define ptx_device void @t4_global_u8(i8 %x) {
-entry:
-;CHECK: mov.u32 r[[R0:[0-9]+]], array_i8;
-;CHECK-NEXT: st.global.u8 [r[[R0]]], rq{{[0-9]+}};
-;CHECK-NEXT: ret;
-  %i = getelementptr [10 x i8]* @array_i8, i8 0, i8 0
-  store i8 %x, i8* %i
-  ret void
-}
-
 define ptx_device void @t4_global_u16(i16 %x) {
 entry:
 ;CHECK: mov.u32 r[[R0:[0-9]+]], array_i16;
@@ -299,16 +251,6 @@ entry:
   ret void
 }
 
-define ptx_device void @t4_local_u8(i8 %x) {
-entry:
-;CHECK: mov.u32 r[[R0:[0-9]+]], array_local_i8;
-;CHECK-NEXT: st.local.u8 [r[[R0]]], rq{{[0-9]+}};
-;CHECK-NEXT: ret;
-  %i = getelementptr [10 x i8] addrspace(2)* @array_local_i8, i32 0, i32 0
-  store i8 %x, i8 addrspace(2)* %i
-  ret void
-}
-
 define ptx_device void @t4_local_u16(i16 %x) {
 entry:
 ;CHECK: mov.u32 r[[R0:[0-9]+]], array_local_i16;
@@ -359,16 +301,6 @@ entry:
   ret void
 }
 
-define ptx_device void @t4_shared_u8(i8 %x) {
-entry:
-;CHECK: mov.u32 r[[R0:[0-9]+]], array_shared_i8;
-;CHECK-NEXT: st.shared.u8 [r[[R0]]], rq{{[0-9]+}};
-;CHECK-NEXT: ret;
-  %i = getelementptr [10 x i8] addrspace(4)* @array_shared_i8, i32 0, i32 0
-  store i8 %x, i8 addrspace(4)* %i
-  ret void
-}
-
 define ptx_device void @t4_shared_u16(i16 %x) {
 entry:
 ;CHECK: mov.u32 r[[R0:[0-9]+]], array_shared_i16;
@@ -419,16 +351,6 @@ entry:
   ret void
 }
 
-define ptx_device void @t5_u8(i8 %x) {
-entry:
-;CHECK: mov.u32 r[[R0:[0-9]+]], array_i8;
-;CHECK-NEXT: st.global.u8 [r[[R0]]+1], rq{{[0-9]+}};
-;CHECK-NEXT: ret;
-  %i = getelementptr [10 x i8]* @array_i8, i32 0, i32 1
-  store i8 %x, i8* %i
-  ret void
-}
-
 define ptx_device void @t5_u16(i16 %x) {
 entry:
 ;CHECK: mov.u32 r[[R0:[0-9]+]], array_i16;