PTX: Finish new calling convention implementation

git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@133172 91177308-0d34-0410-b5e6-96231b3b80d8

14 files changed, 82 insertions, 82 deletions

diff --git a/test/CodeGen/PTX/add.ll b/test/CodeGen/PTX/add.ll
index 235b00e..b89a2f6 100644
--- a/test/CodeGen/PTX/add.ll
+++ b/test/CodeGen/PTX/add.ll
@@ -22,14 +22,14 @@ define ptx_device i64 @t1_u64(i64 %x, i64 %y) {
 }
 
 define ptx_device float @t1_f32(float %x, float %y) {
-; CHECK: add.f32 f0, f1, f2
+; CHECK: add.f32 r0, r1, r2
 ; CHECK-NEXT: ret;
   %z = fadd float %x, %y
   ret float %z
 }
 
 define ptx_device double @t1_f64(double %x, double %y) {
-; CHECK: add.f64 fd0, fd1, fd2
+; CHECK: add.f64 rd0, rd1, rd2
 ; CHECK-NEXT: ret;
   %z = fadd double %x, %y
   ret double %z
@@ -57,14 +57,14 @@ define ptx_device i64 @t2_u64(i64 %x) {
 }
 
 define ptx_device float @t2_f32(float %x) {
-; CHECK: add.f32 f0, f1, 0F3F800000;
+; CHECK: add.f32 r0, r1, 0F3F800000;
 ; CHECK-NEXT: ret;
   %z = fadd float %x, 1.0
   ret float %z
 }
 
 define ptx_device double @t2_f64(double %x) {
-; CHECK: add.f64 fd0, fd1, 0D3FF0000000000000;
+; CHECK: add.f64 rd0, rd1, 0D3FF0000000000000;
 ; CHECK-NEXT: ret;
   %z = fadd double %x, 1.0
   ret double %z
diff --git a/test/CodeGen/PTX/cvt.ll b/test/CodeGen/PTX/cvt.ll
index 2f793de..984cb4d 100644
--- a/test/CodeGen/PTX/cvt.ll
+++ b/test/CodeGen/PTX/cvt.ll
@@ -31,7 +31,7 @@ define ptx_device i32 @cvt_pred_i64(i64 %x, i1 %y) {
 }
 
 define ptx_device i32 @cvt_pred_f32(float %x, i1 %y) {
-; CHECK: cvt.rni.pred.f32 p0, f1;
+; CHECK: cvt.rni.pred.f32 p0, r1;
 ; CHECK: ret;
 	%a = fptoui float %x to i1
 	%b = and i1 %a, %y
@@ -40,7 +40,7 @@ define ptx_device i32 @cvt_pred_f32(float %x, i1 %y) {
 }
 
 define ptx_device i32 @cvt_pred_f64(double %x, i1 %y) {
-; CHECK: cvt.rni.pred.f64 p0, fd1;
+; CHECK: cvt.rni.pred.f64 p0, rd1;
 ; CHECK: ret;
 	%a = fptoui double %x to i1
 	%b = and i1 %a, %y
@@ -72,14 +72,14 @@ define ptx_device i16 @cvt_i16_i64(i64 %x) {
 }
 
 define ptx_device i16 @cvt_i16_f32(float %x) {
-; CHECK: cvt.rni.u16.f32 rh0, f1;
+; CHECK: cvt.rni.u16.f32 rh0, r1;
 ; CHECK: ret;
 	%a = fptoui float %x to i16
 	ret i16 %a
 }
 
 define ptx_device i16 @cvt_i16_f64(double %x) {
-; CHECK: cvt.rni.u16.f64 rh0, fd1;
+; CHECK: cvt.rni.u16.f64 rh0, rd1;
 ; CHECK: ret;
 	%a = fptoui double %x to i16
 	ret i16 %a
@@ -109,14 +109,14 @@ define ptx_device i32 @cvt_i32_i64(i64 %x) {
 }
 
 define ptx_device i32 @cvt_i32_f32(float %x) {
-; CHECK: cvt.rni.u32.f32 r0, f1;
+; CHECK: cvt.rni.u32.f32 r0, r1;
 ; CHECK: ret;
 	%a = fptoui float %x to i32
 	ret i32 %a
 }
 
 define ptx_device i32 @cvt_i32_f64(double %x) {
-; CHECK: cvt.rni.u32.f64 r0, fd1;
+; CHECK: cvt.rni.u32.f64 r0, rd1;
 ; CHECK: ret;
 	%a = fptoui double %x to i32
 	ret i32 %a
@@ -146,14 +146,14 @@ define ptx_device i64 @cvt_i64_i32(i32 %x) {
 }
 
 define ptx_device i64 @cvt_i64_f32(float %x) {
-; CHECK: cvt.rni.u64.f32 rd0, f1;
+; CHECK: cvt.rni.u64.f32 rd0, r1;
 ; CHECK: ret;
 	%a = fptoui float %x to i64
 	ret i64 %a
 }
 
 define ptx_device i64 @cvt_i64_f64(double %x) {
-; CHECK: cvt.rni.u64.f64 rd0, fd1;
+; CHECK: cvt.rni.u64.f64 rd0, rd1;
 ; CHECK: ret;
 	%a = fptoui double %x to i64
 	ret i64 %a
@@ -162,35 +162,35 @@ define ptx_device i64 @cvt_i64_f64(double %x) {
 ; f32
 
 define ptx_device float @cvt_f32_preds(i1 %x) {
-; CHECK: cvt.rn.f32.pred f0, p1;
+; CHECK: cvt.rn.f32.pred r0, p1;
 ; CHECK: ret;
 	%a = uitofp i1 %x to float
 	ret float %a
 }
 
 define ptx_device float @cvt_f32_i16(i16 %x) {
-; CHECK: cvt.rn.f32.u16 f0, rh1;
+; CHECK: cvt.rn.f32.u16 r0, rh1;
 ; CHECK: ret;
 	%a = uitofp i16 %x to float
 	ret float %a
 }
 
 define ptx_device float @cvt_f32_i32(i32 %x) {
-; CHECK: cvt.rn.f32.u32 f0, r1;
+; CHECK: cvt.rn.f32.u32 r0, r1;
 ; CHECK: ret;
 	%a = uitofp i32 %x to float
 	ret float %a
 }
 
 define ptx_device float @cvt_f32_i64(i64 %x) {
-; CHECK: cvt.rn.f32.u64 f0, rd1;
+; CHECK: cvt.rn.f32.u64 r0, rd1;
 ; CHECK: ret;
 	%a = uitofp i64 %x to float
 	ret float %a
 }
 
 define ptx_device float @cvt_f32_f64(double %x) {
-; CHECK: cvt.rn.f32.f64 f0, fd1;
+; CHECK: cvt.rn.f32.f64 r0, rd1;
 ; CHECK: ret;
 	%a = fptrunc double %x to float
 	ret float %a
@@ -199,35 +199,35 @@ define ptx_device float @cvt_f32_f64(double %x) {
 ; f64
 
 define ptx_device double @cvt_f64_preds(i1 %x) {
-; CHECK: cvt.rn.f64.pred fd0, p1;
+; CHECK: cvt.rn.f64.pred rd0, p1;
 ; CHECK: ret;
 	%a = uitofp i1 %x to double
 	ret double %a
 }
 
 define ptx_device double @cvt_f64_i16(i16 %x) {
-; CHECK: cvt.rn.f64.u16 fd0, rh1;
+; CHECK: cvt.rn.f64.u16 rd0, rh1;
 ; CHECK: ret;
 	%a = uitofp i16 %x to double
 	ret double %a
 }
 
 define ptx_device double @cvt_f64_i32(i32 %x) {
-; CHECK: cvt.rn.f64.u32 fd0, r1;
+; CHECK: cvt.rn.f64.u32 rd0, r1;
 ; CHECK: ret;
 	%a = uitofp i32 %x to double
 	ret double %a
 }
 
 define ptx_device double @cvt_f64_i64(i64 %x) {
-; CHECK: cvt.rn.f64.u64 fd0, rd1;
+; CHECK: cvt.rn.f64.u64 rd0, rd1;
 ; CHECK: ret;
 	%a = uitofp i64 %x to double
 	ret double %a
 }
 
 define ptx_device double @cvt_f64_f32(float %x) {
-; CHECK: cvt.f64.f32 fd0, f1;
+; CHECK: cvt.f64.f32 rd0, r1;
 ; CHECK: ret;
 	%a = fpext float %x to double
 	ret double %a
diff --git a/test/CodeGen/PTX/fdiv-sm10.ll b/test/CodeGen/PTX/fdiv-sm10.ll
index 121360c..9aff251 100644
--- a/test/CodeGen/PTX/fdiv-sm10.ll
+++ b/test/CodeGen/PTX/fdiv-sm10.ll
@@ -1,14 +1,14 @@
 ; 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;
+; CHECK: div.approx.f32 r0, r1, r2;
 ; CHECK-NEXT: ret;
 	%a = fdiv float %x, %y
 	ret float %a
 }
 
 define ptx_device double @t1_f64(double %x, double %y) {
-; CHECK: div.f64 fd0, fd1, fd2;
+; CHECK: div.f64 rd0, rd1, rd2;
 ; CHECK-NEXT: ret;
 	%a = fdiv double %x, %y
 	ret double %a
diff --git a/test/CodeGen/PTX/fdiv-sm13.ll b/test/CodeGen/PTX/fdiv-sm13.ll
index 0ec7bae..84e0ada 100644
--- a/test/CodeGen/PTX/fdiv-sm13.ll
+++ b/test/CodeGen/PTX/fdiv-sm13.ll
@@ -1,14 +1,14 @@
 ; 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;
+; CHECK: div.approx.f32 r0, r1, r2;
 ; CHECK-NEXT: ret;
 	%a = fdiv float %x, %y
 	ret float %a
 }
 
 define ptx_device double @t1_f64(double %x, double %y) {
-; CHECK: div.rn.f64 fd0, fd1, fd2;
+; CHECK: div.rn.f64 rd0, rd1, rd2;
 ; CHECK-NEXT: ret;
 	%a = fdiv double %x, %y
 	ret double %a
diff --git a/test/CodeGen/PTX/fneg.ll b/test/CodeGen/PTX/fneg.ll
index 22eeda3..185c37c 100644
--- a/test/CodeGen/PTX/fneg.ll
+++ b/test/CodeGen/PTX/fneg.ll
@@ -1,14 +1,14 @@
 ; RUN: llc < %s -march=ptx32 | FileCheck %s
 
 define ptx_device float @t1_f32(float %x) {
-; CHECK: neg.f32 f0, f1;
+; CHECK: neg.f32 r0, r1;
 ; CHECK-NEXT: ret;
 	%y = fsub float -0.000000e+00, %x
 	ret float %y
 }
 
 define ptx_device double @t1_f64(double %x) {
-; CHECK: neg.f64 fd0, fd1;
+; CHECK: neg.f64 rd0, rd1;
 ; CHECK-NEXT: ret;
 	%y = fsub double -0.000000e+00, %x
 	ret double %y
diff --git a/test/CodeGen/PTX/ld.ll b/test/CodeGen/PTX/ld.ll
index 377a95a..9b75998 100644
--- a/test/CodeGen/PTX/ld.ll
+++ b/test/CodeGen/PTX/ld.ll
@@ -87,7 +87,7 @@ entry:
 
 define ptx_device float @t1_f32(float* %p) {
 entry:
-;CHECK: ld.global.f32 f0, [r1];
+;CHECK: ld.global.f32 r0, [r1];
 ;CHECK-NEXT: ret;
   %x = load float* %p
   ret float %x
@@ -95,7 +95,7 @@ entry:
 
 define ptx_device double @t1_f64(double* %p) {
 entry:
-;CHECK: ld.global.f64 fd0, [r1];
+;CHECK: ld.global.f64 rd0, [r1];
 ;CHECK-NEXT: ret;
   %x = load double* %p
   ret double %x
@@ -130,7 +130,7 @@ entry:
 
 define ptx_device float @t2_f32(float* %p) {
 entry:
-;CHECK: ld.global.f32 f0, [r1+4];
+;CHECK: ld.global.f32 r0, [r1+4];
 ;CHECK-NEXT: ret;
   %i = getelementptr float* %p, i32 1
   %x = load float* %i
@@ -139,7 +139,7 @@ entry:
 
 define ptx_device double @t2_f64(double* %p) {
 entry:
-;CHECK: ld.global.f64 fd0, [r1+8];
+;CHECK: ld.global.f64 rd0, [r1+8];
 ;CHECK-NEXT: ret;
   %i = getelementptr double* %p, i32 1
   %x = load double* %i
@@ -180,7 +180,7 @@ define ptx_device float @t3_f32(float* %p, i32 %q) {
 entry:
 ;CHECK: shl.b32 r0, r2, 2;
 ;CHECK-NEXT: add.u32 r0, r1, r0;
-;CHECK-NEXT: ld.global.f32 f0, [r0];
+;CHECK-NEXT: ld.global.f32 r0, [r0];
   %i = getelementptr float* %p, i32 %q
   %x = load float* %i
   ret float %x
@@ -190,7 +190,7 @@ define ptx_device double @t3_f64(double* %p, i32 %q) {
 entry:
 ;CHECK: shl.b32 r0, r2, 3;
 ;CHECK-NEXT: add.u32 r0, r1, r0;
-;CHECK-NEXT: ld.global.f64 fd0, [r0];
+;CHECK-NEXT: ld.global.f64 rd0, [r0];
   %i = getelementptr double* %p, i32 %q
   %x = load double* %i
   ret double %x
@@ -229,7 +229,7 @@ entry:
 define ptx_device float @t4_global_f32() {
 entry:
 ;CHECK: mov.u32 r0, array_float;
-;CHECK-NEXT: ld.global.f32 f0, [r0];
+;CHECK-NEXT: ld.global.f32 r0, [r0];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x float]* @array_float, i32 0, i32 0
   %x = load float* %i
@@ -239,7 +239,7 @@ entry:
 define ptx_device double @t4_global_f64() {
 entry:
 ;CHECK: mov.u32 r0, array_double;
-;CHECK-NEXT: ld.global.f64 fd0, [r0];
+;CHECK-NEXT: ld.global.f64 rd0, [r0];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x double]* @array_double, i32 0, i32 0
   %x = load double* %i
@@ -279,7 +279,7 @@ entry:
 define ptx_device float @t4_const_f32() {
 entry:
 ;CHECK: mov.u32 r0, array_constant_float;
-;CHECK-NEXT: ld.const.f32 f0, [r0];
+;CHECK-NEXT: ld.const.f32 r0, [r0];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x float] addrspace(1)* @array_constant_float, i32 0, i32 0
   %x = load float addrspace(1)* %i
@@ -289,7 +289,7 @@ entry:
 define ptx_device double @t4_const_f64() {
 entry:
 ;CHECK: mov.u32 r0, array_constant_double;
-;CHECK-NEXT: ld.const.f64 fd0, [r0];
+;CHECK-NEXT: ld.const.f64 rd0, [r0];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x double] addrspace(1)* @array_constant_double, i32 0, i32 0
   %x = load double addrspace(1)* %i
@@ -329,7 +329,7 @@ entry:
 define ptx_device float @t4_local_f32() {
 entry:
 ;CHECK: mov.u32 r0, array_local_float;
-;CHECK-NEXT: ld.local.f32 f0, [r0];
+;CHECK-NEXT: ld.local.f32 r0, [r0];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x float] addrspace(2)* @array_local_float, i32 0, i32 0
   %x = load float addrspace(2)* %i
@@ -339,7 +339,7 @@ entry:
 define ptx_device double @t4_local_f64() {
 entry:
 ;CHECK: mov.u32 r0, array_local_double;
-;CHECK-NEXT: ld.local.f64 fd0, [r0];
+;CHECK-NEXT: ld.local.f64 rd0, [r0];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x double] addrspace(2)* @array_local_double, i32 0, i32 0
   %x = load double addrspace(2)* %i
@@ -379,7 +379,7 @@ entry:
 define ptx_device float @t4_shared_f32() {
 entry:
 ;CHECK: mov.u32 r0, array_shared_float;
-;CHECK-NEXT: ld.shared.f32 f0, [r0];
+;CHECK-NEXT: ld.shared.f32 r0, [r0];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x float] addrspace(4)* @array_shared_float, i32 0, i32 0
   %x = load float addrspace(4)* %i
@@ -389,7 +389,7 @@ entry:
 define ptx_device double @t4_shared_f64() {
 entry:
 ;CHECK: mov.u32 r0, array_shared_double;
-;CHECK-NEXT: ld.shared.f64 fd0, [r0];
+;CHECK-NEXT: ld.shared.f64 rd0, [r0];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x double] addrspace(4)* @array_shared_double, i32 0, i32 0
   %x = load double addrspace(4)* %i
@@ -429,7 +429,7 @@ entry:
 define ptx_device float @t5_f32() {
 entry:
 ;CHECK: mov.u32 r0, array_float;
-;CHECK-NEXT: ld.global.f32 f0, [r0+4];
+;CHECK-NEXT: ld.global.f32 r0, [r0+4];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x float]* @array_float, i32 0, i32 1
   %x = load float* %i
@@ -439,7 +439,7 @@ entry:
 define ptx_device double @t5_f64() {
 entry:
 ;CHECK: mov.u32 r0, array_double;
-;CHECK-NEXT: ld.global.f64 fd0, [r0+8];
+;CHECK-NEXT: ld.global.f64 rd0, [r0+8];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x double]* @array_double, i32 0, i32 1
   %x = load double* %i
diff --git a/test/CodeGen/PTX/llvm-intrinsic.ll b/test/CodeGen/PTX/llvm-intrinsic.ll
index 1e265f5..a317645 100644
--- a/test/CodeGen/PTX/llvm-intrinsic.ll
+++ b/test/CodeGen/PTX/llvm-intrinsic.ll
@@ -2,7 +2,7 @@
 
 define ptx_device float @test_sqrt_f32(float %x) {
 entry:
-; CHECK: sqrt.rn.f32 f0, f1;
+; CHECK: sqrt.rn.f32 r0, r1;
 ; CHECK-NEXT: ret;
   %y = call float @llvm.sqrt.f32(float %x)
   ret float %y
@@ -10,7 +10,7 @@ entry:
 
 define ptx_device double @test_sqrt_f64(double %x) {
 entry:
-; CHECK: sqrt.rn.f64 fd0, fd1;
+; CHECK: sqrt.rn.f64 rd0, rd1;
 ; CHECK-NEXT: ret;
   %y = call double @llvm.sqrt.f64(double %x)
   ret double %y
@@ -18,7 +18,7 @@ entry:
 
 define ptx_device float @test_sin_f32(float %x) {
 entry:
-; CHECK: sin.approx.f32 f0, f1;
+; CHECK: sin.approx.f32 r0, r1;
 ; CHECK-NEXT: ret;
   %y = call float @llvm.sin.f32(float %x)
   ret float %y
@@ -26,7 +26,7 @@ entry:
 
 define ptx_device double @test_sin_f64(double %x) {
 entry:
-; CHECK: sin.approx.f64 fd0, fd1;
+; CHECK: sin.approx.f64 rd0, rd1;
 ; CHECK-NEXT: ret;
   %y = call double @llvm.sin.f64(double %x)
   ret double %y
@@ -34,7 +34,7 @@ entry:
 
 define ptx_device float @test_cos_f32(float %x) {
 entry:
-; CHECK: cos.approx.f32 f0, f1;
+; CHECK: cos.approx.f32 r0, r1;
 ; CHECK-NEXT: ret;
   %y = call float @llvm.cos.f32(float %x)
   ret float %y
@@ -42,7 +42,7 @@ entry:
 
 define ptx_device double @test_cos_f64(double %x) {
 entry:
-; CHECK: cos.approx.f64 fd0, fd1;
+; CHECK: cos.approx.f64 rd0, rd1;
 ; CHECK-NEXT: ret;
   %y = call double @llvm.cos.f64(double %x)
   ret double %y
diff --git a/test/CodeGen/PTX/mad.ll b/test/CodeGen/PTX/mad.ll
index 0c25f2c..56d3811 100644
--- a/test/CodeGen/PTX/mad.ll
+++ b/test/CodeGen/PTX/mad.ll
@@ -1,7 +1,7 @@
 ; 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;
+; CHECK: mad.rn.f32 r0, r1, r2, r3;
 ; CHECK-NEXT: ret;
 	%a = fmul float %x, %y
   %b = fadd float %a, %z
@@ -9,7 +9,7 @@ define ptx_device float @t1_f32(float %x, float %y, float %z) {
 }
 
 define ptx_device double @t1_f64(double %x, double %y, double %z) {
-; CHECK: mad.rn.f64 fd0, fd1, fd2, fd3;
+; CHECK: mad.rn.f64 rd0, rd1, rd2, rd3;
 ; CHECK-NEXT: ret;
 	%a = fmul double %x, %y
   %b = fadd double %a, %z
diff --git a/test/CodeGen/PTX/mov.ll b/test/CodeGen/PTX/mov.ll
index 120572a..05ce4c0 100644
--- a/test/CodeGen/PTX/mov.ll
+++ b/test/CodeGen/PTX/mov.ll
@@ -19,13 +19,13 @@ define ptx_device i64 @t1_u64() {
 }
 
 define ptx_device float @t1_f32() {
-; CHECK: mov.f32 f0, 0F00000000;
+; CHECK: mov.f32 r0, 0F00000000;
 ; CHECK: ret;
 	ret float 0.0
 }
 
 define ptx_device double @t1_f64() {
-; CHECK: mov.f64 fd0, 0D0000000000000000;
+; CHECK: mov.f64 rd0, 0D0000000000000000;
 ; CHECK: ret;
 	ret double 0.0
 }
@@ -49,13 +49,13 @@ define ptx_device i64 @t2_u64(i64 %x) {
 }
 
 define ptx_device float @t3_f32(float %x) {
-; CHECK: mov.f32 f0, f1;
+; CHECK: mov.u32 r0, r1;
 ; CHECK-NEXT: ret;
 	ret float %x
 }
 
 define ptx_device double @t3_f64(double %x) {
-; CHECK: mov.f64 fd0, fd1;
+; CHECK: mov.u64 rd0, rd1;
 ; CHECK-NEXT: ret;
 	ret double %x
 }
diff --git a/test/CodeGen/PTX/mul.ll b/test/CodeGen/PTX/mul.ll
index 5ce0426..93f94e3 100644
--- a/test/CodeGen/PTX/mul.ll
+++ b/test/CodeGen/PTX/mul.ll
@@ -11,28 +11,28 @@
 ;}
 
 define ptx_device float @t1_f32(float %x, float %y) {
-; CHECK: mul.f32 f0, f1, f2
+; CHECK: mul.f32 r0, r1, r2
 ; CHECK-NEXT: ret;
   %z = fmul float %x, %y
   ret float %z
 }
 
 define ptx_device double @t1_f64(double %x, double %y) {
-; CHECK: mul.f64 fd0, fd1, fd2
+; CHECK: mul.f64 rd0, rd1, rd2
 ; CHECK-NEXT: ret;
   %z = fmul double %x, %y
   ret double %z
 }
 
 define ptx_device float @t2_f32(float %x) {
-; CHECK: mul.f32 f0, f1, 0F40A00000;
+; CHECK: mul.f32 r0, r1, 0F40A00000;
 ; CHECK-NEXT: ret;
   %z = fmul float %x, 5.0
   ret float %z
 }
 
 define ptx_device double @t2_f64(double %x) {
-; CHECK: mul.f64 fd0, fd1, 0D4014000000000000;
+; CHECK: mul.f64 rd0, rd1, 0D4014000000000000;
 ; CHECK-NEXT: ret;
   %z = fmul double %x, 5.0
   ret double %z
diff --git a/test/CodeGen/PTX/parameter-order.ll b/test/CodeGen/PTX/parameter-order.ll
index 95d4a32..5486472 100644
--- a/test/CodeGen/PTX/parameter-order.ll
+++ b/test/CodeGen/PTX/parameter-order.ll
@@ -1,8 +1,8 @@
 ; RUN: llc < %s -march=ptx32 | FileCheck %s
 
-; CHECK: .func (.reg .u32 r0) test_parameter_order (.reg .f32 f1, .reg .u32 r1, .reg .u32 r2, .reg .f32 f2)
+; CHECK: .func (.reg .b32 r0) test_parameter_order (.reg .b32 r1, .reg .b32 r2, .reg .b32 r3, .reg .b32 r4)
 define ptx_device i32 @test_parameter_order(float %a, i32 %b, i32 %c, float %d) {
-; CHECK: sub.u32 r0, r1, r2
+; CHECK: sub.u32 r0, r2, r3
 	%result = sub i32 %b, %c
 	ret i32 %result
 }
diff --git a/test/CodeGen/PTX/selp.ll b/test/CodeGen/PTX/selp.ll
index 6f1b03e..19cfa53 100644
--- a/test/CodeGen/PTX/selp.ll
+++ b/test/CodeGen/PTX/selp.ll
@@ -13,13 +13,13 @@ define ptx_device i64 @test_selp_i64(i1 %x, i64 %y, i64 %z) {
 }
 
 define ptx_device float @test_selp_f32(i1 %x, float %y, float %z) {
-; CHECK: selp.f32 f0, f1, f2, p1;
+; CHECK: selp.f32 r0, r1, r2, p1;
 	%a = select i1 %x, float %y, float %z
 	ret float %a
 }
 
 define ptx_device double @test_selp_f64(i1 %x, double %y, double %z) {
-; CHECK: selp.f64 fd0, fd1, fd2, p1;
+; CHECK: selp.f64 rd0, rd1, rd2, p1;
 	%a = select i1 %x, double %y, double %z
 	ret double %a
 }
diff --git a/test/CodeGen/PTX/st.ll b/test/CodeGen/PTX/st.ll
index dee5c61..612967a 100644
--- a/test/CodeGen/PTX/st.ll
+++ b/test/CodeGen/PTX/st.ll
@@ -87,7 +87,7 @@ entry:
 
 define ptx_device void @t1_f32(float* %p, float %x) {
 entry:
-;CHECK: st.global.f32 [r1], f1;
+;CHECK: st.global.f32 [r1], r2;
 ;CHECK-NEXT: ret;
   store float %x, float* %p
   ret void
@@ -95,7 +95,7 @@ entry:
 
 define ptx_device void @t1_f64(double* %p, double %x) {
 entry:
-;CHECK: st.global.f64 [r1], fd1;
+;CHECK: st.global.f64 [r1], rd1;
 ;CHECK-NEXT: ret;
   store double %x, double* %p
   ret void
@@ -130,7 +130,7 @@ entry:
 
 define ptx_device void @t2_f32(float* %p, float %x) {
 entry:
-;CHECK: st.global.f32 [r1+4], f1;
+;CHECK: st.global.f32 [r1+4], r2;
 ;CHECK-NEXT: ret;
   %i = getelementptr float* %p, i32 1
   store float %x, float* %i
@@ -139,7 +139,7 @@ entry:
 
 define ptx_device void @t2_f64(double* %p, double %x) {
 entry:
-;CHECK: st.global.f64 [r1+8], fd1;
+;CHECK: st.global.f64 [r1+8], rd1;
 ;CHECK-NEXT: ret;
   %i = getelementptr double* %p, i32 1
   store double %x, double* %i
@@ -183,7 +183,7 @@ define ptx_device void @t3_f32(float* %p, i32 %q, float %x) {
 entry:
 ;CHECK: shl.b32 r0, r2, 2;
 ;CHECK-NEXT: add.u32 r0, r1, r0;
-;CHECK-NEXT: st.global.f32 [r0], f1;
+;CHECK-NEXT: st.global.f32 [r0], r3;
 ;CHECK-NEXT: ret;
   %i = getelementptr float* %p, i32 %q
   store float %x, float* %i
@@ -194,7 +194,7 @@ define ptx_device void @t3_f64(double* %p, i32 %q, double %x) {
 entry:
 ;CHECK: shl.b32 r0, r2, 3;
 ;CHECK-NEXT: add.u32 r0, r1, r0;
-;CHECK-NEXT: st.global.f64 [r0], fd1;
+;CHECK-NEXT: st.global.f64 [r0], rd1;
 ;CHECK-NEXT: ret;
   %i = getelementptr double* %p, i32 %q
   store double %x, double* %i
@@ -234,7 +234,7 @@ entry:
 define ptx_device void @t4_global_f32(float %x) {
 entry:
 ;CHECK: mov.u32 r0, array_float;
-;CHECK-NEXT: st.global.f32 [r0], f1;
+;CHECK-NEXT: st.global.f32 [r0], r1;
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x float]* @array_float, i32 0, i32 0
   store float %x, float* %i
@@ -244,7 +244,7 @@ entry:
 define ptx_device void @t4_global_f64(double %x) {
 entry:
 ;CHECK: mov.u32 r0, array_double;
-;CHECK-NEXT: st.global.f64 [r0], fd1;
+;CHECK-NEXT: st.global.f64 [r0], rd1;
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x double]* @array_double, i32 0, i32 0
   store double %x, double* %i
@@ -284,7 +284,7 @@ entry:
 define ptx_device void @t4_local_f32(float %x) {
 entry:
 ;CHECK: mov.u32 r0, array_local_float;
-;CHECK-NEXT: st.local.f32 [r0], f1;
+;CHECK-NEXT: st.local.f32 [r0], r1;
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x float] addrspace(2)* @array_local_float, i32 0, i32 0
   store float %x, float addrspace(2)* %i
@@ -294,7 +294,7 @@ entry:
 define ptx_device void @t4_local_f64(double %x) {
 entry:
 ;CHECK: mov.u32 r0, array_local_double;
-;CHECK-NEXT: st.local.f64 [r0], fd1;
+;CHECK-NEXT: st.local.f64 [r0], rd1;
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x double] addrspace(2)* @array_local_double, i32 0, i32 0
   store double %x, double addrspace(2)* %i
@@ -334,7 +334,7 @@ entry:
 define ptx_device void @t4_shared_f32(float %x) {
 entry:
 ;CHECK: mov.u32 r0, array_shared_float;
-;CHECK-NEXT: st.shared.f32 [r0], f1;
+;CHECK-NEXT: st.shared.f32 [r0], r1;
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x float] addrspace(4)* @array_shared_float, i32 0, i32 0
   store float %x, float addrspace(4)* %i
@@ -344,7 +344,7 @@ entry:
 define ptx_device void @t4_shared_f64(double %x) {
 entry:
 ;CHECK: mov.u32 r0, array_shared_double;
-;CHECK-NEXT: st.shared.f64 [r0], fd1;
+;CHECK-NEXT: st.shared.f64 [r0], rd1;
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x double] addrspace(4)* @array_shared_double, i32 0, i32 0
   store double %x, double addrspace(4)* %i
@@ -384,7 +384,7 @@ entry:
 define ptx_device void @t5_f32(float %x) {
 entry:
 ;CHECK: mov.u32 r0, array_float;
-;CHECK-NEXT: st.global.f32 [r0+4], f1;
+;CHECK-NEXT: st.global.f32 [r0+4], r1;
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x float]* @array_float, i32 0, i32 1
   store float %x, float* %i
@@ -394,7 +394,7 @@ entry:
 define ptx_device void @t5_f64(double %x) {
 entry:
 ;CHECK: mov.u32 r0, array_double;
-;CHECK-NEXT: st.global.f64 [r0+8], fd1;
+;CHECK-NEXT: st.global.f64 [r0+8], rd1;
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x double]* @array_double, i32 0, i32 1
   store double %x, double* %i
diff --git a/test/CodeGen/PTX/sub.ll b/test/CodeGen/PTX/sub.ll
index 7dd2c6f..9efeaac 100644
--- a/test/CodeGen/PTX/sub.ll
+++ b/test/CodeGen/PTX/sub.ll
@@ -22,14 +22,14 @@ define ptx_device i64 @t1_u64(i64 %x, i64 %y) {
 }
 
 define ptx_device float @t1_f32(float %x, float %y) {
-; CHECK: sub.f32 f0, f1, f2
+; CHECK: sub.f32 r0, r1, r2
 ; CHECK-NEXT: ret;
   %z = fsub float %x, %y
   ret float %z
 }
 
 define ptx_device double @t1_f64(double %x, double %y) {
-; CHECK: sub.f64 fd0, fd1, fd2
+; CHECK: sub.f64 rd0, rd1, rd2
 ; CHECK-NEXT: ret;
   %z = fsub double %x, %y
   ret double %z
@@ -57,14 +57,14 @@ define ptx_device i64 @t2_u64(i64 %x) {
 }
 
 define ptx_device float @t2_f32(float %x) {
-; CHECK: add.f32 f0, f1, 0FBF800000;
+; CHECK: add.f32 r0, r1, 0FBF800000;
 ; CHECK-NEXT: ret;
   %z = fsub float %x, 1.0
   ret float %z
 }
 
 define ptx_device double @t2_f64(double %x) {
-; CHECK: add.f64 fd0, fd1, 0DBFF0000000000000;
+; CHECK: add.f64 rd0, rd1, 0DBFF0000000000000;
 ; CHECK-NEXT: ret;
   %z = fsub double %x, 1.0
   ret double %z