diff options
Diffstat (limited to 'test')
30 files changed, 0 insertions, 2186 deletions
diff --git a/test/CodeGen/PTX/20110926-sitofp.ll b/test/CodeGen/PTX/20110926-sitofp.ll deleted file mode 100644 index 38d35c5ff4..0000000000 --- a/test/CodeGen/PTX/20110926-sitofp.ll +++ /dev/null @@ -1,24 +0,0 @@ -; RUN: llc < %s -march=ptx32 | FileCheck %s - -@A = common global [1536 x [1536 x float]] zeroinitializer, align 4 -@B = common global [1536 x [1536 x float]] zeroinitializer, align 4 - -define internal ptx_device void @init_array(i32 %x, i32 %y) { - %arrayidx103 = getelementptr [1536 x [1536 x float]]* @A, i32 0, i32 %x, i32 %y - %arrayidx224 = getelementptr [1536 x [1536 x float]]* @B, i32 0, i32 %x, i32 %y - %mul5 = mul i32 %x, %y - %rem = srem i32 %mul5, 1024 - %add = add nsw i32 %rem, 1 -; CHECK: cvt.rn.f64.s32 %fd{{[0-9]+}}, %r{{[0-9]+}} - %conv = sitofp i32 %add to double - %div = fmul double %conv, 5.000000e-01 - %conv7 = fptrunc double %div to float - store float %conv7, float* %arrayidx103, align 4 - %rem14 = srem i32 %mul5, 1024 - %add15 = add nsw i32 %rem14, 1 - %conv16 = sitofp i32 %add15 to double - %div17 = fmul double %conv16, 5.000000e-01 - %conv18 = fptrunc double %div17 to float - store float %conv18, float* %arrayidx224, align 4 - ret void -} diff --git a/test/CodeGen/PTX/add.ll b/test/CodeGen/PTX/add.ll deleted file mode 100644 index 8b10d11cba..0000000000 --- a/test/CodeGen/PTX/add.ll +++ /dev/null @@ -1,71 +0,0 @@ -; RUN: llc < %s -march=ptx32 | FileCheck %s - -define ptx_device i16 @t1_u16(i16 %x, i16 %y) { -; CHECK: add.u16 %ret{{[0-9]+}}, %rh{{[0-9]+}}, %rh{{[0-9]+}}; -; CHECK: ret; - %z = add i16 %x, %y - ret i16 %z -} - -define ptx_device i32 @t1_u32(i32 %x, i32 %y) { -; CHECK: add.u32 %ret{{[0-9]+}}, %r{{[0-9]+}}, %r{{[0-9]+}}; -; CHECK: ret; - %z = add i32 %x, %y - ret i32 %z -} - -define ptx_device i64 @t1_u64(i64 %x, i64 %y) { -; CHECK: add.u64 %ret{{[0-9]+}}, %rd{{[0-9]+}}, %rd{{[0-9]+}}; -; CHECK: ret; - %z = add i64 %x, %y - ret i64 %z -} - -define ptx_device float @t1_f32(float %x, float %y) { -; CHECK: add.rn.f32 %ret{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}} -; CHECK: ret; - %z = fadd float %x, %y - ret float %z -} - -define ptx_device double @t1_f64(double %x, double %y) { -; CHECK: add.rn.f64 %ret{{[0-9]+}}, %fd{{[0-9]+}}, %fd{{[0-9]+}} -; CHECK: ret; - %z = fadd double %x, %y - ret double %z -} - -define ptx_device i16 @t2_u16(i16 %x) { -; CHECK: add.u16 %ret{{[0-9]+}}, %rh{{[0-9]+}}, 1; -; CHECK: ret; - %z = add i16 %x, 1 - ret i16 %z -} - -define ptx_device i32 @t2_u32(i32 %x) { -; CHECK: add.u32 %ret{{[0-9]+}}, %r{{[0-9]+}}, 1; -; CHECK: ret; - %z = add i32 %x, 1 - ret i32 %z -} - -define ptx_device i64 @t2_u64(i64 %x) { -; CHECK: add.u64 %ret{{[0-9]+}}, %rd{{[0-9]+}}, 1; -; CHECK: ret; - %z = add i64 %x, 1 - ret i64 %z -} - -define ptx_device float @t2_f32(float %x) { -; CHECK: add.rn.f32 %ret{{[0-9]+}}, %f{{[0-9]+}}, 0D3FF0000000000000; -; CHECK: ret; - %z = fadd float %x, 1.0 - ret float %z -} - -define ptx_device double @t2_f64(double %x) { -; CHECK: add.rn.f64 %ret{{[0-9]+}}, %fd{{[0-9]+}}, 0D3FF0000000000000; -; CHECK: ret; - %z = fadd double %x, 1.0 - ret double %z -} diff --git a/test/CodeGen/PTX/aggregates.ll b/test/CodeGen/PTX/aggregates.ll deleted file mode 100644 index 3fc0c40881..0000000000 --- a/test/CodeGen/PTX/aggregates.ll +++ /dev/null @@ -1,24 +0,0 @@ -; RUN: llc < %s -march=ptx32 -mattr=sm20 | FileCheck %s -; XFAIL: * - -%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 deleted file mode 100644 index 1403a23d1d..0000000000 --- a/test/CodeGen/PTX/bitwise.ll +++ /dev/null @@ -1,24 +0,0 @@ -; RUN: llc < %s -march=ptx32 | FileCheck %s - -; preds - -define ptx_device i32 @t1_and_preds(i1 %x, i1 %y) { -; 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 %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 %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 deleted file mode 100644 index 464c29cca8..0000000000 --- a/test/CodeGen/PTX/bra.ll +++ /dev/null @@ -1,24 +0,0 @@ -; RUN: llc < %s -march=ptx32 | FileCheck %s - -define ptx_device void @test_bra_direct() { -; CHECK: bra $L__BB0_1; -entry: - br label %loop -loop: - br label %loop -} - -define ptx_device i32 @test_bra_cond_direct(i32 %x, i32 %y) { -entry: -; 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 %ret{{[0-9]+}}, %r[[R0]] - ret i32 %x -clause.else: -; CHECK: mov.u32 %ret{{[0-9]+}}, %r[[R1]] - ret i32 %y -} diff --git a/test/CodeGen/PTX/cvt.ll b/test/CodeGen/PTX/cvt.ll deleted file mode 100644 index f55070af22..0000000000 --- a/test/CodeGen/PTX/cvt.ll +++ /dev/null @@ -1,290 +0,0 @@ -; RUN: llc < %s -march=ptx32 | FileCheck %s - -; preds -; (note: we convert back to i32 to return) - -define ptx_device i32 @cvt_pred_i16(i16 %x, i1 %y) { -; CHECK: setp.gt.u16 %p[[P0:[0-9]+]], %rh{{[0-9]+}}, 0 -; CHECK: and.pred %p2, %p[[P0:[0-9]+]], %p{{[0-9]+}}; -; CHECK: selp.u32 %ret{{[0-9]+}}, 1, 0, %p[[P0:[0-9]+]]; -; CHECK: ret; - %a = trunc i16 %x to i1 - %b = and i1 %a, %y - %c = zext i1 %b to i32 - ret i32 %c -} - -define ptx_device i32 @cvt_pred_i32(i32 %x, i1 %y) { -; CHECK: setp.gt.u32 %p[[P0:[0-9]+]], %r{{[0-9]+}}, 0 -; CHECK: and.pred %p2, %p[[P0:[0-9]+]], %p{{[0-9]+}}; -; CHECK: selp.u32 %ret{{[0-9]+}}, 1, 0, %p[[P0:[0-9]+]]; -; CHECK: ret; - %a = trunc i32 %x to i1 - %b = and i1 %a, %y - %c = zext i1 %b to i32 - ret i32 %c -} - -define ptx_device i32 @cvt_pred_i64(i64 %x, i1 %y) { -; CHECK: setp.gt.u64 %p[[P0:[0-9]+]], %rd{{[0-9]+}}, 0 -; CHECK: and.pred %p2, %p[[P0:[0-9]+]], %p{{[0-9]+}}; -; CHECK: selp.u32 %ret{{[0-9]+}}, 1, 0, %p[[P0:[0-9]+]]; -; CHECK: ret; - %a = trunc i64 %x to i1 - %b = and i1 %a, %y - %c = zext i1 %b to i32 - ret i32 %c -} - -define ptx_device i32 @cvt_pred_f32(float %x, i1 %y) { -; CHECK: setp.gt.u32 %p[[P0:[0-9]+]], %r{{[0-9]+}}, 0 -; CHECK: and.pred %p2, %p[[P0:[0-9]+]], %p{{[0-9]+}}; -; CHECK: selp.u32 %ret{{[0-9]+}}, 1, 0, %p[[P0:[0-9]+]]; -; CHECK: ret; - %a = fptoui float %x to i1 - %b = and i1 %a, %y - %c = zext i1 %b to i32 - ret i32 %c -} - -define ptx_device i32 @cvt_pred_f64(double %x, i1 %y) { -; CHECK: setp.gt.u64 %p[[P0:[0-9]+]], %rd{{[0-9]+}}, 0 -; CHECK: and.pred %p2, %p[[P0:[0-9]+]], %p{{[0-9]+}}; -; CHECK: selp.u32 %ret{{[0-9]+}}, 1, 0, %p[[P0:[0-9]+]]; -; CHECK: ret; - %a = fptoui double %x to i1 - %b = and i1 %a, %y - %c = zext i1 %b to i32 - ret i32 %c -} - -; i16 - -define ptx_device i16 @cvt_i16_preds(i1 %x) { -; CHECK: selp.u16 %ret{{[0-9]+}}, 1, 0, %p{{[0-9]+}}; -; CHECK: ret; - %a = zext i1 %x to i16 - ret i16 %a -} - -define ptx_device i16 @cvt_i16_i32(i32 %x) { -; CHECK: cvt.u16.u32 %ret{{[0-9]+}}, %r{{[0-9]+}}; -; CHECK: ret; - %a = trunc i32 %x to i16 - ret i16 %a -} - -define ptx_device i16 @cvt_i16_i64(i64 %x) { -; CHECK: cvt.u16.u64 %ret{{[0-9]+}}, %rd{{[0-9]+}}; -; CHECK: ret; - %a = trunc i64 %x to i16 - ret i16 %a -} - -define ptx_device i16 @cvt_i16_f32(float %x) { -; CHECK: cvt.rzi.u16.f32 %ret{{[0-9]+}}, %f{{[0-9]+}}; -; CHECK: ret; - %a = fptoui float %x to i16 - ret i16 %a -} - -define ptx_device i16 @cvt_i16_f64(double %x) { -; CHECK: cvt.rzi.u16.f64 %ret{{[0-9]+}}, %fd{{[0-9]+}}; -; CHECK: ret; - %a = fptoui double %x to i16 - ret i16 %a -} - -; i32 - -define ptx_device i32 @cvt_i32_preds(i1 %x) { -; CHECK: selp.u32 %ret{{[0-9]+}}, 1, 0, %p{{[0-9]+}}; -; CHECK: ret; - %a = zext i1 %x to i32 - ret i32 %a -} - -define ptx_device i32 @cvt_i32_i16(i16 %x) { -; CHECK: cvt.u32.u16 %ret{{[0-9]+}}, %rh{{[0-9]+}}; -; CHECK: ret; - %a = zext i16 %x to i32 - ret i32 %a -} - -define ptx_device i32 @cvt_i32_i64(i64 %x) { -; CHECK: cvt.u32.u64 %ret{{[0-9]+}}, %rd{{[0-9]+}}; -; CHECK: ret; - %a = trunc i64 %x to i32 - ret i32 %a -} - -define ptx_device i32 @cvt_i32_f32(float %x) { -; CHECK: cvt.rzi.u32.f32 %ret{{[0-9]+}}, %f{{[0-9]+}}; -; CHECK: ret; - %a = fptoui float %x to i32 - ret i32 %a -} - -define ptx_device i32 @cvt_i32_f64(double %x) { -; CHECK: cvt.rzi.u32.f64 %ret{{[0-9]+}}, %fd{{[0-9]+}}; -; CHECK: ret; - %a = fptoui double %x to i32 - ret i32 %a -} - -; i64 - -define ptx_device i64 @cvt_i64_preds(i1 %x) { -; CHECK: selp.u64 %ret{{[0-9]+}}, 1, 0, %p{{[0-9]+}}; -; CHECK: ret; - %a = zext i1 %x to i64 - ret i64 %a -} - -define ptx_device i64 @cvt_i64_i16(i16 %x) { -; CHECK: cvt.u64.u16 %ret{{[0-9]+}}, %rh{{[0-9]+}}; -; CHECK: ret; - %a = zext i16 %x to i64 - ret i64 %a -} - -define ptx_device i64 @cvt_i64_i32(i32 %x) { -; CHECK: cvt.u64.u32 %ret{{[0-9]+}}, %r{{[0-9]+}}; -; CHECK: ret; - %a = zext i32 %x to i64 - ret i64 %a -} - -define ptx_device i64 @cvt_i64_f32(float %x) { -; CHECK: cvt.rzi.u64.f32 %ret{{[0-9]+}}, %f{{[0-9]+}}; -; CHECK: ret; - %a = fptoui float %x to i64 - ret i64 %a -} - -define ptx_device i64 @cvt_i64_f64(double %x) { -; CHECK: cvt.rzi.u64.f64 %ret{{[0-9]+}}, %fd{{[0-9]+}}; -; CHECK: ret; - %a = fptoui double %x to i64 - ret i64 %a -} - -; f32 - -define ptx_device float @cvt_f32_preds(i1 %x) { -; CHECK: mov.b32 %f0, 0; -; CHECK: mov.b32 %f1, 1065353216; -; CHECK: selp.f32 %ret{{[0-9]+}}, %f1, %f0, %p{{[0-9]+}}; -; 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 %ret{{[0-9]+}}, %rh{{[0-9]+}}; -; 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 %ret{{[0-9]+}}, %r{{[0-9]+}}; -; 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 %ret{{[0-9]+}}, %rd{{[0-9]+}}; -; 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 %ret{{[0-9]+}}, %fd{{[0-9]+}}; -; CHECK: ret; - %a = fptrunc double %x to float - ret float %a -} - -define ptx_device float @cvt_f32_s16(i16 %x) { -; CHECK: cvt.rn.f32.s16 %ret{{[0-9]+}}, %rh{{[0-9]+}} -; CHECK: ret - %a = sitofp i16 %x to float - ret float %a -} - -define ptx_device float @cvt_f32_s32(i32 %x) { -; CHECK: cvt.rn.f32.s32 %ret{{[0-9]+}}, %r{{[0-9]+}} -; CHECK: ret - %a = sitofp i32 %x to float - ret float %a -} - -define ptx_device float @cvt_f32_s64(i64 %x) { -; CHECK: cvt.rn.f32.s64 %ret{{[0-9]+}}, %rd{{[0-9]+}} -; CHECK: ret - %a = sitofp i64 %x to float - ret float %a -} - -; f64 - -define ptx_device double @cvt_f64_preds(i1 %x) { -; CHECK: mov.b64 %fd0, 0; -; CHECK: mov.b64 %fd1, 4575657221408423936; -; CHECK: selp.f64 %ret{{[0-9]+}}, %fd1, %fd0, %p{{[0-9]+}}; -; 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 %ret{{[0-9]+}}, %rh{{[0-9]+}}; -; 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 %ret{{[0-9]+}}, %r{{[0-9]+}}; -; 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 %ret{{[0-9]+}}, %rd{{[0-9]+}}; -; CHECK: ret; - %a = uitofp i64 %x to double - ret double %a -} - -define ptx_device double @cvt_f64_f32(float %x) { -; CHECK: cvt.f64.f32 %ret{{[0-9]+}}, %f{{[0-9]+}}; -; CHECK: ret; - %a = fpext float %x to double - ret double %a -} - -define ptx_device double @cvt_f64_s16(i16 %x) { -; CHECK: cvt.rn.f64.s16 %ret{{[0-9]+}}, %rh{{[0-9]+}} -; CHECK: ret - %a = sitofp i16 %x to double - ret double %a -} - -define ptx_device double @cvt_f64_s32(i32 %x) { -; CHECK: cvt.rn.f64.s32 %ret{{[0-9]+}}, %r{{[0-9]+}} -; CHECK: ret - %a = sitofp i32 %x to double - ret double %a -} - -define ptx_device double @cvt_f64_s64(i64 %x) { -; CHECK: cvt.rn.f64.s64 %ret{{[0-9]+}}, %rd{{[0-9]+}} -; CHECK: ret - %a = sitofp i64 %x to double - ret double %a -} diff --git a/test/CodeGen/PTX/exit.ll b/test/CodeGen/PTX/exit.ll deleted file mode 100644 index 7816c80172..0000000000 --- a/test/CodeGen/PTX/exit.ll +++ /dev/null @@ -1,14 +0,0 @@ -; RUN: llc < %s -march=ptx32 | FileCheck %s - -define ptx_kernel void @t1() { -; CHECK: exit; -; CHECK-NOT: ret; - ret void -} - -define ptx_kernel void @t2(i32* %p, i32 %x) { - store i32 %x, i32* %p -; CHECK: exit; -; CHECK-NOT: ret; - ret void -} diff --git a/test/CodeGen/PTX/fdiv-sm10.ll b/test/CodeGen/PTX/fdiv-sm10.ll deleted file mode 100644 index e1013befa2..0000000000 --- a/test/CodeGen/PTX/fdiv-sm10.ll +++ /dev/null @@ -1,15 +0,0 @@ -; RUN: llc < %s -march=ptx32 -mattr=+sm10 | FileCheck %s - -define ptx_device float @t1_f32(float %x, float %y) { -; CHECK: div.f32 %ret{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}; -; CHECK: ret; - %a = fdiv float %x, %y - ret float %a -} - -define ptx_device double @t1_f64(double %x, double %y) { -; CHECK: div.f64 %ret{{[0-9]+}}, %fd{{[0-9]+}}, %fd{{[0-9]+}}; -; CHECK: 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 deleted file mode 100644 index 1afa2ebd08..0000000000 --- a/test/CodeGen/PTX/fdiv-sm13.ll +++ /dev/null @@ -1,15 +0,0 @@ -; RUN: llc < %s -march=ptx32 -mattr=+sm13 | FileCheck %s - -define ptx_device float @t1_f32(float %x, float %y) { -; CHECK: div.rn.f32 %ret{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}; -; CHECK: ret; - %a = fdiv float %x, %y - ret float %a -} - -define ptx_device double @t1_f64(double %x, double %y) { -; CHECK: div.rn.f64 %ret{{[0-9]+}}, %fd{{[0-9]+}}, %fd{{[0-9]+}}; -; CHECK: ret; - %a = fdiv double %x, %y - ret double %a -} diff --git a/test/CodeGen/PTX/fneg.ll b/test/CodeGen/PTX/fneg.ll deleted file mode 100644 index 2b76e638f6..0000000000 --- a/test/CodeGen/PTX/fneg.ll +++ /dev/null @@ -1,15 +0,0 @@ -; RUN: llc < %s -march=ptx32 | FileCheck %s - -define ptx_device float @t1_f32(float %x) { -; CHECK: neg.f32 %ret{{[0-9]+}}, %f{{[0-9]+}}; -; CHECK: ret; - %y = fsub float -0.000000e+00, %x - ret float %y -} - -define ptx_device double @t1_f64(double %x) { -; CHECK: neg.f64 %ret{{[0-9]+}}, %fd{{[0-9]+}}; -; CHECK: ret; - %y = fsub double -0.000000e+00, %x - ret double %y -} diff --git a/test/CodeGen/PTX/intrinsic.ll b/test/CodeGen/PTX/intrinsic.ll deleted file mode 100644 index 9f37ead38d..0000000000 --- a/test/CodeGen/PTX/intrinsic.ll +++ /dev/null @@ -1,281 +0,0 @@ -; RUN: llc < %s -march=ptx32 -mattr=+ptx20 | FileCheck %s - -define ptx_device i32 @test_tid_x() { -; CHECK: mov.u32 %ret0, %tid.x; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.tid.x() - ret i32 %x -} - -define ptx_device i32 @test_tid_y() { -; CHECK: mov.u32 %ret0, %tid.y; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.tid.y() - ret i32 %x -} - -define ptx_device i32 @test_tid_z() { -; CHECK: mov.u32 %ret0, %tid.z; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.tid.z() - ret i32 %x -} - -define ptx_device i32 @test_tid_w() { -; CHECK: mov.u32 %ret0, %tid.w; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.tid.w() - ret i32 %x -} - -define ptx_device i32 @test_ntid_x() { -; CHECK: mov.u32 %ret0, %ntid.x; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.ntid.x() - ret i32 %x -} - -define ptx_device i32 @test_ntid_y() { -; CHECK: mov.u32 %ret0, %ntid.y; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.ntid.y() - ret i32 %x -} - -define ptx_device i32 @test_ntid_z() { -; CHECK: mov.u32 %ret0, %ntid.z; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.ntid.z() - ret i32 %x -} - -define ptx_device i32 @test_ntid_w() { -; CHECK: mov.u32 %ret0, %ntid.w; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.ntid.w() - ret i32 %x -} - -define ptx_device i32 @test_laneid() { -; CHECK: mov.u32 %ret0, %laneid; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.laneid() - ret i32 %x -} - -define ptx_device i32 @test_warpid() { -; CHECK: mov.u32 %ret0, %warpid; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.warpid() - ret i32 %x -} - -define ptx_device i32 @test_nwarpid() { -; CHECK: mov.u32 %ret0, %nwarpid; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.nwarpid() - ret i32 %x -} - -define ptx_device i32 @test_ctaid_x() { -; CHECK: mov.u32 %ret0, %ctaid.x; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.ctaid.x() - ret i32 %x -} - -define ptx_device i32 @test_ctaid_y() { -; CHECK: mov.u32 %ret0, %ctaid.y; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.ctaid.y() - ret i32 %x -} - -define ptx_device i32 @test_ctaid_z() { -; CHECK: mov.u32 %ret0, %ctaid.z; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.ctaid.z() - ret i32 %x -} - -define ptx_device i32 @test_ctaid_w() { -; CHECK: mov.u32 %ret0, %ctaid.w; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.ctaid.w() - ret i32 %x -} - -define ptx_device i32 @test_nctaid_x() { -; CHECK: mov.u32 %ret0, %nctaid.x; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.nctaid.x() - ret i32 %x -} - -define ptx_device i32 @test_nctaid_y() { -; CHECK: mov.u32 %ret0, %nctaid.y; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.nctaid.y() - ret i32 %x -} - -define ptx_device i32 @test_nctaid_z() { -; CHECK: mov.u32 %ret0, %nctaid.z; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.nctaid.z() - ret i32 %x -} - -define ptx_device i32 @test_nctaid_w() { -; CHECK: mov.u32 %ret0, %nctaid.w; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.nctaid.w() - ret i32 %x -} - -define ptx_device i32 @test_smid() { -; CHECK: mov.u32 %ret0, %smid; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.smid() - ret i32 %x -} - -define ptx_device i32 @test_nsmid() { -; CHECK: mov.u32 %ret0, %nsmid; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.nsmid() - ret i32 %x -} - -define ptx_device i32 @test_gridid() { -; CHECK: mov.u32 %ret0, %gridid; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.gridid() - ret i32 %x -} - -define ptx_device i32 @test_lanemask_eq() { -; CHECK: mov.u32 %ret0, %lanemask_eq; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.lanemask.eq() - ret i32 %x -} - -define ptx_device i32 @test_lanemask_le() { -; CHECK: mov.u32 %ret0, %lanemask_le; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.lanemask.le() - ret i32 %x -} - -define ptx_device i32 @test_lanemask_lt() { -; CHECK: mov.u32 %ret0, %lanemask_lt; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.lanemask.lt() - ret i32 %x -} - -define ptx_device i32 @test_lanemask_ge() { -; CHECK: mov.u32 %ret0, %lanemask_ge; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.lanemask.ge() - ret i32 %x -} - -define ptx_device i32 @test_lanemask_gt() { -; CHECK: mov.u32 %ret0, %lanemask_gt; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.lanemask.gt() - ret i32 %x -} - -define ptx_device i32 @test_clock() { -; CHECK: mov.u32 %ret0, %clock; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.clock() - ret i32 %x -} - -define ptx_device i64 @test_clock64() { -; CHECK: mov.u64 %ret0, %clock64; -; CHECK: ret; - %x = call i64 @llvm.ptx.read.clock64() - ret i64 %x -} - -define ptx_device i32 @test_pm0() { -; CHECK: mov.u32 %ret0, %pm0; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.pm0() - ret i32 %x -} - -define ptx_device i32 @test_pm1() { -; CHECK: mov.u32 %ret0, %pm1; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.pm1() - ret i32 %x -} - -define ptx_device i32 @test_pm2() { -; CHECK: mov.u32 %ret0, %pm2; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.pm2() - ret i32 %x -} - -define ptx_device i32 @test_pm3() { -; CHECK: mov.u32 %ret0, %pm3; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.pm3() - ret i32 %x -} - -define ptx_device void @test_bar_sync() { -; CHECK: bar.sync 0 -; CHECK: ret; - call void @llvm.ptx.bar.sync(i32 0) - ret void -} - -declare i32 @llvm.ptx.read.tid.x() -declare i32 @llvm.ptx.read.tid.y() -declare i32 @llvm.ptx.read.tid.z() -declare i32 @llvm.ptx.read.tid.w() -declare i32 @llvm.ptx.read.ntid.x() -declare i32 @llvm.ptx.read.ntid.y() -declare i32 @llvm.ptx.read.ntid.z() -declare i32 @llvm.ptx.read.ntid.w() - -declare i32 @llvm.ptx.read.laneid() -declare i32 @llvm.ptx.read.warpid() -declare i32 @llvm.ptx.read.nwarpid() - -declare i32 @llvm.ptx.read.ctaid.x() -declare i32 @llvm.ptx.read.ctaid.y() -declare i32 @llvm.ptx.read.ctaid.z() -declare i32 @llvm.ptx.read.ctaid.w() -declare i32 @llvm.ptx.read.nctaid.x() -declare i32 @llvm.ptx.read.nctaid.y() -declare i32 @llvm.ptx.read.nctaid.z() -declare i32 @llvm.ptx.read.nctaid.w() - -declare i32 @llvm.ptx.read.smid() -declare i32 @llvm.ptx.read.nsmid() -declare i32 @llvm.ptx.read.gridid() - -declare i32 @llvm.ptx.read.lanemask.eq() -declare i32 @llvm.ptx.read.lanemask.le() -declare i32 @llvm.ptx.read.lanemask.lt() -declare i32 @llvm.ptx.read.lanemask.ge() -declare i32 @llvm.ptx.read.lanemask.gt() - -declare i32 @llvm.ptx.read.clock() -declare i64 @llvm.ptx.read.clock64() - -declare i32 @llvm.ptx.read.pm0() -declare i32 @llvm.ptx.read.pm1() -declare i32 @llvm.ptx.read.pm2() -declare i32 @llvm.ptx.read.pm3() - -declare void @llvm.ptx.bar.sync(i32 %i) diff --git a/test/CodeGen/PTX/ld.ll b/test/CodeGen/PTX/ld.ll deleted file mode 100644 index e55820dfb0..0000000000 --- a/test/CodeGen/PTX/ld.ll +++ /dev/null @@ -1,382 +0,0 @@ -; RUN: llc < %s -march=ptx32 | FileCheck %s - -;CHECK: .extern .global .b16 array_i16[10]; -@array_i16 = external global [10 x i16] - -;CHECK: .extern .const .b16 array_constant_i16[10]; -@array_constant_i16 = external addrspace(1) constant [10 x i16] - -;CHECK: .extern .shared .b16 array_shared_i16[10]; -@array_shared_i16 = external addrspace(4) global [10 x i16] - -;CHECK: .extern .global .b32 array_i32[10]; -@array_i32 = external global [10 x i32] - -;CHECK: .extern .const .b32 array_constant_i32[10]; -@array_constant_i32 = external addrspace(1) constant [10 x i32] - -;CHECK: .extern .shared .b32 array_shared_i32[10]; -@array_shared_i32 = external addrspace(4) global [10 x i32] - -;CHECK: .extern .global .b64 array_i64[10]; -@array_i64 = external global [10 x i64] - -;CHECK: .extern .const .b64 array_constant_i64[10]; -@array_constant_i64 = external addrspace(1) constant [10 x i64] - -;CHECK: .extern .shared .b64 array_shared_i64[10]; -@array_shared_i64 = external addrspace(4) global [10 x i64] - -;CHECK: .extern .global .b32 array_float[10]; -@array_float = external global [10 x float] - -;CHECK: .extern .const .b32 array_constant_float[10]; -@array_constant_float = external addrspace(1) constant [10 x float] - -;CHECK: .extern .shared .b32 array_shared_float[10]; -@array_shared_float = external addrspace(4) global [10 x float] - -;CHECK: .extern .global .b64 array_double[10]; -@array_double = external global [10 x double] - -;CHECK: .extern .const .b64 array_constant_double[10]; -@array_constant_double = external addrspace(1) constant [10 x double] - -;CHECK: .extern .shared .b64 array_shared_double[10]; -@array_shared_double = external addrspace(4) global [10 x double] - - -define ptx_device i16 @t1_u16(i16* %p) { -entry: -;CHECK: ld.global.u16 %ret{{[0-9]+}}, [%r{{[0-9]+}}]; -;CHECK: ret; - %x = load i16* %p - ret i16 %x -} - -define ptx_device i32 @t1_u32(i32* %p) { -entry: -;CHECK: ld.global.u32 %ret{{[0-9]+}}, [%r{{[0-9]+}}]; -;CHECK: ret; - %x = load i32* %p - ret i32 %x -} - -define ptx_device i64 @t1_u64(i64* %p) { -entry: -;CHECK: ld.global.u64 %ret{{[0-9]+}}, [%r{{[0-9]+}}]; -;CHECK: ret; - %x = load i64* %p - ret i64 %x -} - -define ptx_device float @t1_f32(float* %p) { -entry: -;CHECK: ld.global.f32 %ret{{[0-9]+}}, [%r{{[0-9]+}}]; -;CHECK: ret; - %x = load float* %p - ret float %x -} - -define ptx_device double @t1_f64(double* %p) { -entry: -;CHECK: ld.global.f64 %ret{{[0-9]+}}, [%r{{[0-9]+}}]; -;CHECK: ret; - %x = load double* %p - ret double %x -} - -define ptx_device i16 @t2_u16(i16* %p) { -entry: -;CHECK: ld.global.u16 %ret{{[0-9]+}}, [%r{{[0-9]+}}+2]; -;CHECK: ret; - %i = getelementptr i16* %p, i32 1 - %x = load i16* %i - ret i16 %x -} - -define ptx_device i32 @t2_u32(i32* %p) { -entry: -;CHECK: ld.global.u32 %ret{{[0-9]+}}, [%r{{[0-9]+}}+4]; -;CHECK: ret; - %i = getelementptr i32* %p, i32 1 - %x = load i32* %i - ret i32 %x -} - -define ptx_device i64 @t2_u64(i64* %p) { -entry: -;CHECK: ld.global.u64 %ret{{[0-9]+}}, [%r{{[0-9]+}}+8]; -;CHECK: ret; - %i = getelementptr i64* %p, i32 1 - %x = load i64* %i - ret i64 %x -} - -define ptx_device float @t2_f32(float* %p) { -entry: -;CHECK: ld.global.f32 %ret{{[0-9]+}}, [%r{{[0-9]+}}+4]; -;CHECK: ret; - %i = getelementptr float* %p, i32 1 |