PTX: Always use registers for return values, but use .param space for device

     parameters if SM >= 2.0

- Update test cases to be more robust against register allocation changes
- Bump up the number of registers to 128 per type
- Include Python script to re-generate register file with any number of
  registers

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

20 files changed, 318 insertions, 295 deletions

diff --git a/test/CodeGen/PTX/add.ll b/test/CodeGen/PTX/add.ll
index c16be4933c..293aebe51e 100644
--- a/test/CodeGen/PTX/add.ll
+++ b/test/CodeGen/PTX/add.ll
@@ -1,70 +1,70 @@
 ; RUN: llc < %s -march=ptx32 | FileCheck %s
 
 define ptx_device i16 @t1_u16(i16 %x, i16 %y) {
-; CHECK: add.u16 rh0, rh1, rh2;
+; CHECK: add.u16 rh{{[0-9]+}}, rh{{[0-9]+}}, rh{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%z = add i16 %x, %y
 	ret i16 %z
 }
 
 define ptx_device i32 @t1_u32(i32 %x, i32 %y) {
-; CHECK: add.u32 r0, r1, r2;
+; CHECK: add.u32 r{{[0-9]+}}, r{{[0-9]+}}, r{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%z = add i32 %x, %y
 	ret i32 %z
 }
 
 define ptx_device i64 @t1_u64(i64 %x, i64 %y) {
-; CHECK: add.u64 rd0, rd1, rd2;
+; CHECK: add.u64 rd{{[0-9]+}}, rd{{[0-9]+}}, rd{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%z = add i64 %x, %y
 	ret i64 %z
 }
 
 define ptx_device float @t1_f32(float %x, float %y) {
-; CHECK: add.rn.f32 r0, r1, r2
+; CHECK: add.rn.f32 r{{[0-9]+}}, r{{[0-9]+}}, r{{[0-9]+}}
 ; CHECK-NEXT: ret;
   %z = fadd float %x, %y
   ret float %z
 }
 
 define ptx_device double @t1_f64(double %x, double %y) {
-; CHECK: add.rn.f64 rd0, rd1, rd2
+; CHECK: add.rn.f64 rd{{[0-9]+}}, rd{{[0-9]+}}, rd{{[0-9]+}}
 ; CHECK-NEXT: ret;
   %z = fadd double %x, %y
   ret double %z
 }
 
 define ptx_device i16 @t2_u16(i16 %x) {
-; CHECK: add.u16 rh0, rh1, 1;
+; CHECK: add.u16 rh{{[0-9]+}}, rh{{[0-9]+}}, 1;
 ; CHECK-NEXT: ret;
 	%z = add i16 %x, 1
 	ret i16 %z
 }
 
 define ptx_device i32 @t2_u32(i32 %x) {
-; CHECK: add.u32 r0, r1, 1;
+; CHECK: add.u32 r{{[0-9]+}}, r{{[0-9]+}}, 1;
 ; CHECK-NEXT: ret;
 	%z = add i32 %x, 1
 	ret i32 %z
 }
 
 define ptx_device i64 @t2_u64(i64 %x) {
-; CHECK: add.u64 rd0, rd1, 1;
+; CHECK: add.u64 rd{{[0-9]+}}, rd{{[0-9]+}}, 1;
 ; CHECK-NEXT: ret;
 	%z = add i64 %x, 1
 	ret i64 %z
 }
 
 define ptx_device float @t2_f32(float %x) {
-; CHECK: add.rn.f32 r0, r1, 0F3F800000;
+; CHECK: add.rn.f32 r{{[0-9]+}}, r{{[0-9]+}}, 0F3F800000;
 ; CHECK-NEXT: ret;
   %z = fadd float %x, 1.0
   ret float %z
 }
 
 define ptx_device double @t2_f64(double %x) {
-; CHECK: add.rn.f64 rd0, rd1, 0D3FF0000000000000;
+; CHECK: add.rn.f64 rd{{[0-9]+}}, rd{{[0-9]+}}, 0D3FF0000000000000;
 ; CHECK-NEXT: ret;
   %z = fadd double %x, 1.0
   ret double %z
diff --git a/test/CodeGen/PTX/aggregates.ll b/test/CodeGen/PTX/aggregates.ll
new file mode 100644
index 0000000000..23f28a79c0
--- /dev/null
+++ b/test/CodeGen/PTX/aggregates.ll
@@ -0,0 +1,23 @@
+; RUN: llc < %s -march=ptx32 -mattr=sm20 | FileCheck %s
+
+%complex = type { float, float }
+
+define ptx_device %complex @complex_add(%complex %a, %complex %b) {
+entry:
+; CHECK:      ld.param.f32	r[[R0:[0-9]+]], [__param_1];
+; CHECK-NEXT:	ld.param.f32	r[[R2:[0-9]+]], [__param_3];
+; CHECK-NEXT:	ld.param.f32	r[[R1:[0-9]+]], [__param_2];
+; CHECK-NEXT:	ld.param.f32	r[[R3:[0-9]+]], [__param_4];
+; CHECK-NEXT:	add.rn.f32	r[[R0]], r[[R0]], r[[R2]];
+; CHECK-NEXT:	add.rn.f32	r[[R1]], r[[R1]], r[[R3]];
+; CHECK-NEXT:	ret;
+  %a.real = extractvalue %complex %a, 0
+  %a.imag = extractvalue %complex %a, 1
+  %b.real = extractvalue %complex %b, 0
+  %b.imag = extractvalue %complex %b, 1
+  %ret.real = fadd float %a.real, %b.real
+  %ret.imag = fadd float %a.imag, %b.imag
+  %ret.0 = insertvalue %complex undef, float %ret.real, 0
+  %ret.1 = insertvalue %complex %ret.0, float %ret.imag, 1
+  ret %complex %ret.1
+}
diff --git a/test/CodeGen/PTX/bitwise.ll b/test/CodeGen/PTX/bitwise.ll
index dbc77e5333..3859280735 100644
--- a/test/CodeGen/PTX/bitwise.ll
+++ b/test/CodeGen/PTX/bitwise.ll
@@ -3,21 +3,21 @@
 ; preds
 
 define ptx_device i32 @t1_and_preds(i1 %x, i1 %y) {
-; CHECK: and.pred p0, p1, p2
+; CHECK: and.pred p{{[0-9]+}}, p{{[0-9]+}}, p{{[0-9]+}}
   %c = and i1 %x, %y
   %d = zext i1 %c to i32 
   ret i32 %d
 }
 
 define ptx_device i32 @t1_or_preds(i1 %x, i1 %y) {
-; CHECK: or.pred p0, p1, p2
+; CHECK: or.pred p{{[0-9]+}}, p{{[0-9]+}}, p{{[0-9]+}}
   %a = or i1 %x, %y
   %b = zext i1 %a to i32 
   ret i32 %b
 }
 
 define ptx_device i32 @t1_xor_preds(i1 %x, i1 %y) {
-; CHECK: xor.pred p0, p1, p2
+; CHECK: xor.pred p{{[0-9]+}}, p{{[0-9]+}}, p{{[0-9]+}}
   %a = xor i1 %x, %y
   %b = zext i1 %a to i32 
   ret i32 %b
diff --git a/test/CodeGen/PTX/bra.ll b/test/CodeGen/PTX/bra.ll
index 49383eb3cf..7cc944466d 100644
--- a/test/CodeGen/PTX/bra.ll
+++ b/test/CodeGen/PTX/bra.ll
@@ -10,15 +10,15 @@ loop:
 
 define ptx_device i32 @test_bra_cond_direct(i32 %x, i32 %y) {
 entry:
-; CHECK: setp.le.u32 p0, r1, r2
+; CHECK: setp.le.u32 p0, r[[R0:[0-9]+]], r[[R1:[0-9]+]]
 	%p = icmp ugt i32 %x, %y
 ; CHECK-NEXT: @p0 bra
 ; CHECK-NOT: bra
 	br i1 %p, label %clause.if, label %clause.else
 clause.if:
-; CHECK: mov.u32 r0, r1
+; CHECK: mov.u32 r{{[0-9]+}}, r[[R0]]
 	ret i32 %x
 clause.else:
-; CHECK: mov.u32 r0, r2
+; CHECK: mov.u32 r{{[0-9]+}}, r[[R1]]
 	ret i32 %y
 }
diff --git a/test/CodeGen/PTX/cvt.ll b/test/CodeGen/PTX/cvt.ll
index f7233697c0..18f7ef365b 100644
--- a/test/CodeGen/PTX/cvt.ll
+++ b/test/CodeGen/PTX/cvt.ll
@@ -4,9 +4,9 @@
 ; (note: we convert back to i32 to return)
 
 define ptx_device i32 @cvt_pred_i16(i16 %x, i1 %y) {
-; CHECK: setp.gt.b16 p0, rh1, 0
-; CHECK-NEXT: and.pred p0, p0, p1;
-; CHECK-NEXT: selp.u32 r0, 1, 0, p0;
+; CHECK: setp.gt.b16 p[[P0:[0-9]+]], rh{{[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 i16 %x to i1
 	%b = and i1 %a, %y
@@ -15,9 +15,9 @@ define ptx_device i32 @cvt_pred_i16(i16 %x, i1 %y) {
 }
 
 define ptx_device i32 @cvt_pred_i32(i32 %x, i1 %y) {
-; CHECK: setp.gt.b32 p0, r1, 0
-; CHECK-NEXT: and.pred p0, p0, p1;
-; CHECK-NEXT: selp.u32 r0, 1, 0, p0;
+; CHECK: setp.gt.b32 p[[P0:[0-9]+]], r{{[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 i32 %x to i1
 	%b = and i1 %a, %y
@@ -26,9 +26,9 @@ define ptx_device i32 @cvt_pred_i32(i32 %x, i1 %y) {
 }
 
 define ptx_device i32 @cvt_pred_i64(i64 %x, i1 %y) {
-; CHECK: setp.gt.b64 p0, rd1, 0
-; CHECK-NEXT: and.pred p0, p0, p1;
-; CHECK-NEXT: selp.u32 r0, 1, 0, p0;
+; CHECK: setp.gt.b64 p[[P0:[0-9]+]], rd{{[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 i64 %x to i1
 	%b = and i1 %a, %y
@@ -37,9 +37,9 @@ define ptx_device i32 @cvt_pred_i64(i64 %x, i1 %y) {
 }
 
 define ptx_device i32 @cvt_pred_f32(float %x, i1 %y) {
-; CHECK: setp.gt.b32 p0, r1, 0
-; CHECK-NEXT: and.pred p0, p0, p1;
-; CHECK-NEXT: selp.u32 r0, 1, 0, p0;
+; CHECK: setp.gt.b32 p[[P0:[0-9]+]], r{{[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 = fptoui float %x to i1
 	%b = and i1 %a, %y
@@ -48,9 +48,9 @@ define ptx_device i32 @cvt_pred_f32(float %x, i1 %y) {
 }
 
 define ptx_device i32 @cvt_pred_f64(double %x, i1 %y) {
-; CHECK: setp.gt.b64 p0, rd1, 0
-; CHECK-NEXT: and.pred p0, p0, p1;
-; CHECK-NEXT: selp.u32 r0, 1, 0, p0;
+; CHECK: setp.gt.b64 p[[P0:[0-9]+]], rd{{[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 = fptoui double %x to i1
 	%b = and i1 %a, %y
@@ -61,35 +61,35 @@ define ptx_device i32 @cvt_pred_f64(double %x, i1 %y) {
 ; i16
 
 define ptx_device i16 @cvt_i16_preds(i1 %x) {
-; CHECK: selp.u16 rh0, 1, 0, p1;
+; CHECK: selp.u16 rh{{[0-9]+}}, 1, 0, p{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%a = zext i1 %x to i16
 	ret i16 %a
 }
 
 define ptx_device i16 @cvt_i16_i32(i32 %x) {
-; CHECK: cvt.u16.u32 rh0, r1;
+; CHECK: cvt.u16.u32 rh{{[0-9]+}}, r{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%a = trunc i32 %x to i16
 	ret i16 %a
 }
 
 define ptx_device i16 @cvt_i16_i64(i64 %x) {
-; CHECK: cvt.u16.u64 rh0, rd1;
+; CHECK: cvt.u16.u64 rh{{[0-9]+}}, rd{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%a = trunc i64 %x to i16
 	ret i16 %a
 }
 
 define ptx_device i16 @cvt_i16_f32(float %x) {
-; CHECK: cvt.rzi.u16.f32 rh0, r1;
+; CHECK: cvt.rzi.u16.f32 rh{{[0-9]+}}, r{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%a = fptoui float %x to i16
 	ret i16 %a
 }
 
 define ptx_device i16 @cvt_i16_f64(double %x) {
-; CHECK: cvt.rzi.u16.f64 rh0, rd1;
+; CHECK: cvt.rzi.u16.f64 rh{{[0-9]+}}, rd{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%a = fptoui double %x to i16
 	ret i16 %a
@@ -98,35 +98,35 @@ define ptx_device i16 @cvt_i16_f64(double %x) {
 ; i32
 
 define ptx_device i32 @cvt_i32_preds(i1 %x) {
-; CHECK: selp.u32 r0, 1, 0, p1;
+; CHECK: selp.u32 r{{[0-9]+}}, 1, 0, p{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%a = zext i1 %x to i32
 	ret i32 %a
 }
 
 define ptx_device i32 @cvt_i32_i16(i16 %x) {
-; CHECK: cvt.u32.u16 r0, rh1;
+; CHECK: cvt.u32.u16 r{{[0-9]+}}, rh{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%a = zext i16 %x to i32
 	ret i32 %a
 }
 
 define ptx_device i32 @cvt_i32_i64(i64 %x) {
-; CHECK: cvt.u32.u64 r0, rd1;
+; CHECK: cvt.u32.u64 r{{[0-9]+}}, rd{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%a = trunc i64 %x to i32
 	ret i32 %a
 }
 
 define ptx_device i32 @cvt_i32_f32(float %x) {
-; CHECK: cvt.rzi.u32.f32 r0, r1;
+; CHECK: cvt.rzi.u32.f32 r{{[0-9]+}}, r{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%a = fptoui float %x to i32
 	ret i32 %a
 }
 
 define ptx_device i32 @cvt_i32_f64(double %x) {
-; CHECK: cvt.rzi.u32.f64 r0, rd1;
+; CHECK: cvt.rzi.u32.f64 r{{[0-9]+}}, rd{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%a = fptoui double %x to i32
 	ret i32 %a
@@ -135,35 +135,35 @@ define ptx_device i32 @cvt_i32_f64(double %x) {
 ; i64
 
 define ptx_device i64 @cvt_i64_preds(i1 %x) {
-; CHECK: selp.u64 rd0, 1, 0, p1;
+; CHECK: selp.u64 rd{{[0-9]+}}, 1, 0, p{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%a = zext i1 %x to i64
 	ret i64 %a
 }
 
 define ptx_device i64 @cvt_i64_i16(i16 %x) {
-; CHECK: cvt.u64.u16 rd0, rh1;
+; CHECK: cvt.u64.u16 rd{{[0-9]+}}, rh{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%a = zext i16 %x to i64
 	ret i64 %a
 }
 
 define ptx_device i64 @cvt_i64_i32(i32 %x) {
-; CHECK: cvt.u64.u32 rd0, r1;
+; CHECK: cvt.u64.u32 rd{{[0-9]+}}, r{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%a = zext i32 %x to i64
 	ret i64 %a
 }
 
 define ptx_device i64 @cvt_i64_f32(float %x) {
-; CHECK: cvt.rzi.u64.f32 rd0, r1;
+; CHECK: cvt.rzi.u64.f32 rd{{[0-9]+}}, r{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%a = fptoui float %x to i64
 	ret i64 %a
 }
 
 define ptx_device i64 @cvt_i64_f64(double %x) {
-; CHECK: cvt.rzi.u64.f64 rd0, rd1;
+; CHECK: cvt.rzi.u64.f64 rd{{[0-9]+}}, rd{{[0-9]+}};
 ; CHECK: ret;
 	%a = fptoui double %x to i64
 	ret i64 %a
@@ -172,35 +172,35 @@ define ptx_device i64 @cvt_i64_f64(double %x) {
 ; f32
 
 define ptx_device float @cvt_f32_preds(i1 %x) {
-; CHECK: selp.f32 r0, 0F3F800000, 0F00000000, p1;
+; CHECK: selp.f32 r{{[0-9]+}}, 0F3F800000, 0F00000000, p{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%a = uitofp i1 %x to float
 	ret float %a
 }
 
 define ptx_device float @cvt_f32_i16(i16 %x) {
-; CHECK: cvt.rn.f32.u16 r0, rh1;
+; CHECK: cvt.rn.f32.u16 r{{[0-9]+}}, rh{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%a = uitofp i16 %x to float
 	ret float %a
 }
 
 define ptx_device float @cvt_f32_i32(i32 %x) {
-; CHECK: cvt.rn.f32.u32 r0, r1;
+; CHECK: cvt.rn.f32.u32 r{{[0-9]+}}, r{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%a = uitofp i32 %x to float
 	ret float %a
 }
 
 define ptx_device float @cvt_f32_i64(i64 %x) {
-; CHECK: cvt.rn.f32.u64 r0, rd1;
+; CHECK: cvt.rn.f32.u64 r{{[0-9]+}}, rd{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%a = uitofp i64 %x to float
 	ret float %a
 }
 
 define ptx_device float @cvt_f32_f64(double %x) {
-; CHECK: cvt.rn.f32.f64 r0, rd1;
+; CHECK: cvt.rn.f32.f64 r{{[0-9]+}}, rd{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%a = fptrunc double %x to float
 	ret float %a
@@ -209,35 +209,35 @@ define ptx_device float @cvt_f32_f64(double %x) {
 ; f64
 
 define ptx_device double @cvt_f64_preds(i1 %x) {
-; CHECK: selp.f64 rd0, 0D3F80000000000000, 0D0000000000000000, p1;
+; CHECK: selp.f64 rd{{[0-9]+}}, 0D3F80000000000000, 0D0000000000000000, p{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%a = uitofp i1 %x to double
 	ret double %a
 }
 
 define ptx_device double @cvt_f64_i16(i16 %x) {
-; CHECK: cvt.rn.f64.u16 rd0, rh1;
+; CHECK: cvt.rn.f64.u16 rd{{[0-9]+}}, rh{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%a = uitofp i16 %x to double
 	ret double %a
 }
 
 define ptx_device double @cvt_f64_i32(i32 %x) {
-; CHECK: cvt.rn.f64.u32 rd0, r1;
+; CHECK: cvt.rn.f64.u32 rd{{[0-9]+}}, r{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%a = uitofp i32 %x to double
 	ret double %a
 }
 
 define ptx_device double @cvt_f64_i64(i64 %x) {
-; CHECK: cvt.rn.f64.u64 rd0, rd1;
+; CHECK: cvt.rn.f64.u64 rd{{[0-9]+}}, rd{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%a = uitofp i64 %x to double
 	ret double %a
 }
 
 define ptx_device double @cvt_f64_f32(float %x) {
-; CHECK: cvt.f64.f32 rd0, r1;
+; CHECK: cvt.f64.f32 rd{{[0-9]+}}, r{{[0-9]+}};
 ; CHECK-NEXT: 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 eb32222f3a..049d8913b3 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.f32 r0, r1, r2;
+; CHECK: div.f32 r{{[0-9]+}}, r{{[0-9]+}}, r{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%a = fdiv float %x, %y
 	ret float %a
 }
 
 define ptx_device double @t1_f64(double %x, double %y) {
-; CHECK: div.f64 rd0, rd1, rd2;
+; CHECK: div.f64 rd{{[0-9]+}}, rd{{[0-9]+}}, rd{{[0-9]+}};
 ; 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 ad24f35b20..2d953397d3 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.rn.f32 r0, r1, r2;
+; CHECK: div.rn.f32 r{{[0-9]+}}, r{{[0-9]+}}, r{{[0-9]+}};
 ; 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 rd0, rd1, rd2;
+; CHECK: div.rn.f64 rd{{[0-9]+}}, rd{{[0-9]+}}, rd{{[0-9]+}};
 ; 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 185c37c879..66ca74a6ff 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 r0, r1;
+; CHECK: neg.f32 r{{[0-9]+}}, r{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%y = fsub float -0.000000e+00, %x
 	ret float %y
 }
 
 define ptx_device double @t1_f64(double %x) {
-; CHECK: neg.f64 rd0, rd1;
+; CHECK: neg.f64 rd{{[0-9]+}}, rd{{[0-9]+}};
 ; 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 9b759987f8..d184d1243a 100644
--- a/test/CodeGen/PTX/ld.ll
+++ b/test/CodeGen/PTX/ld.ll
@@ -63,7 +63,7 @@
 
 define ptx_device i16 @t1_u16(i16* %p) {
 entry:
-;CHECK: ld.global.u16 rh0, [r1];
+;CHECK: ld.global.u16 rh{{[0-9]+}}, [r{{[0-9]+}}];
 ;CHECK-NEXT: ret;
   %x = load i16* %p
   ret i16 %x
@@ -71,7 +71,7 @@ entry:
 
 define ptx_device i32 @t1_u32(i32* %p) {
 entry:
-;CHECK: ld.global.u32 r0, [r1];
+;CHECK: ld.global.u32 r{{[0-9]+}}, [r{{[0-9]+}}];
 ;CHECK-NEXT: ret;
   %x = load i32* %p
   ret i32 %x
@@ -79,7 +79,7 @@ entry:
 
 define ptx_device i64 @t1_u64(i64* %p) {
 entry:
-;CHECK: ld.global.u64 rd0, [r1];
+;CHECK: ld.global.u64 rd{{[0-9]+}}, [r{{[0-9]+}}];
 ;CHECK-NEXT: ret;
   %x = load i64* %p
   ret i64 %x
@@ -87,7 +87,7 @@ entry:
 
 define ptx_device float @t1_f32(float* %p) {
 entry:
-;CHECK: ld.global.f32 r0, [r1];
+;CHECK: ld.global.f32 r{{[0-9]+}}, [r{{[0-9]+}}];
 ;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 rd0, [r1];
+;CHECK: ld.global.f64 rd{{[0-9]+}}, [r{{[0-9]+}}];
 ;CHECK-NEXT: ret;
   %x = load double* %p
   ret double %x
@@ -103,7 +103,7 @@ entry:
 
 define ptx_device i16 @t2_u16(i16* %p) {
 entry:
-;CHECK: ld.global.u16 rh0, [r1+2];
+;CHECK: ld.global.u16 rh{{[0-9]+}}, [r{{[0-9]+}}+2];
 ;CHECK-NEXT: ret;
   %i = getelementptr i16* %p, i32 1
   %x = load i16* %i
@@ -112,7 +112,7 @@ entry:
 
 define ptx_device i32 @t2_u32(i32* %p) {
 entry:
-;CHECK: ld.global.u32 r0, [r1+4];
+;CHECK: ld.global.u32 r{{[0-9]+}}, [r{{[0-9]+}}+4];
 ;CHECK-NEXT: ret;
   %i = getelementptr i32* %p, i32 1
   %x = load i32* %i
@@ -121,7 +121,7 @@ entry:
 
 define ptx_device i64 @t2_u64(i64* %p) {
 entry:
-;CHECK: ld.global.u64 rd0, [r1+8];
+;CHECK: ld.global.u64 rd{{[0-9]+}}, [r{{[0-9]+}}+8];
 ;CHECK-NEXT: ret;
   %i = getelementptr i64* %p, i32 1
   %x = load i64* %i
@@ -130,7 +130,7 @@ entry:
 
 define ptx_device float @t2_f32(float* %p) {
 entry:
-;CHECK: ld.global.f32 r0, [r1+4];
+;CHECK: ld.global.f32 r{{[0-9]+}}, [r{{[0-9]+}}+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 rd0, [r1+8];
+;CHECK: ld.global.f64 rd{{[0-9]+}}, [r{{[0-9]+}}+8];
 ;CHECK-NEXT: ret;
   %i = getelementptr double* %p, i32 1
   %x = load double* %i
@@ -148,9 +148,9 @@ entry:
 
 define ptx_device i16 @t3_u16(i16* %p, i32 %q) {
 entry:
-;CHECK: shl.b32 r0, r2, 1;
-;CHECK-NEXT: add.u32 r0, r1, r0;
-;CHECK-NEXT: ld.global.u16 rh0, [r0];
+;CHECK: shl.b32 r[[R0:[0-9]+]], r{{[0-9]+}}, 1;
+;CHECK-NEXT: add.u32 r[[R0]], r{{[0-9]+}}, r[[R0]];
+;CHECK-NEXT: ld.global.u16 rh{{[0-9]+}}, [r[[R0]]];
   %i = getelementptr i16* %p, i32 %q
   %x = load i16* %i
   ret i16 %x
@@ -158,9 +158,9 @@ entry:
 
 define ptx_device i32 @t3_u32(i32* %p, i32 %q) {
 entry:
-;CHECK: shl.b32 r0, r2, 2;
-;CHECK-NEXT: add.u32 r0, r1, r0;
-;CHECK-NEXT: ld.global.u32 r0, [r0];
+;CHECK: shl.b32 r[[R0:[0-9]+]], r{{[0-9]+}}, 2;
+;CHECK-NEXT: add.u32 r[[R0]], r{{[0-9]+}}, r[[R0]];
+;CHECK-NEXT: ld.global.u32 r{{[0-9]+}}, [r[[R0]]];
   %i = getelementptr i32* %p, i32 %q
   %x = load i32* %i
   ret i32 %x
@@ -168,9 +168,9 @@ entry:
 
 define ptx_device i64 @t3_u64(i64* %p, i32 %q) {
 entry:
-;CHECK: shl.b32 r0, r2, 3;
-;CHECK-NEXT: add.u32 r0, r1, r0;
-;CHECK-NEXT: ld.global.u64 rd0, [r0];
+;CHECK: shl.b32 r[[R0:[0-9]+]], r{{[0-9]+}}, 3;
+;CHECK-NEXT: add.u32 r[[R0]], r{{[0-9]+}}, r[[R0]];
+;CHECK-NEXT: ld.global.u64 rd{{[0-9]+}}, [r[[R0]]];
   %i = getelementptr i64* %p, i32 %q
   %x = load i64* %i
   ret i64 %x
@@ -178,9 +178,9 @@ entry:
 
 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 r0, [r0];
+;CHECK: shl.b32 r[[R0:[0-9]+]], r{{[0-9]+}}, 2;
+;CHECK-NEXT: add.u32 r[[R0]], r{{[0-9]+}}, r[[R0]];
+;CHECK-NEXT: ld.global.f32 r{{[0-9]+}}, [r[[R0]]];
   %i = getelementptr float* %p, i32 %q
   %x = load float* %i
   ret float %x
@@ -188,9 +188,9 @@ entry:
 
 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 rd0, [r0];
+;CHECK: shl.b32 r[[R0:[0-9]+]], r{{[0-9]+}}, 3;
+;CHECK-NEXT: add.u32 r[[R0]], r{{[0-9]+}}, r[[R0]];
+;CHECK-NEXT: ld.global.f64 rd{{[0-9]+}}, [r[[R0]]];
   %i = getelementptr double* %p, i32 %q
   %x = load double* %i
   ret double %x
@@ -198,8 +198,8 @@ entry:
 
 define ptx_device i16 @t4_global_u16() {
 entry:
-;CHECK: mov.u32 r0, array_i16;
-;CHECK-NEXT: ld.global.u16 rh0, [r0];
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_i16;
+;CHECK-NEXT: ld.global.u16 rh{{[0-9]+}}, [r[[R0]]];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x i16]* @array_i16, i32 0, i32 0
   %x = load i16* %i
@@ -208,8 +208,8 @@ entry:
 
 define ptx_device i32 @t4_global_u32() {
 entry:
-;CHECK: mov.u32 r0, array_i32;
-;CHECK-NEXT: ld.global.u32 r0, [r0];
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_i32;
+;CHECK-NEXT: ld.global.u32 r{{[0-9]+}}, [r[[R0]]];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x i32]* @array_i32, i32 0, i32 0
   %x = load i32* %i
@@ -218,8 +218,8 @@ entry:
 
 define ptx_device i64 @t4_global_u64() {
 entry:
-;CHECK: mov.u32 r0, array_i64;
-;CHECK-NEXT: ld.global.u64 rd0, [r0];
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_i64;
+;CHECK-NEXT: ld.global.u64 rd{{[0-9]+}}, [r[[R0]]];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x i64]* @array_i64, i32 0, i32 0
   %x = load i64* %i
@@ -228,8 +228,8 @@ entry:
 
 define ptx_device float @t4_global_f32() {
 entry:
-;CHECK: mov.u32 r0, array_float;
-;CHECK-NEXT: ld.global.f32 r0, [r0];
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_float;
+;CHECK-NEXT: ld.global.f32 r{{[0-9]+}}, [r[[R0]]];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x float]* @array_float, i32 0, i32 0
   %x = load float* %i
@@ -238,8 +238,8 @@ entry:
 
 define ptx_device double @t4_global_f64() {
 entry:
-;CHECK: mov.u32 r0, array_double;
-;CHECK-NEXT: ld.global.f64 rd0, [r0];
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_double;
+;CHECK-NEXT: ld.global.f64 rd{{[0-9]+}}, [r[[R0]]];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x double]* @array_double, i32 0, i32 0
   %x = load double* %i
@@ -248,8 +248,8 @@ entry:
 
 define ptx_device i16 @t4_const_u16() {
 entry:
-;CHECK: mov.u32 r0, array_constant_i16;
-;CHECK-NEXT: ld.const.u16 rh0, [r0];
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_constant_i16;
+;CHECK-NEXT: ld.const.u16 rh{{[0-9]+}}, [r[[R0]]];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x i16] addrspace(1)* @array_constant_i16, i32 0, i32 0
   %x = load i16 addrspace(1)* %i
@@ -258,8 +258,8 @@ entry:
 
 define ptx_device i32 @t4_const_u32() {
 entry:
-;CHECK: mov.u32 r0, array_constant_i32;
-;CHECK-NEXT: ld.const.u32 r0, [r0];
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_constant_i32;
+;CHECK-NEXT: ld.const.u32 r{{[0-9]+}}, [r[[R0]]];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x i32] addrspace(1)* @array_constant_i32, i32 0, i32 0
   %x = load i32 addrspace(1)* %i
@@ -268,8 +268,8 @@ entry:
 
 define ptx_device i64 @t4_const_u64() {
 entry:
-;CHECK: mov.u32 r0, array_constant_i64;
-;CHECK-NEXT: ld.const.u64 rd0, [r0];
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_constant_i64;
+;CHECK-NEXT: ld.const.u64 rd{{[0-9]+}}, [r[[R0]]];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x i64] addrspace(1)* @array_constant_i64, i32 0, i32 0
   %x = load i64 addrspace(1)* %i
@@ -278,8 +278,8 @@ entry:
 
 define ptx_device float @t4_const_f32() {
 entry:
-;CHECK: mov.u32 r0, array_constant_float;
-;CHECK-NEXT: ld.const.f32 r0, [r0];
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_constant_float;
+;CHECK-NEXT: ld.const.f32 r{{[0-9]+}}, [r[[R0]]];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x float] addrspace(1)* @array_constant_float, i32 0, i32 0
   %x = load float addrspace(1)* %i
@@ -288,8 +288,8 @@ entry:
 
 define ptx_device double @t4_const_f64() {
 entry:
-;CHECK: mov.u32 r0, array_constant_double;
-;CHECK-NEXT: ld.const.f64 rd0, [r0];
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_constant_double;
+;CHECK-NEXT: ld.const.f64 rd{{[0-9]+}}, [r[[R0]]];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x double] addrspace(1)* @array_constant_double, i32 0, i32 0
   %x = load double addrspace(1)* %i
@@ -298,8 +298,8 @@ entry:
 
 define ptx_device i16 @t4_local_u16() {
 entry:
-;CHECK: mov.u32 r0, array_local_i16;
-;CHECK-NEXT: ld.local.u16 rh0, [r0];
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_local_i16;
+;CHECK-NEXT: ld.local.u16 rh{{[0-9]+}}, [r[[R0]]];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x i16] addrspace(2)* @array_local_i16, i32 0, i32 0
   %x = load i16 addrspace(2)* %i
@@ -308,8 +308,8 @@ entry:
 
 define ptx_device i32 @t4_local_u32() {
 entry:
-;CHECK: mov.u32 r0, array_local_i32;
-;CHECK-NEXT: ld.local.u32 r0, [r0];
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_local_i32;
+;CHECK-NEXT: ld.local.u32 r{{[0-9]+}}, [r[[R0]]];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x i32] addrspace(2)* @array_local_i32, i32 0, i32 0
   %x = load i32 addrspace(2)* %i
@@ -318,8 +318,8 @@ entry:
 
 define ptx_device i64 @t4_local_u64() {
 entry:
-;CHECK: mov.u32 r0, array_local_i64;
-;CHECK-NEXT: ld.local.u64 rd0, [r0];
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_local_i64;
+;CHECK-NEXT: ld.local.u64 rd{{[0-9]+}}, [r[[R0]]];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x i64] addrspace(2)* @array_local_i64, i32 0, i32 0
   %x = load i64 addrspace(2)* %i
@@ -328,8 +328,8 @@ entry:
 
 define ptx_device float @t4_local_f32() {
 entry:
-;CHECK: mov.u32 r0, array_local_float;
-;CHECK-NEXT: ld.local.f32 r0, [r0];
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_local_float;
+;CHECK-NEXT: ld.local.f32 r{{[0-9]+}}, [r[[R0]]];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x float] addrspace(2)* @array_local_float, i32 0, i32 0
   %x = load float addrspace(2)* %i
@@ -338,8 +338,8 @@ entry:
 
 define ptx_device double @t4_local_f64() {
 entry:
-;CHECK: mov.u32 r0, array_local_double;
-;CHECK-NEXT: ld.local.f64 rd0, [r0];
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_local_double;
+;CHECK-NEXT: ld.local.f64 rd{{[0-9]+}}, [r[[R0]]];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x double] addrspace(2)* @array_local_double, i32 0, i32 0
   %x = load double addrspace(2)* %i
@@ -348,8 +348,8 @@ entry:
 
 define ptx_device i16 @t4_shared_u16() {
 entry:
-;CHECK: mov.u32 r0, array_shared_i16;
-;CHECK-NEXT: ld.shared.u16 rh0, [r0];
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_shared_i16;
+;CHECK-NEXT: ld.shared.u16 rh{{[0-9]+}}, [r[[R0]]];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x i16] addrspace(4)* @array_shared_i16, i32 0, i32 0
   %x = load i16 addrspace(4)* %i
@@ -358,8 +358,8 @@ entry:
 
 define ptx_device i32 @t4_shared_u32() {
 entry:
-;CHECK: mov.u32 r0, array_shared_i32;
-;CHECK-NEXT: ld.shared.u32 r0, [r0];
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_shared_i32;
+;CHECK-NEXT: ld.shared.u32 r{{[0-9]+}}, [r[[R0]]];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x i32] addrspace(4)* @array_shared_i32, i32 0, i32 0
   %x = load i32 addrspace(4)* %i
@@ -368,8 +368,8 @@ entry:
 
 define ptx_device i64 @t4_shared_u64() {
 entry:
-;CHECK: mov.u32 r0, array_shared_i64;
-;CHECK-NEXT: ld.shared.u64 rd0, [r0];
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_shared_i64;
+;CHECK-NEXT: ld.shared.u64 rd{{[0-9]+}}, [r[[R0]]];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x i64] addrspace(4)* @array_shared_i64, i32 0, i32 0
   %x = load i64 addrspace(4)* %i
@@ -378,8 +378,8 @@ entry:
 
 define ptx_device float @t4_shared_f32() {
 entry:
-;CHECK: mov.u32 r0, array_shared_float;
-;CHECK-NEXT: ld.shared.f32 r0, [r0];
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_shared_float;
+;CHECK-NEXT: ld.shared.f32 r{{[0-9]+}}, [r[[R0]]];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x float] addrspace(4)* @array_shared_float, i32 0, i32 0
   %x = load float addrspace(4)* %i
@@ -388,8 +388,8 @@ entry:
 
 define ptx_device double @t4_shared_f64() {
 entry:
-;CHECK: mov.u32 r0, array_shared_double;
-;CHECK-NEXT: ld.shared.f64 rd0, [r0];
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_shared_double;
+;CHECK-NEXT: ld.shared.f64 rd{{[0-9]+}}, [r[[R0]]];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x double] addrspace(4)* @array_shared_double, i32 0, i32 0
   %x = load double addrspace(4)* %i
@@ -398,8 +398,8 @@ entry:
 
 define ptx_device i16 @t5_u16() {
 entry:
-;CHECK: mov.u32 r0, array_i16;
-;CHECK-NEXT: ld.global.u16 rh0, [r0+2];
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_i16;
+;CHECK-NEXT: ld.global.u16 rh{{[0-9]+}}, [r[[R0]]+2];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x i16]* @array_i16, i32 0, i32 1
   %x = load i16* %i
@@ -408,8 +408,8 @@ entry:
 
 define ptx_device i32 @t5_u32() {
 entry:
-;CHECK: mov.u32 r0, array_i32;
-;CHECK-NEXT: ld.global.u32 r0, [r0+4];
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_i32;
+;CHECK-NEXT: ld.global.u32 r{{[0-9]+}}, [r[[R0]]+4];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x i32]* @array_i32, i32 0, i32 1
   %x = load i32* %i
@@ -418,8 +418,8 @@ entry:
 
 define ptx_device i64 @t5_u64() {
 entry:
-;CHECK: mov.u32 r0, array_i64;
-;CHECK-NEXT: ld.global.u64 rd0, [r0+8];
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_i64;
+;CHECK-NEXT: ld.global.u64 rd{{[0-9]+}}, [r[[R0]]+8];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x i64]* @array_i64, i32 0, i32 1
   %x = load i64* %i
@@ -428,8 +428,8 @@ entry:
 
 define ptx_device float @t5_f32() {
 entry:
-;CHECK: mov.u32 r0, array_float;
-;CHECK-NEXT: ld.global.f32 r0, [r0+4];
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_float;
+;CHECK-NEXT: ld.global.f32 r{{[0-9]+}}, [r[[R0]]+4];
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x float]* @array_float, i32 0, i32 1
   %x = load float* %i
@@ -438,8 +438,8 @@ entry:
 
 define ptx_device double @t5_f64() {
 entry:
-;CHECK: mov.u32 r0, array_double;
-;CHECK-NEXT: ld.global.f64 rd0, [r0+8];
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_double;
+;CHECK-NEXT: ld.global.f64 rd{{[0-9]+}}, [r[[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 9c532c1a3c..4611c54be8 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 r0, r1;
+; CHECK: sqrt.rn.f32 r{{[0-9]+}}, r{{[0-9]+}};
 ; 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 rd0, rd1;
+; CHECK: sqrt.rn.f64 rd{{[0-9]+}}, rd{{[0-9]+}};
 ; 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 r0, r1;
+; CHECK: sin.approx.f32 r{{[0-9]+}}, r{{[0-9]+}};
 ; 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 rd0, rd1;
+; CHECK: sin.approx.f64 rd{{[0-9]+}}, rd{{[0-9]+}};
 ; 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 r0, r1;
+; CHECK: cos.approx.f32 r{{[0-9]+}}, r{{[0-9]+}};
 ; 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 rd0, rd1;
+; CHECK: cos.approx.f64 rd{{[0-9]+}}, rd{{[0-9]+}};
 ; 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 56d3811aa7..0e4d3f9953 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 r0, r1, r2, r3;
+; CHECK: mad.rn.f32 r{{[0-9]+}}, r{{[0-9]+}}, r{{[0-9]+}}, r{{[0-9]+}};
 ; 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 rd0, rd1, rd2, rd3;
+; CHECK: mad.rn.f64 rd{{[0-9]+}}, rd{{[0-9]+}}, rd{{[0-9]+}}, rd{{[0-9]+}};
 ; 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 05ce4c0a88..cce6a5b897 100644
--- a/test/CodeGen/PTX/mov.ll
+++ b/test/CodeGen/PTX/mov.ll
@@ -1,61 +1,61 @@
 ; RUN: llc < %s -march=ptx32 | FileCheck %s
 
 define ptx_device i16 @t1_u16() {
-; CHECK: mov.u16 rh0, 0;
+; CHECK: mov.u16 rh{{[0-9]+}}, 0;
 ; CHECK: ret;
 	ret i16 0
 }
 
 define ptx_device i32 @t1_u32() {
-; CHECK: mov.u32 r0, 0;
+; CHECK: mov.u32 r{{[0-9]+}}, 0;
 ; CHECK: ret;
 	ret i32 0
 }
 
 define ptx_device i64 @t1_u64() {
-; CHECK: mov.u64 rd0, 0;
+; CHECK: mov.u64 rd{{[0-9]+}}, 0;
 ; CHECK: ret;
 	ret i64 0
 }
 
 define ptx_device float @t1_f32() {
-; CHECK: mov.f32 r0, 0F00000000;
+; CHECK: mov.f32 r{{[0-9]+}}, 0F00000000;
 ; CHECK: ret;
 	ret float 0.0
 }
 
 define ptx_device double @t1_f64() {
-; CHECK: mov.f64 rd0, 0D0000000000000000;
+; CHECK: mov.f64 rd{{[0-9]+}}, 0D0000000000000000;
 ; CHECK: ret;
 	ret double 0.0
 }
 
 define ptx_device i16 @t2_u16(i16 %x) {
-; CHECK: mov.u16 rh0, rh1;
+; CHECK: mov.u16 rh{{[0-9]+}}, rh{{[0-9]+}};
 ; CHECK: ret;
 	ret i16 %x
 }
 
 define ptx_device i32 @t2_u32(i32 %x) {
-; CHECK: mov.u32 r0, r1;
+; CHECK: mov.u32 r{{[0-9]+}}, r{{[0-9]+}};
 ; CHECK: ret;
 	ret i32 %x
 }
 
 define ptx_device i64 @t2_u64(i64 %x) {
-; CHECK: mov.u64 rd0, rd1;
+; CHECK: mov.u64 rd{{[0-9]+}}, rd{{[0-9]+}};
 ; CHECK: ret;
 	ret i64 %x
 }
 
 define ptx_device float @t3_f32(float %x) {
-; CHECK: mov.u32 r0, r1;
+; CHECK: mov.u32 r{{[0-9]+}}, r{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	ret float %x
 }
 
 define ptx_device double @t3_f64(double %x) {
-; CHECK: mov.u64 rd0, rd1;
+; CHECK: mov.u64 rd{{[0-9]+}}, rd{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	ret double %x
 }
diff --git a/test/CodeGen/PTX/mul.ll b/test/CodeGen/PTX/mul.ll
index 2093556dac..491cc747af 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.rn.f32 r0, r1, r2
+; CHECK: mul.rn.f32 r{{[0-9]+}}, r{{[0-9]+}}, r{{[0-9]+}}
 ; CHECK-NEXT: ret;
   %z = fmul float %x, %y
   ret float %z
 }
 
 define ptx_device double @t1_f64(double %x, double %y) {
-; CHECK: mul.rn.f64 rd0, rd1, rd2
+; CHECK: mul.rn.f64 rd{{[0-9]+}}, rd{{[0-9]+}}, rd{{[0-9]+}}
 ; CHECK-NEXT: ret;
   %z = fmul double %x, %y
   ret double %z
 }
 
 define ptx_device float @t2_f32(float %x) {
-; CHECK: mul.rn.f32 r0, r1, 0F40A00000;
+; CHECK: mul.rn.f32 r{{[0-9]+}}, r{{[0-9]+}}, 0F40A00000;
 ; CHECK-NEXT: ret;
   %z = fmul float %x, 5.0
   ret float %z
 }
 
 define ptx_device double @t2_f64(double %x) {
-; CHECK: mul.rn.f64 rd0, rd1, 0D4014000000000000;
+; CHECK: mul.rn.f64 rd{{[0-9]+}}, rd{{[0-9]+}}, 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 5486472099..b16556e066 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 .b32 r0) test_parameter_order (.reg .b32 r1, .reg .b32 r2, .reg .b32 r3, .reg .b32 r4)
+; CHECK: .func (.reg .b32 r{{[0-9]+}}) test_parameter_order (.reg .b32 r{{[0-9]+}}, .reg .b32 r{{[0-9]+}}, .reg .b32 r{{[0-9]+}}, .reg .b32 r{{[0-9]+}})
 define ptx_device i32 @test_parameter_order(float %a, i32 %b, i32 %c, float %d) {
-; CHECK: sub.u32 r0, r2, r3
+; CHECK: sub.u32 r{{[0-9]+}}, r{{[0-9]+}}, r{{[0-9]+}}
 	%result = sub i32 %b, %c
 	ret i32 %result
 }
diff --git a/test/CodeGen/PTX/selp.ll b/test/CodeGen/PTX/selp.ll
index 19cfa5399d..e705fbea27 100644
--- a/test/CodeGen/PTX/selp.ll
+++ b/test/CodeGen/PTX/selp.ll
@@ -1,25 +1,25 @@
 ; RUN: llc < %s -march=ptx32 | FileCheck %s
 
 define ptx_device i32 @test_selp_i32(i1 %x, i32 %y, i32 %z) {
-; CHECK: selp.u32 r0, r1, r2, p1;
+; CHECK: selp.u32 r{{[0-9]+}}, r{{[0-9]+}}, r{{[0-9]+}}, p{{[0-9]+}};
 	%a = select i1 %x, i32 %y, i32 %z
 	ret i32 %a
 }
 
 define ptx_device i64 @test_selp_i64(i1 %x, i64 %y, i64 %z) {
-; CHECK: selp.u64 rd0, rd1, rd2, p1;
+; CHECK: selp.u64 rd{{[0-9]+}}, rd{{[0-9]+}}, rd{{[0-9]+}}, p{{[0-9]+}};
 	%a = select i1 %x, i64 %y, i64 %z
 	ret i64 %a
 }
 
 define ptx_device float @test_selp_f32(i1 %x, float %y, float %z) {
-; CHECK: selp.f32 r0, r1, r2, p1;
+; CHECK: selp.f32 r{{[0-9]+}}, r{{[0-9]+}}, r{{[0-9]+}}, p{{[0-9]+}};
 	%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 rd0, rd1, rd2, p1;
+; CHECK: selp.f64 rd{{[0-9]+}}, rd{{[0-9]+}}, rd{{[0-9]+}}, p{{[0-9]+}};
 	%a = select i1 %x, double %y, double %z
 	ret double %a
 }
diff --git a/test/CodeGen/PTX/setp.ll b/test/CodeGen/PTX/setp.ll
index 3e01a75a70..b8cb364550 100644
--- a/test/CodeGen/PTX/setp.ll
+++ b/test/CodeGen/PTX/setp.ll
@@ -1,8 +1,8 @@
 ; 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;
-; CHECK-NEXT: selp.u32 r0, 1, 0, p0;
+; CHECK: setp.eq.u32 p[[P0:[0-9]+]], r{{[0-9]+}}, r{{[0-9]+}};
+; CHECK-NEXT: selp.u32 r{{[0-9]+}}, 1, 0, p[[P0]];
 ; CHECK-NEXT: ret;
 	%p = icmp eq i32 %x, %y
 	%z = zext i1 %p to i32
@@ -10,8 +10,8 @@ define ptx_device i32 @test_setp_eq_u32_rr(i32 %x, i32 %y) {
 }
 
 define ptx_device i32 @test_setp_ne_u32_rr(i32 %x, i32 %y) {
-; CHECK: setp.ne.u32 p0, r1, r2;
-; CHECK-NEXT: selp.u32 r0, 1, 0, p0;
+; CHECK: setp.ne.u32 p[[P0:[0-9]+]], r{{[0-9]+}}, r{{[0-9]+}};
+; CHECK-NEXT: selp.u32 r{{[0-9]+}}, 1, 0, p[[P0]];
 ; CHECK-NEXT: ret;
 	%p = icmp ne i32 %x, %y
 	%z = zext i1 %p to i32
@@ -19,8 +19,8 @@ define ptx_device i32 @test_setp_ne_u32_rr(i32 %x, i32 %y) {
 }
 
 define ptx_device i32 @test_setp_lt_u32_rr(i32 %x, i32 %y) {
-; CHECK: setp.lt.u32 p0, r1, r2;
-; CHECK-NEXT: selp.u32 r0, 1, 0, p0;
+; CHECK: setp.lt.u32 p[[P0:[0-9]+]], r{{[0-9]+}}, r{{[0-9]+}};
+; CHECK-NEXT: selp.u32 r{{[0-9]+}}, 1, 0, p[[P0]];
 ; CHECK-NEXT: ret;
 	%p = icmp ult i32 %x, %y
 	%z = zext i1 %p to i32
@@ -28,8 +28,8 @@ define ptx_device i32 @test_setp_lt_u32_rr(i32 %x, i32 %y) {
 }
 
 define ptx_device i32 @test_setp_le_u32_rr(i32 %x, i32 %y) {
-; CHECK: setp.le.u32 p0, r1, r2;
-; CHECK-NEXT: selp.u32 r0, 1, 0, p0;
+; CHECK: setp.le.u32 p[[P0:[0-9]+]], r{{[0-9]+}}, r{{[0-9]+}};
+; CHECK-NEXT: selp.u32 r{{[0-9]+}}, 1, 0, p[[P0]];
 ; CHECK-NEXT: ret;
 	%p = icmp ule i32 %x, %y
 	%z = zext i1 %p to i32
@@ -37,8 +37,8 @@ define ptx_device i32 @test_setp_le_u32_rr(i32 %x, i32 %y) {
 }
 
 define ptx_device i32 @test_setp_gt_u32_rr(i32 %x, i32 %y) {
-; CHECK: setp.gt.u32 p0, r1, r2;
-; CHECK-NEXT: selp.u32 r0, 1, 0, p0;
+; CHECK: setp.gt.u32 p[[P0:[0-9]+]], r{{[0-9]+}}, r{{[0-9]+}};
+; CHECK-NEXT: selp.u32 r{{[0-9]+}}, 1, 0, p[[P0]];
 ; CHECK-NEXT: ret;
 	%p = icmp ugt i32 %x, %y
 	%z = zext i1 %p to i32
@@ -46,8 +46,8 @@ define ptx_device i32 @test_setp_gt_u32_rr(i32 %x, i32 %y) {
 }
 
 define ptx_device i32 @test_setp_ge_u32_rr(i32 %x, i32 %y) {
-; CHECK: setp.ge.u32 p0, r1, r2;
-; CHECK-NEXT: selp.u32 r0, 1, 0, p0;
+; CHECK: setp.ge.u32 p[[P0:[0-9]+]], r{{[0-9]+}}, r{{[0-9]+}};
+; CHECK-NEXT: selp.u32 r{{[0-9]+}}, 1, 0, p[[P0]];
 ; CHECK-NEXT: ret;
 	%p = icmp uge i32 %x, %y
 	%z = zext i1 %p to i32
@@ -55,8 +55,8 @@ define ptx_device i32 @test_setp_ge_u32_rr(i32 %x, i32 %y) {
 }
 
 define ptx_device i32 @test_setp_lt_s32_rr(i32 %x, i32 %y) {
-; CHECK: setp.lt.s32 p0, r1, r2;
-; CHECK-NEXT: selp.u32 r0, 1, 0, p0;
+; CHECK: setp.lt.s32 p[[P0:[0-9]+]], r{{[0-9]+}}, r{{[0-9]+}};
+; CHECK-NEXT: selp.u32 r{{[0-9]+}}, 1, 0, p[[P0]];
 ; CHECK-NEXT: ret;
 	%p = icmp slt i32 %x, %y
 	%z = zext i1 %p to i32
@@ -64,8 +64,8 @@ define ptx_device i32 @test_setp_lt_s32_rr(i32 %x, i32 %y) {
 }
 
 define ptx_device i32 @test_setp_le_s32_rr(i32 %x, i32 %y) {
-; CHECK: setp.le.s32 p0, r1, r2;
-; CHECK-NEXT: selp.u32 r0, 1, 0, p0;
+; CHECK: setp.le.s32 p[[P0:[0-9]+]], r{{[0-9]+}}, r{{[0-9]+}};
+; CHECK-NEXT: selp.u32 r{{[0-9]+}}, 1, 0, p[[P0]];
 ; CHECK-NEXT: ret;
 	%p = icmp sle i32 %x, %y
 	%z = zext i1 %p to i32
@@ -73,8 +73,8 @@ define ptx_device i32 @test_setp_le_s32_rr(i32 %x, i32 %y) {
 }
 
 define ptx_device i32 @test_setp_gt_s32_rr(i32 %x, i32 %y) {
-; CHECK: setp.gt.s32 p0, r1, r2;
-; CHECK-NEXT: selp.u32 r0, 1, 0, p0;
+; CHECK: setp.gt.s32 p[[P0:[0-9]+]], r{{[0-9]+}}, r{{[0-9]+}};
+; CHECK-NEXT: selp.u32 r{{[0-9]+}}, 1, 0, p[[P0]];
 ; CHECK-NEXT: ret;
 	%p = icmp sgt i32 %x, %y
 	%z = zext i1 %p to i32
@@ -82,8 +82,8 @@ define ptx_device i32 @test_setp_gt_s32_rr(i32 %x, i32 %y) {
 }
 
 define ptx_device i32 @test_setp_ge_s32_rr(i32 %x, i32 %y) {
-; CHECK: setp.ge.s32 p0, r1, r2;
-; CHECK-NEXT: selp.u32 r0, 1, 0, p0;
+; CHECK: setp.ge.s32 p[[P0:[0-9]+]], r{{[0-9]+}}, r{{[0-9]+}};
+; CHECK-NEXT: selp.u32 r{{[0-9]+}}, 1, 0, p[[P0]];
 ; CHECK-NEXT: ret;
 	%p = icmp sge i32 %x, %y
 	%z = zext i1 %p to i32
@@ -91,8 +91,8 @@ define ptx_device i32 @test_setp_ge_s32_rr(i32 %x, i32 %y) {
 }
 
 define ptx_device i32 @test_setp_eq_u32_ri(i32 %x) {
-; CHECK: setp.eq.u32 p0, r1, 1;
-; CHECK-NEXT: selp.u32 r0, 1, 0, p0;
+; CHECK: setp.eq.u32 p[[P0:[0-9]+]], r{{[0-9]+}}, 1;
+; CHECK-NEXT: selp.u32 r{{[0-9]+}}, 1, 0, p[[P0]];
 ; CHECK-NEXT: ret;
 	%p = icmp eq i32 %x, 1
 	%z = zext i1 %p to i32
@@ -100,8 +100,8 @@ define ptx_device i32 @test_setp_eq_u32_ri(i32 %x) {
 }
 
 define ptx_device i32 @test_setp_ne_u32_ri(i32 %x) {
-; CHECK: setp.ne.u32 p0, r1, 1;
-; CHECK-NEXT: selp.u32 r0, 1, 0, p0;
+; CHECK: setp.ne.u32 p[[P0:[0-9]+]], r{{[0-9]+}}, 1;
+; CHECK-NEXT: selp.u32 r{{[0-9]+}}, 1, 0, p[[P0]];
 ; CHECK-NEXT: ret;
 	%p = icmp ne i32 %x, 1
 	%z = zext i1 %p to i32
@@ -109,8 +109,8 @@ define ptx_device i32 @test_setp_ne_u32_ri(i32 %x) {
 }
 
 define ptx_device i32 @test_setp_lt_u32_ri(i32 %x) {
-; CHECK: setp.eq.u32 p0, r1, 0;
-; CHECK-NEXT: selp.u32 r0, 1, 0, p0;
+; CHECK: setp.eq.u32 p[[P0:[0-9]+]], r{{[0-9]+}}, 0;
+; CHECK-NEXT: selp.u32 r{{[0-9]+}}, 1, 0, p[[P0]];
 ; CHECK-NEXT: ret;
 	%p = icmp ult i32 %x, 1
 	%z = zext i1 %p to i32
@@ -118,8 +118,8 @@ define ptx_device i32 @test_setp_lt_u32_ri(i32 %x) {
 }
 
 define ptx_device i32 @test_setp_le_u32_ri(i32 %x) {
-; CHECK: setp.lt.u32 p0, r1, 2;
-; CHECK-NEXT: selp.u32 r0, 1, 0, p0;
+; CHECK: setp.lt.u32 p[[P0:[0-9]+]], r{{[0-9]+}}, 2;
+; CHECK-NEXT: selp.u32 r{{[0-9]+}}, 1, 0, p[[P0]];
 ; CHECK-NEXT: ret;
 	%p = icmp ule i32 %x, 1
 	%z = zext i1 %p to i32
@@ -127,8 +127,8 @@ define ptx_device i32 @test_setp_le_u32_ri(i32 %x) {
 }
 
 define ptx_device i32 @test_setp_gt_u32_ri(i32 %x) {
-; CHECK: setp.gt.u32 p0, r1, 1;
-; CHECK-NEXT: selp.u32 r0, 1, 0, p0;
+; CHECK: setp.gt.u32 p[[P0:[0-9]+]], r{{[0-9]+}}, 1;
+; CHECK-NEXT: selp.u32 r{{[0-9]+}}, 1, 0, p[[P0]];
 ; CHECK-NEXT: ret;
 	%p = icmp ugt i32 %x, 1
 	%z = zext i1 %p to i32
@@ -136,8 +136,8 @@ define ptx_device i32 @test_setp_gt_u32_ri(i32 %x) {
 }
 
 define ptx_device i32 @test_setp_ge_u32_ri(i32 %x) {
-; CHECK: setp.ne.u32 p0, r1, 0;
-; CHECK-NEXT: selp.u32 r0, 1, 0, p0;
+; CHECK: setp.ne.u32 p[[P0:[0-9]+]], r{{[0-9]+}}, 0;
+; CHECK-NEXT: selp.u32 r{{[0-9]+}}, 1, 0, p[[P0]];
 ; CHECK-NEXT: ret;
 	%p = icmp uge i32 %x, 1
 	%z = zext i1 %p to i32
@@ -145,8 +145,8 @@ define ptx_device i32 @test_setp_ge_u32_ri(i32 %x) {
 }
 
 define ptx_device i32 @test_setp_lt_s32_ri(i32 %x) {
-; CHECK: setp.lt.s32 p0, r1, 1;
-; CHECK-NEXT: selp.u32 r0, 1, 0, p0;
+; CHECK: setp.lt.s32 p[[P0:[0-9]+]], r{{[0-9]+}}, 1;
+; CHECK-NEXT: selp.u32 r{{[0-9]+}}, 1, 0, p[[P0]];
 ; CHECK-NEXT: ret;
 	%p = icmp slt i32 %x, 1
 	%z = zext i1 %p to i32
@@ -154,8 +154,8 @@ define ptx_device i32 @test_setp_lt_s32_ri(i32 %x) {
 }
 
 define ptx_device i32 @test_setp_le_s32_ri(i32 %x) {
-; CHECK: setp.lt.s32 p0, r1, 2;
-; CHECK-NEXT: selp.u32 r0, 1, 0, p0;
+; CHECK: setp.lt.s32 p[[P0:[0-9]+]], r{{[0-9]+}}, 2;
+; CHECK-NEXT: selp.u32 r{{[0-9]+}}, 1, 0, p[[P0]];
 ; CHECK-NEXT: ret;
 	%p = icmp sle i32 %x, 1
 	%z = zext i1 %p to i32
@@ -163,8 +163,8 @@ define ptx_device i32 @test_setp_le_s32_ri(i32 %x) {
 }
 
 define ptx_device i32 @test_setp_gt_s32_ri(i32 %x) {
-; CHECK: setp.gt.s32 p0, r1, 1;
-; CHECK-NEXT: selp.u32 r0, 1, 0, p0;
+; CHECK: setp.gt.s32 p[[P0:[0-9]+]], r{{[0-9]+}}, 1;
+; CHECK-NEXT: selp.u32 r{{[0-9]+}}, 1, 0, p[[P0]];
 ; CHECK-NEXT: ret;
 	%p = icmp sgt i32 %x, 1
 	%z = zext i1 %p to i32
@@ -172,8 +172,8 @@ define ptx_device i32 @test_setp_gt_s32_ri(i32 %x) {
 }
 
 define ptx_device i32 @test_setp_ge_s32_ri(i32 %x) {
-; CHECK: setp.gt.s32 p0, r1, 0;
-; CHECK-NEXT: selp.u32 r0, 1, 0, p0;
+; CHECK: setp.gt.s32 p[[P0:[0-9]+]], r{{[0-9]+}}, 0;
+; CHECK-NEXT: selp.u32 r{{[0-9]+}}, 1, 0, p[[P0]];
 ; CHECK-NEXT: ret;
 	%p = icmp sge i32 %x, 1
 	%z = zext i1 %p to i32
@@ -181,9 +181,9 @@ define ptx_device i32 @test_setp_ge_s32_ri(i32 %x) {
 }
 
 define ptx_device i32 @test_setp_4_op_format_1(i32 %x, i32 %y, i32 %u, i32 %v) {
-; CHECK: setp.gt.u32 p0, r3, r4;
-; CHECK-NEXT: setp.eq.and.u32 p0, r1, r2, p0;
-; CHECK-NEXT: selp.u32 r0, 1, 0, p0;
+; CHECK: setp.gt.u32 p[[P0:[0-9]+]], r{{[0-9]+}}, r{{[0-9]+}};
+; CHECK-NEXT: setp.eq.and.u32 p[[P0]], r{{[0-9]+}}, r{{[0-9]+}}, p[[P0]];
+; CHECK-NEXT: selp.u32 r{{[0-9]+}}, 1, 0, p[[P0]];
 ; CHECK-NEXT: ret;
 	%c = icmp eq i32 %x, %y
 	%d = icmp ugt i32 %u, %v
@@ -193,9 +193,9 @@ define ptx_device i32 @test_setp_4_op_format_1(i32 %x, i32 %y, i32 %u, i32 %v) {
 }
 
 define ptx_device i32 @test_setp_4_op_format_2(i32 %x, i32 %y, i32 %w) {
-; CHECK: setp.gt.b32 p0, r3, 0;
-; CHECK-NEXT: setp.eq.and.u32 p0, r1, r2, !p0;
-; CHECK-NEXT: selp.u32 r0, 1, 0, p0;
+; CHECK: setp.gt.b32 p[[P0:[0-9]+]], r{{[0-9]+}}, 0;
+; CHECK-NEXT: setp.eq.and.u32 p[[P0]], r{{[0-9]+}}, r{{[0-9]+}}, !p[[P0]];
+; CHECK-NEXT: selp.u32 r{{[0-9]+}}, 1, 0, p[[P0]];
 ; CHECK-NEXT: ret;
 	%c = trunc i32 %w to i1
 	%d = icmp eq i32 %x, %y
diff --git a/test/CodeGen/PTX/shl.ll b/test/CodeGen/PTX/shl.ll
index 6e72c92213..b3818e1e76 100644
--- a/test/CodeGen/PTX/shl.ll
+++ b/test/CodeGen/PTX/shl.ll
@@ -1,21 +1,21 @@
 ; RUN: llc < %s -march=ptx32 | FileCheck %s
 
 define ptx_device i32 @t1(i32 %x, i32 %y) {
-; CHECK: shl.b32 r0, r1, r2
+; CHECK: shl.b32 r{{[0-9]+}}, r{{[0-9]+}}, r{{[0-9]+}}
 	%z = shl i32 %x, %y
 ; CHECK: ret;
 	ret i32 %z
 }
 
 define ptx_device i32 @t2(i32 %x) {
-; CHECK: shl.b32 r0, r1, 3
+; CHECK: shl.b32 r{{[0-9]+}}, r{{[0-9]+}}, 3
 	%z = shl i32 %x, 3
 ; CHECK: ret;
 	ret i32 %z
 }
 
 define ptx_device i32 @t3(i32 %x) {
-; CHECK: shl.b32 r0, 3, r1
+; CHECK: shl.b32 r{{[0-9]+}}, 3, r{{[0-9]+}}
 	%z = shl i32 3, %x
 ; CHECK: ret;
 	ret i32 %z
diff --git a/test/CodeGen/PTX/shr.ll b/test/CodeGen/PTX/shr.ll
index 8693e0ecf4..cb57546dca 100644
--- a/test/CodeGen/PTX/shr.ll
+++ b/test/CodeGen/PTX/shr.ll
@@ -1,42 +1,42 @@
 ; RUN: llc < %s -march=ptx32 | FileCheck %s
 
 define ptx_device i32 @t1(i32 %x, i32 %y) {
-; CHECK: shr.u32 r0, r1, r2
+; CHECK: shr.u32 r{{[0-9]+}}, r{{[0-9]+}}, r{{[0-9]+}}
 	%z = lshr i32 %x, %y
 ; CHECK: ret;
 	ret i32 %z
 }
 
 define ptx_device i32 @t2(i32 %x) {
-; CHECK: shr.u32 r0, r1, 3
+; CHECK: shr.u32 r{{[0-9]+}}, r{{[0-9]+}}, 3
 	%z = lshr i32 %x, 3
 ; CHECK: ret;
 	ret i32 %z
 }
 
 define ptx_device i32 @t3(i32 %x) {
-; CHECK: shr.u32 r0, 3, r1
+; CHECK: shr.u32 r{{[0-9]+}}, 3, r{{[0-9]+}}
 	%z = lshr i32 3, %x
 ; CHECK: ret;
 	ret i32 %z
 }
 
 define ptx_device i32 @t4(i32 %x, i32 %y) {
-; CHECK: shr.s32 r0, r1, r2
+; CHECK: shr.s32 r{{[0-9]+}}, r{{[0-9]+}}, r{{[0-9]+}}
 	%z = ashr i32 %x, %y
 ; CHECK: ret;
 	ret i32 %z
 }
 
 define ptx_device i32 @t5(i32 %x) {
-; CHECK: shr.s32 r0, r1, 3
+; CHECK: shr.s32 r{{[0-9]+}}, r{{[0-9]+}}, 3
 	%z = ashr i32 %x, 3
 ; CHECK: ret;
 	ret i32 %z
 }
 
 define ptx_device i32 @t6(i32 %x) {
-; CHECK: shr.s32 r0, -3, r1
+; CHECK: shr.s32 r{{[0-9]+}}, -3, r{{[0-9]+}}
 	%z = ashr i32 -3, %x
 ; CHECK: ret;
 	ret i32 %z
diff --git a/test/CodeGen/PTX/st.ll b/test/CodeGen/PTX/st.ll
index 612967ac83..b08528e1c3 100644
--- a/test/CodeGen/PTX/st.ll
+++ b/test/CodeGen/PTX/st.ll
@@ -63,7 +63,7 @@
 
 define ptx_device void @t1_u16(i16* %p, i16 %x) {
 entry:
-;CHECK: st.global.u16 [r1], rh1;
+;CHECK: st.global.u16 [r{{[0-9]+}}], rh{{[0-9]+}};
 ;CHECK-NEXT: ret;
   store i16 %x, i16* %p
   ret void
@@ -71,7 +71,7 @@ entry:
 
 define ptx_device void @t1_u32(i32* %p, i32 %x) {
 entry:
-;CHECK: st.global.u32 [r1], r2;
+;CHECK: st.global.u32 [r{{[0-9]+}}], r{{[0-9]+}};
 ;CHECK-NEXT: ret;
   store i32 %x, i32* %p
   ret void
@@ -79,7 +79,7 @@ entry:
 
 define ptx_device void @t1_u64(i64* %p, i64 %x) {
 entry:
-;CHECK: st.global.u64 [r1], rd1;
+;CHECK: st.global.u64 [r{{[0-9]+}}], rd{{[0-9]+}};
 ;CHECK-NEXT: ret;
   store i64 %x, i64* %p
   ret void
@@ -87,7 +87,7 @@ entry:
 
 define ptx_device void @t1_f32(float* %p, float %x) {
 entry:
-;CHECK: st.global.f32 [r1], r2;
+;CHECK: st.global.f32 [r{{[0-9]+}}], r{{[0-9]+}};
 ;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], rd1;
+;CHECK: st.global.f64 [r{{[0-9]+}}], rd{{[0-9]+}};
 ;CHECK-NEXT: ret;
   store double %x, double* %p
   ret void
@@ -103,7 +103,7 @@ entry:
 
 define ptx_device void @t2_u16(i16* %p, i16 %x) {
 entry:
-;CHECK: st.global.u16 [r1+2], rh1;
+;CHECK: st.global.u16 [r{{[0-9]+}}+2], rh{{[0-9]+}};
 ;CHECK-NEXT: ret;
   %i = getelementptr i16* %p, i32 1
   store i16 %x, i16* %i
@@ -112,7 +112,7 @@ entry:
 
 define ptx_device void @t2_u32(i32* %p, i32 %x) {
 entry:
-;CHECK: st.global.u32 [r1+4], r2;
+;CHECK: st.global.u32 [r{{[0-9]+}}+4], r{{[0-9]+}};
 ;CHECK-NEXT: ret;
   %i = getelementptr i32* %p, i32 1
   store i32 %x, i32* %i
@@ -121,7 +121,7 @@ entry:
 
 define ptx_device void @t2_u64(i64* %p, i64 %x) {
 entry:
-;CHECK: st.global.u64 [r1+8], rd1;
+;CHECK: st.global.u64 [r{{[0-9]+}}+8], rd{{[0-9]+}};
 ;CHECK-NEXT: ret;
   %i = getelementptr i64* %p, i32 1
   store i64 %x, i64* %i
@@ -130,7 +130,7 @@ entry:
 
 define ptx_device void @t2_f32(float* %p, float %x) {
 entry:
-;CHECK: st.global.f32 [r1+4], r2;
+;CHECK: st.global.f32 [r{{[0-9]+}}+4], r{{[0-9]+}};
 ;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], rd1;
+;CHECK: st.global.f64 [r{{[0-9]+}}+8], rd{{[0-9]+}};
 ;CHECK-NEXT: ret;
   %i = getelementptr double* %p, i32 1
   store double %x, double* %i
@@ -148,9 +148,9 @@ entry:
 
 define ptx_device void @t3_u16(i16* %p, i32 %q, i16 %x) {
 entry:
-;CHECK: shl.b32 r0, r2, 1;
-;CHECK-NEXT: add.u32 r0, r1, r0;
-;CHECK-NEXT: st.global.u16 [r0], rh1;
+;CHECK: shl.b32 r[[R0:[0-9]+]], r{{[0-9]+}}, 1;
+;CHECK-NEXT: add.u32 r[[R0]], r{{[0-9]+}}, r[[R0]];
+;CHECK-NEXT: st.global.u16 [r[[R0]]], rh{{[0-9]+}};
 ;CHECK-NEXT: ret;
   %i = getelementptr i16* %p, i32 %q
   store i16 %x, i16* %i
@@ -159,9 +159,9 @@ entry:
 
 define ptx_device void @t3_u32(i32* %p, i32 %q, i32 %x) {
 entry:
-;CHECK: shl.b32 r0, r2, 2;
-;CHECK-NEXT: add.u32 r0, r1, r0;
-;CHECK-NEXT: st.global.u32 [r0], r3;
+;CHECK: shl.b32 r[[R0:[0-9]+]], r{{[0-9]+}}, 2;
+;CHECK-NEXT: add.u32 r[[R0]], r{{[0-9]+}}, r[[R0]];
+;CHECK-NEXT: st.global.u32 [r[[R0]]], r{{[0-9]+}};
 ;CHECK-NEXT: ret;
   %i = getelementptr i32* %p, i32 %q
   store i32 %x, i32* %i
@@ -170,9 +170,9 @@ entry:
 
 define ptx_device void @t3_u64(i64* %p, i32 %q, i64 %x) {
 entry:
-;CHECK: shl.b32 r0, r2, 3;
-;CHECK-NEXT: add.u32 r0, r1, r0;
-;CHECK-NEXT: st.global.u64 [r0], rd1;
+;CHECK: shl.b32 r[[R0:[0-9]+]], r{{[0-9]+}}, 3;
+;CHECK-NEXT: add.u32 r[[R0]], r{{[0-9]+}}, r[[R0]];
+;CHECK-NEXT: st.global.u64 [r[[R0]]], rd{{[0-9]+}};
 ;CHECK-NEXT: ret;
   %i = getelementptr i64* %p, i32 %q
   store i64 %x, i64* %i
@@ -181,9 +181,9 @@ entry:
 
 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], r3;
+;CHECK: shl.b32 r[[R0:[0-9]+]], r{{[0-9]+}}, 2;
+;CHECK-NEXT: add.u32 r[[R0]], r{{[0-9]+}}, r[[R0]];
+;CHECK-NEXT: st.global.f32 [r[[R0]]], r{{[0-9]+}};
 ;CHECK-NEXT: ret;
   %i = getelementptr float* %p, i32 %q
   store float %x, float* %i
@@ -192,9 +192,9 @@ entry:
 
 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], rd1;
+;CHECK: shl.b32 r[[R0:[0-9]+]], r{{[0-9]+}}, 3;
+;CHECK-NEXT: add.u32 r[[R0]], r{{[0-9]+}}, r[[R0]];
+;CHECK-NEXT: st.global.f64 [r[[R0]]], rd{{[0-9]+}};
 ;CHECK-NEXT: ret;
   %i = getelementptr double* %p, i32 %q
   store double %x, double* %i
@@ -203,8 +203,8 @@ entry:
 
 define ptx_device void @t4_global_u16(i16 %x) {
 entry:
-;CHECK: mov.u32 r0, array_i16;
-;CHECK-NEXT: st.global.u16 [r0], rh1;
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_i16;
+;CHECK-NEXT: st.global.u16 [r[[R0]]], rh{{[0-9]+}};
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x i16]* @array_i16, i16 0, i16 0
   store i16 %x, i16* %i
@@ -213,8 +213,8 @@ entry:
 
 define ptx_device void @t4_global_u32(i32 %x) {
 entry:
-;CHECK: mov.u32 r0, array_i32;
-;CHECK-NEXT: st.global.u32 [r0], r1;
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_i32;
+;CHECK-NEXT: st.global.u32 [r[[R0]]], r{{[0-9]+}};
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x i32]* @array_i32, i32 0, i32 0
   store i32 %x, i32* %i
@@ -223,8 +223,8 @@ entry:
 
 define ptx_device void @t4_global_u64(i64 %x) {
 entry:
-;CHECK: mov.u32 r0, array_i64;
-;CHECK-NEXT: st.global.u64 [r0], rd1;
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_i64;
+;CHECK-NEXT: st.global.u64 [r[[R0]]], rd{{[0-9]+}};
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x i64]* @array_i64, i32 0, i32 0
   store i64 %x, i64* %i
@@ -233,8 +233,8 @@ entry:
 
 define ptx_device void @t4_global_f32(float %x) {
 entry:
-;CHECK: mov.u32 r0, array_float;
-;CHECK-NEXT: st.global.f32 [r0], r1;
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_float;
+;CHECK-NEXT: st.global.f32 [r[[R0]]], r{{[0-9]+}};
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x float]* @array_float, i32 0, i32 0
   store float %x, float* %i
@@ -243,8 +243,8 @@ entry:
 
 define ptx_device void @t4_global_f64(double %x) {
 entry:
-;CHECK: mov.u32 r0, array_double;
-;CHECK-NEXT: st.global.f64 [r0], rd1;
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_double;
+;CHECK-NEXT: st.global.f64 [r[[R0]]], rd{{[0-9]+}};
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x double]* @array_double, i32 0, i32 0
   store double %x, double* %i
@@ -253,8 +253,8 @@ entry:
 
 define ptx_device void @t4_local_u16(i16 %x) {
 entry:
-;CHECK: mov.u32 r0, array_local_i16;
-;CHECK-NEXT: st.local.u16 [r0], rh1;
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_local_i16;
+;CHECK-NEXT: st.local.u16 [r[[R0]]], rh{{[0-9]+}};
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x i16] addrspace(2)* @array_local_i16, i32 0, i32 0
   store i16 %x, i16 addrspace(2)* %i
@@ -263,8 +263,8 @@ entry:
 
 define ptx_device void @t4_local_u32(i32 %x) {
 entry:
-;CHECK: mov.u32 r0, array_local_i32;
-;CHECK-NEXT: st.local.u32 [r0], r1;
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_local_i32;
+;CHECK-NEXT: st.local.u32 [r[[R0]]], r{{[0-9]+}};
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x i32] addrspace(2)* @array_local_i32, i32 0, i32 0
   store i32 %x, i32 addrspace(2)* %i
@@ -273,8 +273,8 @@ entry:
 
 define ptx_device void @t4_local_u64(i64 %x) {
 entry:
-;CHECK: mov.u32 r0, array_local_i64;
-;CHECK-NEXT: st.local.u64 [r0], rd1;
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_local_i64;
+;CHECK-NEXT: st.local.u64 [r[[R0]]], rd{{[0-9]+}};
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x i64] addrspace(2)* @array_local_i64, i32 0, i32 0
   store i64 %x, i64 addrspace(2)* %i
@@ -283,8 +283,8 @@ entry:
 
 define ptx_device void @t4_local_f32(float %x) {
 entry:
-;CHECK: mov.u32 r0, array_local_float;
-;CHECK-NEXT: st.local.f32 [r0], r1;
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_local_float;
+;CHECK-NEXT: st.local.f32 [r[[R0]]], r{{[0-9]+}};
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x float] addrspace(2)* @array_local_float, i32 0, i32 0
   store float %x, float addrspace(2)* %i
@@ -293,8 +293,8 @@ entry:
 
 define ptx_device void @t4_local_f64(double %x) {
 entry:
-;CHECK: mov.u32 r0, array_local_double;
-;CHECK-NEXT: st.local.f64 [r0], rd1;
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_local_double;
+;CHECK-NEXT: st.local.f64 [r[[R0]]], rd{{[0-9]+}};
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x double] addrspace(2)* @array_local_double, i32 0, i32 0
   store double %x, double addrspace(2)* %i
@@ -303,8 +303,8 @@ entry:
 
 define ptx_device void @t4_shared_u16(i16 %x) {
 entry:
-;CHECK: mov.u32 r0, array_shared_i16;
-;CHECK-NEXT: st.shared.u16 [r0], rh1;
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_shared_i16;
+;CHECK-NEXT: st.shared.u16 [r[[R0]]], rh{{[0-9]+}};
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x i16] addrspace(4)* @array_shared_i16, i32 0, i32 0
   store i16 %x, i16 addrspace(4)* %i
@@ -313,8 +313,8 @@ entry:
 
 define ptx_device void @t4_shared_u32(i32 %x) {
 entry:
-;CHECK: mov.u32 r0, array_shared_i32;
-;CHECK-NEXT: st.shared.u32 [r0], r1;
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_shared_i32;
+;CHECK-NEXT: st.shared.u32 [r[[R0]]], r{{[0-9]+}};
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x i32] addrspace(4)* @array_shared_i32, i32 0, i32 0
   store i32 %x, i32 addrspace(4)* %i
@@ -323,8 +323,8 @@ entry:
 
 define ptx_device void @t4_shared_u64(i64 %x) {
 entry:
-;CHECK: mov.u32 r0, array_shared_i64;
-;CHECK-NEXT: st.shared.u64 [r0], rd1;
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_shared_i64;
+;CHECK-NEXT: st.shared.u64 [r[[R0]]], rd{{[0-9]+}};
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x i64] addrspace(4)* @array_shared_i64, i32 0, i32 0
   store i64 %x, i64 addrspace(4)* %i
@@ -333,8 +333,8 @@ entry:
 
 define ptx_device void @t4_shared_f32(float %x) {
 entry:
-;CHECK: mov.u32 r0, array_shared_float;
-;CHECK-NEXT: st.shared.f32 [r0], r1;
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_shared_float;
+;CHECK-NEXT: st.shared.f32 [r[[R0]]], r{{[0-9]+}};
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x float] addrspace(4)* @array_shared_float, i32 0, i32 0
   store float %x, float addrspace(4)* %i
@@ -343,8 +343,8 @@ entry:
 
 define ptx_device void @t4_shared_f64(double %x) {
 entry:
-;CHECK: mov.u32 r0, array_shared_double;
-;CHECK-NEXT: st.shared.f64 [r0], rd1;
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_shared_double;
+;CHECK-NEXT: st.shared.f64 [r[[R0]]], rd{{[0-9]+}};
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x double] addrspace(4)* @array_shared_double, i32 0, i32 0
   store double %x, double addrspace(4)* %i
@@ -353,8 +353,8 @@ entry:
 
 define ptx_device void @t5_u16(i16 %x) {
 entry:
-;CHECK: mov.u32 r0, array_i16;
-;CHECK-NEXT: st.global.u16 [r0+2], rh1;
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_i16;
+;CHECK-NEXT: st.global.u16 [r[[R0]]+2], rh{{[0-9]+}};
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x i16]* @array_i16, i32 0, i32 1
   store i16 %x, i16* %i
@@ -363,8 +363,8 @@ entry:
 
 define ptx_device void @t5_u32(i32 %x) {
 entry:
-;CHECK: mov.u32 r0, array_i32;
-;CHECK-NEXT: st.global.u32 [r0+4], r1;
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_i32;
+;CHECK-NEXT: st.global.u32 [r[[R0]]+4], r{{[0-9]+}};
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x i32]* @array_i32, i32 0, i32 1
   store i32 %x, i32* %i
@@ -373,8 +373,8 @@ entry:
 
 define ptx_device void @t5_u64(i64 %x) {
 entry:
-;CHECK: mov.u32 r0, array_i64;
-;CHECK-NEXT: st.global.u64 [r0+8], rd1;
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_i64;
+;CHECK-NEXT: st.global.u64 [r[[R0]]+8], rd{{[0-9]+}};
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x i64]* @array_i64, i32 0, i32 1
   store i64 %x, i64* %i
@@ -383,8 +383,8 @@ entry:
 
 define ptx_device void @t5_f32(float %x) {
 entry:
-;CHECK: mov.u32 r0, array_float;
-;CHECK-NEXT: st.global.f32 [r0+4], r1;
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_float;
+;CHECK-NEXT: st.global.f32 [r[[R0]]+4], r{{[0-9]+}};
 ;CHECK-NEXT: ret;
   %i = getelementptr [10 x float]* @array_float, i32 0, i32 1
   store float %x, float* %i
@@ -393,8 +393,8 @@ entry:
 
 define ptx_device void @t5_f64(double %x) {
 entry:
-;CHECK: mov.u32 r0, array_double;
-;CHECK-NEXT: st.global.f64 [r0+8], rd1;
+;CHECK: mov.u32 r[[R0:[0-9]+]], array_double;
+;CHECK-NEXT: st.global.f64 [r[[R0]]+8], rd{{[0-9]+}};
 ;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 4d552801a4..acef3961bf 100644
--- a/test/CodeGen/PTX/sub.ll
+++ b/test/CodeGen/PTX/sub.ll
@@ -1,70 +1,70 @@
 ; RUN: llc < %s -march=ptx32 | FileCheck %s
 
 define ptx_device i16 @t1_u16(i16 %x, i16 %y) {
-; CHECK: sub.u16 rh0, rh1, rh2;
+; CHECK: sub.u16 rh{{[0-9]+}}, rh{{[0-9]+}}, rh{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%z = sub i16 %x, %y
 	ret i16 %z
 }
 
 define ptx_device i32 @t1_u32(i32 %x, i32 %y) {
-; CHECK: sub.u32 r0, r1, r2;
+; CHECK: sub.u32 r{{[0-9]+}}, r{{[0-9]+}}, r{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%z = sub i32 %x, %y
 	ret i32 %z
 }
 
 define ptx_device i64 @t1_u64(i64 %x, i64 %y) {
-; CHECK: sub.u64 rd0, rd1, rd2;
+; CHECK: sub.u64 rd{{[0-9]+}}, rd{{[0-9]+}}, rd{{[0-9]+}};
 ; CHECK-NEXT: ret;
 	%z = sub i64 %x, %y
 	ret i64 %z
 }
 
 define ptx_device float @t1_f32(float %x, float %y) {
-; CHECK: sub.rn.f32 r0, r1, r2
+; CHECK: sub.rn.f32 r{{[0-9]+}}, r{{[0-9]+}}, r{{[0-9]+}}
 ; CHECK-NEXT: ret;
   %z = fsub float %x, %y
   ret float %z
 }
 
 define ptx_device double @t1_f64(double %x, double %y) {
-; CHECK: sub.rn.f64 rd0, rd1, rd2
+; CHECK: sub.rn.f64 rd{{[0-9]+}}, rd{{[0-9]+}}, rd{{[0-9]+}}
 ; CHECK-NEXT: ret;
   %z = fsub double %x, %y
   ret double %z
 }
 
 define ptx_device i16 @t2_u16(i16 %x) {
-; CHECK: add.u16 rh0, rh1, -1;
+; CHECK: add.u16 rh{{[0-9]+}}, rh{{[0-9]+}}, -1;
 ; CHECK-NEXT: ret;
 	%z = sub i16 %x, 1
 	ret i16 %z
 }
 
 define ptx_device i32 @t2_u32(i32 %x) {
-; CHECK: add.u32 r0, r1, -1;
+; CHECK: add.u32 r{{[0-9]+}}, r{{[0-9]+}}, -1;
 ; CHECK-NEXT: ret;
 	%z = sub i32 %x, 1
 	ret i32 %z
 }
 
 define ptx_device i64 @t2_u64(i64 %x) {
-; CHECK: add.u64 rd0, rd1, -1;
+; CHECK: add.u64 rd{{[0-9]+}}, rd{{[0-9]+}}, -1;
 ; CHECK-NEXT: ret;
 	%z = sub i64 %x, 1
 	ret i64 %z
 }
 
 define ptx_device float @t2_f32(float %x) {
-; CHECK: add.rn.f32 r0, r1, 0FBF800000;
+; CHECK: add.rn.f32 r{{[0-9]+}}, r{{[0-9]+}}, 0FBF800000;
 ; CHECK-NEXT: ret;
   %z = fsub float %x, 1.0
   ret float %z
 }
 
 define ptx_device double @t2_f64(double %x) {
-; CHECK: add.rn.f64 rd0, rd1, 0DBFF0000000000000;
+; CHECK: add.rn.f64 rd{{[0-9]+}}, rd{{[0-9]+}}, 0DBFF0000000000000;
 ; CHECK-NEXT: ret;
   %z = fsub double %x, 1.0
   ret double %z