diff options
Diffstat (limited to 'test/CodeGen/NVPTX')
58 files changed, 1750 insertions, 151 deletions
diff --git a/test/CodeGen/NVPTX/access-non-generic.ll b/test/CodeGen/NVPTX/access-non-generic.ll new file mode 100644 index 000000000000..c225abf0fd85 --- /dev/null +++ b/test/CodeGen/NVPTX/access-non-generic.ll @@ -0,0 +1,91 @@ +; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix PTX +; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefix PTX +; RUN: opt < %s -S -nvptx-favor-non-generic -dce | FileCheck %s --check-prefix IR + +@array = internal addrspace(3) global [10 x float] zeroinitializer, align 4 +@scalar = internal addrspace(3) global float 0.000000e+00, align 4 + +; Verifies nvptx-favor-non-generic correctly optimizes generic address space +; usage to non-generic address space usage for the patterns we claim to handle: +; 1. load cast +; 2. store cast +; 3. load gep cast +; 4. store gep cast +; gep and cast can be an instruction or a constant expression. This function +; tries all possible combinations. +define float @ld_st_shared_f32(i32 %i, float %v) { +; IR-LABEL: @ld_st_shared_f32 +; IR-NOT: addrspacecast +; PTX-LABEL: ld_st_shared_f32( + ; load cast + %1 = load float* addrspacecast (float addrspace(3)* @scalar to float*), align 4 +; PTX: ld.shared.f32 %f{{[0-9]+}}, [scalar]; + ; store cast + store float %v, float* addrspacecast (float addrspace(3)* @scalar to float*), align 4 +; PTX: st.shared.f32 [scalar], %f{{[0-9]+}}; + ; use syncthreads to disable optimizations across components + call void @llvm.cuda.syncthreads() +; PTX: bar.sync 0; + + ; cast; load + %2 = addrspacecast float addrspace(3)* @scalar to float* + %3 = load float* %2, align 4 +; PTX: ld.shared.f32 %f{{[0-9]+}}, [scalar]; + ; cast; store + store float %v, float* %2, align 4 +; PTX: st.shared.f32 [scalar], %f{{[0-9]+}}; + call void @llvm.cuda.syncthreads() +; PTX: bar.sync 0; + + ; load gep cast + %4 = load float* getelementptr inbounds ([10 x float]* addrspacecast ([10 x float] addrspace(3)* @array to [10 x float]*), i32 0, i32 5), align 4 +; PTX: ld.shared.f32 %f{{[0-9]+}}, [array+20]; + ; store gep cast + store float %v, float* getelementptr inbounds ([10 x float]* addrspacecast ([10 x float] addrspace(3)* @array to [10 x float]*), i32 0, i32 5), align 4 +; PTX: st.shared.f32 [array+20], %f{{[0-9]+}}; + call void @llvm.cuda.syncthreads() +; PTX: bar.sync 0; + + ; gep cast; load + %5 = getelementptr inbounds [10 x float]* addrspacecast ([10 x float] addrspace(3)* @array to [10 x float]*), i32 0, i32 5 + %6 = load float* %5, align 4 +; PTX: ld.shared.f32 %f{{[0-9]+}}, [array+20]; + ; gep cast; store + store float %v, float* %5, align 4 +; PTX: st.shared.f32 [array+20], %f{{[0-9]+}}; + call void @llvm.cuda.syncthreads() +; PTX: bar.sync 0; + + ; cast; gep; load + %7 = addrspacecast [10 x float] addrspace(3)* @array to [10 x float]* + %8 = getelementptr inbounds [10 x float]* %7, i32 0, i32 %i + %9 = load float* %8, align 4 +; PTX: ld.shared.f32 %f{{[0-9]+}}, [%{{(r|rl|rd)[0-9]+}}]; + ; cast; gep; store + store float %v, float* %8, align 4 +; PTX: st.shared.f32 [%{{(r|rl|rd)[0-9]+}}], %f{{[0-9]+}}; + call void @llvm.cuda.syncthreads() +; PTX: bar.sync 0; + + %sum2 = fadd float %1, %3 + %sum3 = fadd float %sum2, %4 + %sum4 = fadd float %sum3, %6 + %sum5 = fadd float %sum4, %9 + ret float %sum5 +} + +; When hoisting an addrspacecast between different pointer types, replace the +; addrspacecast with a bitcast. +define i32 @ld_int_from_float() { +; IR-LABEL: @ld_int_from_float +; IR: load i32 addrspace(3)* bitcast (float addrspace(3)* @scalar to i32 addrspace(3)*) +; PTX-LABEL: ld_int_from_float( +; PTX: ld.shared.u{{(32|64)}} + %1 = load i32* addrspacecast(float addrspace(3)* @scalar to i32*), align 4 + ret i32 %1 +} + +declare void @llvm.cuda.syncthreads() #3 + +attributes #3 = { noduplicate nounwind } + diff --git a/test/CodeGen/NVPTX/addrspacecast-gvar.ll b/test/CodeGen/NVPTX/addrspacecast-gvar.ll new file mode 100644 index 000000000000..6afbdb8a429f --- /dev/null +++ b/test/CodeGen/NVPTX/addrspacecast-gvar.ll @@ -0,0 +1,9 @@ +; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s + +; CHECK: .visible .global .align 4 .u32 g = 42; +; CHECK: .visible .global .align 4 .u32 g2 = generic(g); +; CHECK: .visible .global .align 4 .u32 g3 = g; + +@g = addrspace(1) global i32 42 +@g2 = addrspace(1) global i32* addrspacecast (i32 addrspace(1)* @g to i32*) +@g3 = addrspace(1) global i32 addrspace(1)* @g diff --git a/test/CodeGen/NVPTX/addrspacecast.ll b/test/CodeGen/NVPTX/addrspacecast.ll new file mode 100644 index 000000000000..03b9a9844752 --- /dev/null +++ b/test/CodeGen/NVPTX/addrspacecast.ll @@ -0,0 +1,99 @@ +; RUN: llc < %s -march=nvptx -mcpu=sm_20 -disable-nvptx-favor-non-generic | FileCheck %s -check-prefix=PTX32 +; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -disable-nvptx-favor-non-generic | FileCheck %s -check-prefix=PTX64 + + +define i32 @conv1(i32 addrspace(1)* %ptr) { +; PTX32: conv1 +; PTX32: cvta.global.u32 +; PTX32: ld.u32 +; PTX64: conv1 +; PTX64: cvta.global.u64 +; PTX64: ld.u32 + %genptr = addrspacecast i32 addrspace(1)* %ptr to i32* + %val = load i32* %genptr + ret i32 %val +} + +define i32 @conv2(i32 addrspace(3)* %ptr) { +; PTX32: conv2 +; PTX32: cvta.shared.u32 +; PTX32: ld.u32 +; PTX64: conv2 +; PTX64: cvta.shared.u64 +; PTX64: ld.u32 + %genptr = addrspacecast i32 addrspace(3)* %ptr to i32* + %val = load i32* %genptr + ret i32 %val +} + +define i32 @conv3(i32 addrspace(4)* %ptr) { +; PTX32: conv3 +; PTX32: cvta.const.u32 +; PTX32: ld.u32 +; PTX64: conv3 +; PTX64: cvta.const.u64 +; PTX64: ld.u32 + %genptr = addrspacecast i32 addrspace(4)* %ptr to i32* + %val = load i32* %genptr + ret i32 %val +} + +define i32 @conv4(i32 addrspace(5)* %ptr) { +; PTX32: conv4 +; PTX32: cvta.local.u32 +; PTX32: ld.u32 +; PTX64: conv4 +; PTX64: cvta.local.u64 +; PTX64: ld.u32 + %genptr = addrspacecast i32 addrspace(5)* %ptr to i32* + %val = load i32* %genptr + ret i32 %val +} + +define i32 @conv5(i32* %ptr) { +; PTX32: conv5 +; PTX32: cvta.to.global.u32 +; PTX32: ld.global.u32 +; PTX64: conv5 +; PTX64: cvta.to.global.u64 +; PTX64: ld.global.u32 + %specptr = addrspacecast i32* %ptr to i32 addrspace(1)* + %val = load i32 addrspace(1)* %specptr + ret i32 %val +} + +define i32 @conv6(i32* %ptr) { +; PTX32: conv6 +; PTX32: cvta.to.shared.u32 +; PTX32: ld.shared.u32 +; PTX64: conv6 +; PTX64: cvta.to.shared.u64 +; PTX64: ld.shared.u32 + %specptr = addrspacecast i32* %ptr to i32 addrspace(3)* + %val = load i32 addrspace(3)* %specptr + ret i32 %val +} + +define i32 @conv7(i32* %ptr) { +; PTX32: conv7 +; PTX32: cvta.to.const.u32 +; PTX32: ld.const.u32 +; PTX64: conv7 +; PTX64: cvta.to.const.u64 +; PTX64: ld.const.u32 + %specptr = addrspacecast i32* %ptr to i32 addrspace(4)* + %val = load i32 addrspace(4)* %specptr + ret i32 %val +} + +define i32 @conv8(i32* %ptr) { +; PTX32: conv8 +; PTX32: cvta.to.local.u32 +; PTX32: ld.local.u32 +; PTX64: conv8 +; PTX64: cvta.to.local.u64 +; PTX64: ld.local.u32 + %specptr = addrspacecast i32* %ptr to i32 addrspace(5)* + %val = load i32 addrspace(5)* %specptr + ret i32 %val +} diff --git a/test/CodeGen/NVPTX/aggr-param.ll b/test/CodeGen/NVPTX/aggr-param.ll new file mode 100644 index 000000000000..21deb7ebce87 --- /dev/null +++ b/test/CodeGen/NVPTX/aggr-param.ll @@ -0,0 +1,20 @@ +; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s + +; Make sure aggregate param types get emitted properly. + +%struct.float4 = type { float, float, float, float } + +; CHECK: .visible .func bar +; CHECK: .param .align 4 .b8 bar_param_0[16] +define void @bar(%struct.float4 %f) { +entry: + ret void +} + +; CHECK: .visible .func foo +; CHECK: .param .align 4 .b8 foo_param_0[20] +define void @foo([5 x i32] %f) { +entry: + ret void +} + diff --git a/test/CodeGen/NVPTX/arg-lowering.ll b/test/CodeGen/NVPTX/arg-lowering.ll new file mode 100644 index 000000000000..f7b8a1491d37 --- /dev/null +++ b/test/CodeGen/NVPTX/arg-lowering.ll @@ -0,0 +1,13 @@ +; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s + +; CHECK: .visible .func (.param .align 16 .b8 func_retval0[16]) foo0( +; CHECK: .param .align 4 .b8 foo0_param_0[8] +define <4 x float> @foo0({float, float} %arg0) { + ret <4 x float> <float 1.0, float 1.0, float 1.0, float 1.0> +} + +; CHECK: .visible .func (.param .align 8 .b8 func_retval0[8]) foo1( +; CHECK: .param .align 8 .b8 foo1_param_0[16] +define <2 x float> @foo1({float, float, i64} %arg0) { + ret <2 x float> <float 1.0, float 1.0> +} diff --git a/test/CodeGen/NVPTX/arithmetic-fp-sm20.ll b/test/CodeGen/NVPTX/arithmetic-fp-sm20.ll index e474fa4df5ce..c167db4b46dc 100644 --- a/test/CodeGen/NVPTX/arithmetic-fp-sm20.ll +++ b/test/CodeGen/NVPTX/arithmetic-fp-sm20.ll @@ -1,5 +1,5 @@ -; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s -; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s +; RUN: llc < %s -march=nvptx -mcpu=sm_20 -fp-contract=fast | FileCheck %s +; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -fp-contract=fast | FileCheck %s ;; These tests should run for all targets @@ -9,28 +9,28 @@ ;;; f64 define double @fadd_f64(double %a, double %b) { -; CHECK: add.f64 %fl{{[0-9]+}}, %fl{{[0-9]+}}, %fl{{[0-9]+}} +; CHECK: add.f64 %fd{{[0-9]+}}, %fd{{[0-9]+}}, %fd{{[0-9]+}} ; CHECK: ret %ret = fadd double %a, %b ret double %ret } define double @fsub_f64(double %a, double %b) { -; CHECK: sub.f64 %fl{{[0-9]+}}, %fl{{[0-9]+}}, %fl{{[0-9]+}} +; CHECK: sub.f64 %fd{{[0-9]+}}, %fd{{[0-9]+}}, %fd{{[0-9]+}} ; CHECK: ret %ret = fsub double %a, %b ret double %ret } define double @fmul_f64(double %a, double %b) { -; CHECK: mul.f64 %fl{{[0-9]+}}, %fl{{[0-9]+}}, %fl{{[0-9]+}} +; CHECK: mul.f64 %fd{{[0-9]+}}, %fd{{[0-9]+}}, %fd{{[0-9]+}} ; CHECK: ret %ret = fmul double %a, %b ret double %ret } define double @fdiv_f64(double %a, double %b) { -; CHECK: div.rn.f64 %fl{{[0-9]+}}, %fl{{[0-9]+}}, %fl{{[0-9]+}} +; CHECK: div.rn.f64 %fd{{[0-9]+}}, %fd{{[0-9]+}}, %fd{{[0-9]+}} ; CHECK: ret %ret = fdiv double %a, %b ret double %ret diff --git a/test/CodeGen/NVPTX/arithmetic-int.ll b/test/CodeGen/NVPTX/arithmetic-int.ll index 8d73b7e6c4c6..b5a2872299b7 100644 --- a/test/CodeGen/NVPTX/arithmetic-int.ll +++ b/test/CodeGen/NVPTX/arithmetic-int.ll @@ -9,70 +9,70 @@ ;;; i64 define i64 @add_i64(i64 %a, i64 %b) { -; CHECK: add.s64 %rl{{[0-9]+}}, %rl{{[0-9]+}}, %rl{{[0-9]+}} +; CHECK: add.s64 %rd{{[0-9]+}}, %rd{{[0-9]+}}, %rd{{[0-9]+}} ; CHECK: ret %ret = add i64 %a, %b ret i64 %ret } define i64 @sub_i64(i64 %a, i64 %b) { -; CHECK: sub.s64 %rl{{[0-9]+}}, %rl{{[0-9]+}}, %rl{{[0-9]+}} +; CHECK: sub.s64 %rd{{[0-9]+}}, %rd{{[0-9]+}}, %rd{{[0-9]+}} ; CHECK: ret %ret = sub i64 %a, %b ret i64 %ret } define i64 @mul_i64(i64 %a, i64 %b) { -; CHECK: mul.lo.s64 %rl{{[0-9]+}}, %rl{{[0-9]+}}, %rl{{[0-9]+}} +; CHECK: mul.lo.s64 %rd{{[0-9]+}}, %rd{{[0-9]+}}, %rd{{[0-9]+}} ; CHECK: ret %ret = mul i64 %a, %b ret i64 %ret } define i64 @sdiv_i64(i64 %a, i64 %b) { -; CHECK: div.s64 %rl{{[0-9]+}}, %rl{{[0-9]+}}, %rl{{[0-9]+}} +; CHECK: div.s64 %rd{{[0-9]+}}, %rd{{[0-9]+}}, %rd{{[0-9]+}} ; CHECK: ret %ret = sdiv i64 %a, %b ret i64 %ret } define i64 @udiv_i64(i64 %a, i64 %b) { -; CHECK: div.u64 %rl{{[0-9]+}}, %rl{{[0-9]+}}, %rl{{[0-9]+}} +; CHECK: div.u64 %rd{{[0-9]+}}, %rd{{[0-9]+}}, %rd{{[0-9]+}} ; CHECK: ret %ret = udiv i64 %a, %b ret i64 %ret } define i64 @srem_i64(i64 %a, i64 %b) { -; CHECK: rem.s64 %rl{{[0-9]+}}, %rl{{[0-9]+}}, %rl{{[0-9]+}} +; CHECK: rem.s64 %rd{{[0-9]+}}, %rd{{[0-9]+}}, %rd{{[0-9]+}} ; CHECK: ret %ret = srem i64 %a, %b ret i64 %ret } define i64 @urem_i64(i64 %a, i64 %b) { -; CHECK: rem.u64 %rl{{[0-9]+}}, %rl{{[0-9]+}}, %rl{{[0-9]+}} +; CHECK: rem.u64 %rd{{[0-9]+}}, %rd{{[0-9]+}}, %rd{{[0-9]+}} ; CHECK: ret %ret = urem i64 %a, %b ret i64 %ret } define i64 @and_i64(i64 %a, i64 %b) { -; CHECK: and.b64 %rl{{[0-9]+}}, %rl{{[0-9]+}}, %rl{{[0-9]+}} +; CHECK: and.b64 %rd{{[0-9]+}}, %rd{{[0-9]+}}, %rd{{[0-9]+}} ; CHECK: ret %ret = and i64 %a, %b ret i64 %ret } define i64 @or_i64(i64 %a, i64 %b) { -; CHECK: or.b64 %rl{{[0-9]+}}, %rl{{[0-9]+}}, %rl{{[0-9]+}} +; CHECK: or.b64 %rd{{[0-9]+}}, %rd{{[0-9]+}}, %rd{{[0-9]+}} ; CHECK: ret %ret = or i64 %a, %b ret i64 %ret } define i64 @xor_i64(i64 %a, i64 %b) { -; CHECK: xor.b64 %rl{{[0-9]+}}, %rl{{[0-9]+}}, %rl{{[0-9]+}} +; CHECK: xor.b64 %rd{{[0-9]+}}, %rd{{[0-9]+}}, %rd{{[0-9]+}} ; CHECK: ret %ret = xor i64 %a, %b ret i64 %ret @@ -80,7 +80,7 @@ define i64 @xor_i64(i64 %a, i64 %b) { define i64 @shl_i64(i64 %a, i64 %b) { ; PTX requires 32-bit shift amount -; CHECK: shl.b64 %rl{{[0-9]+}}, %rl{{[0-9]+}}, %r{{[0-9]+}} +; CHECK: shl.b64 %rd{{[0-9]+}}, %rd{{[0-9]+}}, %r{{[0-9]+}} ; CHECK: ret %ret = shl i64 %a, %b ret i64 %ret @@ -88,7 +88,7 @@ define i64 @shl_i64(i64 %a, i64 %b) { define i64 @ashr_i64(i64 %a, i64 %b) { ; PTX requires 32-bit shift amount -; CHECK: shr.s64 %rl{{[0-9]+}}, %rl{{[0-9]+}}, %r{{[0-9]+}} +; CHECK: shr.s64 %rd{{[0-9]+}}, %rd{{[0-9]+}}, %r{{[0-9]+}} ; CHECK: ret %ret = ashr i64 %a, %b ret i64 %ret @@ -96,7 +96,7 @@ define i64 @ashr_i64(i64 %a, i64 %b) { define i64 @lshr_i64(i64 %a, i64 %b) { ; PTX requires 32-bit shift amount -; CHECK: shr.u64 %rl{{[0-9]+}}, %rl{{[0-9]+}}, %r{{[0-9]+}} +; CHECK: shr.u64 %rd{{[0-9]+}}, %rd{{[0-9]+}}, %r{{[0-9]+}} ; CHECK: ret %ret = lshr i64 %a, %b ret i64 %ret diff --git a/test/CodeGen/NVPTX/atomics.ll b/test/CodeGen/NVPTX/atomics.ll new file mode 100644 index 000000000000..daadb6e9c1a0 --- /dev/null +++ b/test/CodeGen/NVPTX/atomics.ll @@ -0,0 +1,182 @@ +; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s + + +; CHECK-LABEL: atom0 +define i32 @atom0(i32* %addr, i32 %val) { +; CHECK: atom.add.u32 + %ret = atomicrmw add i32* %addr, i32 %val seq_cst + ret i32 %ret +} + +; CHECK-LABEL: atom1 +define i64 @atom1(i64* %addr, i64 %val) { +; CHECK: atom.add.u64 + %ret = atomicrmw add i64* %addr, i64 %val seq_cst + ret i64 %ret +} + +; CHECK-LABEL: atom2 +define i32 @atom2(i32* %subr, i32 %val) { +; CHECK: neg.s32 +; CHECK: atom.add.u32 + %ret = atomicrmw sub i32* %subr, i32 %val seq_cst + ret i32 %ret +} + +; CHECK-LABEL: atom3 +define i64 @atom3(i64* %subr, i64 %val) { +; CHECK: neg.s64 +; CHECK: atom.add.u64 + %ret = atomicrmw sub i64* %subr, i64 %val seq_cst + ret i64 %ret +} + +; CHECK-LABEL: atom4 +define i32 @atom4(i32* %subr, i32 %val) { +; CHECK: atom.and.b32 + %ret = atomicrmw and i32* %subr, i32 %val seq_cst + ret i32 %ret +} + +; CHECK-LABEL: atom5 +define i64 @atom5(i64* %subr, i64 %val) { +; CHECK: atom.and.b64 + %ret = atomicrmw and i64* %subr, i64 %val seq_cst + ret i64 %ret +} + +;; NAND not yet supported +;define i32 @atom6(i32* %subr, i32 %val) { +; %ret = atomicrmw nand i32* %subr, i32 %val seq_cst +; ret i32 %ret +;} + +;define i64 @atom7(i64* %subr, i64 %val) { +; %ret = atomicrmw nand i64* %subr, i64 %val seq_cst +; ret i64 %ret +;} + +; CHECK-LABEL: atom8 +define i32 @atom8(i32* %subr, i32 %val) { +; CHECK: atom.or.b32 + %ret = atomicrmw or i32* %subr, i32 %val seq_cst + ret i32 %ret +} + +; CHECK-LABEL: atom9 +define i64 @atom9(i64* %subr, i64 %val) { +; CHECK: atom.or.b64 + %ret = atomicrmw or i64* %subr, i64 %val seq_cst + ret i64 %ret +} + +; CHECK-LABEL: atom10 +define i32 @atom10(i32* %subr, i32 %val) { +; CHECK: atom.xor.b32 + %ret = atomicrmw xor i32* %subr, i32 %val seq_cst + ret i32 %ret +} + +; CHECK-LABEL: atom11 +define i64 @atom11(i64* %subr, i64 %val) { +; CHECK: atom.xor.b64 + %ret = atomicrmw xor i64* %subr, i64 %val seq_cst + ret i64 %ret +} + +; CHECK-LABEL: atom12 +define i32 @atom12(i32* %subr, i32 %val) { +; CHECK: atom.max.s32 + %ret = atomicrmw max i32* %subr, i32 %val seq_cst + ret i32 %ret +} + +; CHECK-LABEL: atom13 +define i64 @atom13(i64* %subr, i64 %val) { +; CHECK: atom.max.s64 + %ret = atomicrmw max i64* %subr, i64 %val seq_cst + ret i64 %ret +} + +; CHECK-LABEL: atom14 +define i32 @atom14(i32* %subr, i32 %val) { +; CHECK: atom.min.s32 + %ret = atomicrmw min i32* %subr, i32 %val seq_cst + ret i32 %ret +} + +; CHECK-LABEL: atom15 +define i64 @atom15(i64* %subr, i64 %val) { +; CHECK: atom.min.s64 + %ret = atomicrmw min i64* %subr, i64 %val seq_cst + ret i64 %ret +} + +; CHECK-LABEL: atom16 +define i32 @atom16(i32* %subr, i32 %val) { +; CHECK: atom.max.u32 + %ret = atomicrmw umax i32* %subr, i32 %val seq_cst + ret i32 %ret +} + +; CHECK-LABEL: atom17 +define i64 @atom17(i64* %subr, i64 %val) { +; CHECK: atom.max.u64 + %ret = atomicrmw umax i64* %subr, i64 %val seq_cst + ret i64 %ret +} + +; CHECK-LABEL: atom18 +define i32 @atom18(i32* %subr, i32 %val) { +; CHECK: atom.min.u32 + %ret = atomicrmw umin i32* %subr, i32 %val seq_cst + ret i32 %ret +} + +; CHECK-LABEL: atom19 +define i64 @atom19(i64* %subr, i64 %val) { +; CHECK: atom.min.u64 + %ret = atomicrmw umin i64* %subr, i64 %val seq_cst + ret i64 %ret +} + +declare float @llvm.nvvm.atomic.load.add.f32.p0f32(float* %addr, float %val) + +; CHECK-LABEL: atomic_add_f32_generic +define float @atomic_add_f32_generic(float* %addr, float %val) { +; CHECK: atom.add.f32 + %ret = call float @llvm.nvvm.atomic.load.add.f32.p0f32(float* %addr, float %val) + ret float %ret +} + +declare float @llvm.nvvm.atomic.load.add.f32.p1f32(float addrspace(1)* %addr, float %val) + +; CHECK-LABEL: atomic_add_f32_addrspace1 +define float @atomic_add_f32_addrspace1(float addrspace(1)* %addr, float %val) { +; CHECK: atom.global.add.f32 + %ret = call float @llvm.nvvm.atomic.load.add.f32.p1f32(float addrspace(1)* %addr, float %val) + ret float %ret +} + +declare float @llvm.nvvm.atomic.load.add.f32.p3f32(float addrspace(3)* %addr, float %val) + +; CHECK-LABEL: atomic_add_f32_addrspace3 +define float @atomic_add_f32_addrspace3(float addrspace(3)* %addr, float %val) { +; CHECK: atom.shared.add.f32 + %ret = call float @llvm.nvvm.atomic.load.add.f32.p3f32(float addrspace(3)* %addr, float %val) + ret float %ret +} + +; CHECK-LABEL: atomic_cmpxchg_i32 +define i32 @atomic_cmpxchg_i32(i32* %addr, i32 %cmp, i32 %new) { +; CHECK: atom.cas.b32 + %pairold = cmpxchg i32* %addr, i32 %cmp, i32 %new seq_cst seq_cst + ret i32 %new +} + +; CHECK-LABEL: atomic_cmpxchg_i64 +define i64 @atomic_cmpxchg_i64(i64* %addr, i64 %cmp, i64 %new) { +; CHECK: atom.cas.b64 + %pairold = cmpxchg i64* %addr, i64 %cmp, i64 %new seq_cst seq_cst + ret i64 %new +} diff --git a/test/CodeGen/NVPTX/bfe.ll b/test/CodeGen/NVPTX/bfe.ll new file mode 100644 index 000000000000..2e816fec2c59 --- /dev/null +++ b/test/CodeGen/NVPTX/bfe.ll @@ -0,0 +1,32 @@ +; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s + + +; CHECK: bfe0 +define i32 @bfe0(i32 %a) { +; CHECK: bfe.u32 %r{{[0-9]+}}, %r{{[0-9]+}}, 4, 4 +; CHECK-NOT: shr +; CHECK-NOT: and + %val0 = ashr i32 %a, 4 + %val1 = and i32 %val0, 15 + ret i32 %val1 +} + +; CHECK: bfe1 +define i32 @bfe1(i32 %a) { +; CHECK: bfe.u32 %r{{[0-9]+}}, %r{{[0-9]+}}, 3, 3 +; CHECK-NOT: shr +; CHECK-NOT: and + %val0 = ashr i32 %a, 3 + %val1 = and i32 %val0, 7 + ret i32 %val1 +} + +; CHECK: bfe2 +define i32 @bfe2(i32 %a) { +; CHECK: bfe.u32 %r{{[0-9]+}}, %r{{[0-9]+}}, 5, 3 +; CHECK-NOT: shr +; CHECK-NOT: and + %val0 = ashr i32 %a, 5 + %val1 = and i32 %val0, 7 + ret i32 %val1 +} diff --git a/test/CodeGen/NVPTX/bug17709.ll b/test/CodeGen/NVPTX/bug17709.ll index 92f0fcb11e41..076c44684579 100644 --- a/test/CodeGen/NVPTX/bug17709.ll +++ b/test/CodeGen/NVPTX/bug17709.ll @@ -4,7 +4,7 @@ target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64" target triple = "nvptx64-nvidia-cuda" -define linker_private ptx_device { double, double } @__utils1_MOD_trace(%"struct.array2_complex(kind=8).43.5.57"* noalias %m) { +define private ptx_device { double, double } @__utils1_MOD_trace(%"struct.array2_complex(kind=8).43.5.57"* noalias %m) { entry: ;unreachable %t0 = insertvalue {double, double} undef, double 1.0, 0 diff --git a/test/CodeGen/NVPTX/call-with-alloca-buffer.ll b/test/CodeGen/NVPTX/call-with-alloca-buffer.ll new file mode 100644 index 000000000000..83d491637041 --- /dev/null +++ b/test/CodeGen/NVPTX/call-with-alloca-buffer.ll @@ -0,0 +1,66 @@ +; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s + +; Checks how NVPTX lowers alloca buffers and their passing to functions. +; +; Produced with the following CUDA code: +; extern "C" __attribute__((device)) void callee(float* f, char* buf); +; +; extern "C" __attribute__((global)) void kernel_func(float* a) { +; char buf[4 * sizeof(float)]; +; *(reinterpret_cast<float*>(&buf[0])) = a[0]; +; *(reinterpret_cast<float*>(&buf[1])) = a[1]; +; *(reinterpret_cast<float*>(&buf[2])) = a[2]; +; *(reinterpret_cast<float*>(&buf[3])) = a[3]; +; callee(a, buf); +; } + +; CHECK: .visible .entry kernel_func +define void @kernel_func(float* %a) { +entry: + %buf = alloca [16 x i8], align 4 + +; CHECK: .local .align 4 .b8 __local_depot0[16] +; CHECK: mov.u64 %rd[[BUF_REG:[0-9]+]] +; CHECK: cvta.local.u64 %SP, %rd[[BUF_REG]] + +; CHECK: ld.param.u64 %rd[[A_REG:[0-9]+]], [kernel_func_param_0] +; CHECK: ld.f32 %f[[A0_REG:[0-9]+]], [%rd[[A_REG]]] +; CHECK: st.f32 [%SP+0], %f[[A0_REG]] + + %0 = load float* %a, align 4 + %1 = bitcast [16 x i8]* %buf to float* + store float %0, float* %1, align 4 + %arrayidx2 = getelementptr inbounds float* %a, i64 1 + %2 = load float* %arrayidx2, align 4 + %arrayidx3 = getelementptr inbounds [16 x i8]* %buf, i64 0, i64 1 + %3 = bitcast i8* %arrayidx3 to float* + store float %2, float* %3, align 4 + %arrayidx4 = getelementptr inbounds float* %a, i64 2 + %4 = load float* %arrayidx4, align 4 + %arrayidx5 = getelementptr inbounds [16 x i8]* %buf, i64 0, i64 2 + %5 = bitcast i8* %arrayidx5 to float* + store float %4, float* %5, align 4 + %arrayidx6 = getelementptr inbounds float* %a, i64 3 + %6 = load float* %arrayidx6, align 4 + %arrayidx7 = getelementptr inbounds [16 x i8]* %buf, i64 0, i64 3 + %7 = bitcast i8* %arrayidx7 to float* + store float %6, float* %7, align 4 + +; CHECK: add.u64 %rd[[SP_REG:[0-9]+]], %SP, 0 +; CHECK: .param .b64 param0; +; CHECK-NEXT: st.param.b64 [param0+0], %rd[[A_REG]] +; CHECK-NEXT: .param .b64 param1; +; CHECK-NEXT: st.param.b64 [param1+0], %rd[[SP_REG]] +; CHECK-NEXT: call.uni +; CHECK-NEXT: callee, + + %arraydecay = getelementptr inbounds [16 x i8]* %buf, i64 0, i64 0 + call void @callee(float* %a, i8* %arraydecay) #2 + ret void +} + +declare void @callee(float*, i8*) + +!nvvm.annotations = !{!0} + +!0 = metadata !{void (float*)* @kernel_func, metadata !"kernel", i32 1} diff --git a/test/CodeGen/NVPTX/compare-int.ll b/test/CodeGen/NVPTX/compare-int.ll index c595f215f6f1..e4e0601db59f 100644 --- a/test/CodeGen/NVPTX/compare-int.ll +++ b/test/CodeGen/NVPTX/compare-int.ll @@ -9,8 +9,8 @@ ;;; i64 define i64 @icmp_eq_i64(i64 %a, i64 %b) { -; CHECK: setp.eq.s64 %p[[P0:[0-9]+]], %rl{{[0-9]+}}, %rl{{[0-9]+}} -; CHECK: selp.u64 %rl{{[0-9]+}}, 1, 0, %p[[P0]] +; CHECK: setp.eq.s64 %p[[P0:[0-9]+]], %rd{{[0-9]+}}, %rd{{[0-9]+}} +; CHECK: selp.u64 %rd{{[0-9]+}}, 1, 0, %p[[P0]] ; CHECK: ret %cmp = icmp eq i64 %a, %b %ret = zext i1 %cmp to i64 @@ -18,8 +18,8 @@ define i64 @icmp_eq_i64(i64 %a, i64 %b) { } define i64 @icmp_ne_i64(i64 %a, i64 %b) { -; CHECK: setp.ne.s64 %p[[P0:[0-9]+]], %rl{{[0-9]+}}, %rl{{[0-9]+}} -; CHECK: selp.u64 %rl{{[0-9]+}}, 1, 0, %p[[P0]] +; CHECK: setp.ne.s64 %p[[P0:[0-9]+]], %rd{{[0-9]+}}, %rd{{[0-9]+}} +; CHECK: selp.u64 %rd{{[0-9]+}}, 1, 0, %p[[P0]] ; CHECK: ret %cmp = icmp ne i64 %a, %b %ret = zext i1 %cmp to i64 @@ -27,8 +27,8 @@ define i64 @icmp_ne_i64(i64 %a, i64 %b) { } define i64 @icmp_ugt_i64(i64 %a, i64 %b) { -; CHECK: setp.gt.u64 %p[[P0:[0-9]+]], %rl{{[0-9]+}}, %rl{{[0-9]+}} -; CHECK: selp.u64 %rl{{[0-9]+}}, 1, 0, %p[[P0]] +; CHECK: setp.gt.u64 %p[[P0:[0-9]+]], %rd{{[0-9]+}}, %rd{{[0-9]+}} +; CHECK: selp.u64 %rd{{[0-9]+}}, 1, 0, %p[[P0]] ; CHECK: ret %cmp = icmp ugt i64 %a, %b %ret = zext i1 %cmp to i64 @@ -36,8 +36,8 @@ define i64 @icmp_ugt_i64(i64 %a, i64 %b) { } define i64 @icmp_uge_i64(i64 %a, i64 %b) { -; CHECK: setp.ge.u64 %p[[P0:[0-9]+]], %rl{{[0-9]+}}, %rl{{[0-9]+}} -; CHECK: selp.u64 %rl{{[0-9]+}}, 1, 0, %p[[P0]] +; CHECK: setp.ge.u64 %p[[P0:[0-9]+]], %rd{{[0-9]+}}, %rd{{[0-9]+}} +; CHECK: selp.u64 %rd{{[0-9]+}}, 1, 0, %p[[P0]] ; CHECK: ret %cmp = icmp uge i64 %a, %b %ret = zext i1 %cmp to i64 @@ -45,8 +45,8 @@ define i64 @icmp_uge_i64(i64 %a, i64 %b) { } define i64 @icmp_ult_i64(i64 %a, i64 %b) { -; CHECK: setp.lt.u64 %p[[P0:[0-9]+]], %rl{{[0-9]+}}, %rl{{[0-9]+}} -; CHECK: selp.u64 %rl{{[0-9]+}}, 1, 0, %p[[P0]] +; CHECK: setp.lt.u64 %p[[P0:[0-9]+]], %rd{{[0-9]+}}, %rd{{[0-9]+}} +; CHECK: selp.u64 %rd{{[0-9]+}}, 1, 0, %p[[P0]] ; CHECK: ret %cmp = icmp ult i64 %a, %b %ret = zext i1 %cmp to i64 @@ -54,8 +54,8 @@ define i64 @icmp_ult_i64(i64 %a, i64 %b) { } define i64 @icmp_ule_i64(i64 %a, i64 %b) { -; CHECK: setp.le.u64 %p[[P0:[0-9]+]], %rl{{[0-9]+}}, %rl{{[0-9]+}} -; CHECK: selp.u64 %rl{{[0-9]+}}, 1, 0, %p[[P0]] +; CHECK: setp.le.u64 %p[[P0:[0-9]+]], %rd{{[0-9]+}}, %rd{{[0-9]+}} +; CHECK: selp.u64 %rd{{[0-9]+}}, 1, 0, %p[[P0]] ; CHECK: ret %cmp = icmp ule i64 %a, %b %ret = zext i1 %cmp to i64 @@ -63,8 +63,8 @@ define i64 @icmp_ule_i64(i64 %a, i64 %b) { } define i64 @icmp_sgt_i64(i64 %a, i64 %b) { -; CHECK: setp.gt.s64 %p[[P0:[0-9]+]], %rl{{[0-9]+}}, %rl{{[0-9]+}} -; CHECK: selp.u64 %rl{{[0-9]+}}, 1, 0, %p[[P0]] +; CHECK: setp.gt.s64 %p[[P0:[0-9]+]], %rd{{[0-9]+}}, %rd{{[0-9]+}} +; CHECK: selp.u64 %rd{{[0-9]+}}, 1, 0, %p[[P0]] ; CHECK: ret %cmp = icmp sgt i64 %a, %b %ret = zext i1 %cmp to i64 @@ -72,8 +72,8 @@ define i64 @icmp_sgt_i64(i64 %a, i64 %b) { } define i64 @icmp_sge_i64(i64 %a, i64 %b) { -; CHECK: setp.ge.s64 %p[[P0:[0-9]+]], %rl{{[0-9]+}}, %rl{{[0-9]+}} -; CHECK: selp.u64 %rl{{[0-9]+}}, 1, 0, %p[[P0]] +; CHECK: setp.ge.s64 %p[[P0:[0-9]+]], %rd{{[0-9]+}}, %rd{{[0-9]+}} +; CHECK: selp.u64 %rd{{[0-9]+}}, 1, 0, %p[[P0]] ; CHECK: ret %cmp = icmp sge i64 %a, %b %ret = zext i1 %cmp to i64 @@ -81,8 +81,8 @@ define i64 @icmp_sge_i64(i64 %a, i64 %b) { } define i64 @icmp_slt_i64(i64 %a, i64 %b) { -; CHECK: setp.lt.s64 %p[[P0:[0-9]+]], %rl{{[0-9]+}}, %rl{{[0-9]+}} -; CHECK: selp.u64 %rl{{[0-9]+}}, 1, 0, %p[[P0]] +; CHECK: setp.lt.s64 %p[[P0:[0-9]+]], %rd{{[0-9]+}}, %rd{{[0-9]+}} +; CHECK: selp.u64 %rd{{[0-9]+}}, 1, 0, %p[[P0]] ; CHECK: ret %cmp = icmp slt i64 %a, %b %ret = zext i1 %cmp to i64 @@ -90,8 +90,8 @@ define i64 @icmp_slt_i64(i64 %a, i64 %b) { } define i64 @icmp_sle_i64(i64 %a, i64 %b) { -; CHECK: setp.le.s64 %p[[P0:[0-9]+]], %rl{{[0-9]+}}, %rl{{[0-9]+}} -; CHECK: selp.u64 %rl{{[0-9]+}}, 1, 0, %p[[P0]] +; CHECK: setp.le.s64 %p[[P0:[0-9]+]], %rd{{[0-9]+}}, %rd{{[0-9]+}} +; CHECK: selp.u64 %rd{{[0-9]+}}, 1, 0, %p[[P0]] ; CHECK: ret %cmp = icmp sle i64 %a, %b %ret = zext i1 %cmp to i64 diff --git a/test/CodeGen/NVPTX/convert-fp.ll b/test/CodeGen/NVPTX/convert-fp.ll index 1882121fa724..4b5446e317f4 100644 --- a/test/CodeGen/NVPTX/convert-fp.ll +++ b/test/CodeGen/NVPTX/convert-fp.ll @@ -10,7 +10,7 @@ define i16 @cvt_i16_f32(float %x) { } define i16 @cvt_i16_f64(double %x) { -; CHECK: cvt.rzi.u16.f64 %rs{{[0-9]+}}, %fl{{[0-9]+}}; +; CHECK: cvt.rzi.u16.f64 %rs{{[0-9]+}}, %fd{{[0-9]+}}; ; CHECK: ret; %a = fptoui double %x to i16 ret i16 %a @@ -24,7 +24,7 @@ define i32 @cvt_i32_f32(float %x) { } define i32 @cvt_i32_f64(double %x) { -; CHECK: cvt.rzi.u32.f64 %r{{[0-9]+}}, %fl{{[0-9]+}}; +; CHECK: cvt.rzi.u32.f64 %r{{[0-9]+}}, %fd{{[0-9]+}}; ; CHECK: ret; %a = fptoui double %x to i32 ret i32 %a @@ -32,14 +32,14 @@ define i32 @cvt_i32_f64(double %x) { define i64 @cvt_i64_f32(float %x) { -; CHECK: cvt.rzi.u64.f32 %rl{{[0-9]+}}, %f{{[0-9]+}}; +; CHECK: cvt.rzi.u64.f32 %rd{{[0-9]+}}, %f{{[0-9]+}}; ; CHECK: ret; %a = fptoui float %x to i64 ret i64 %a } define i64 @cvt_i64_f64(double %x) { -; CHECK: cvt.rzi.u64.f64 %rl{{[0-9]+}}, %fl{{[0-9]+}}; +; CHECK: cvt.rzi.u64.f64 %rd{{[0-9]+}}, %fd{{[0-9]+}}; ; CHECK: ret; %a = fptoui double %x to i64 ret i64 %a @@ -60,14 +60,14 @@ define float @cvt_f32_i32(i32 %x) { } define float @cvt_f32_i64(i64 %x) { -; CHECK: cvt.rn.f32.u64 %f{{[0-9]+}}, %rl{{[0-9]+}}; +; CHECK: cvt.rn.f32.u64 %f{{[0-9]+}}, %rd{{[0-9]+}}; ; CHECK: ret; %a = uitofp i64 %x to float ret float %a } define float @cvt_f32_f64(double %x) { -; CHECK: cvt.rn.f32.f64 %f{{[0-9]+}}, %fl{{[0-9]+}}; +; CHECK: cvt.rn.f32.f64 %f{{[0-9]+}}, %fd{{[0-9]+}}; ; CHECK: ret; %a = fptrunc double %x to float ret float %a @@ -88,56 +88,56 @@ define float @cvt_f32_s32(i32 %x) { } define float @cvt_f32_s64(i64 %x) { -; CHECK: cvt.rn.f32.s64 %f{{[0-9]+}}, %rl{{[0-9]+}} +; CHECK: cvt.rn.f32.s64 %f{{[0-9]+}}, %rd{{[0-9]+}} ; CHECK: ret %a = sitofp i64 %x to float ret float %a } define double @cvt_f64_i16(i16 %x) { -; CHECK: cvt.rn.f64.u16 %fl{{[0-9]+}}, %rs{{[0-9]+}}; +; CHECK: cvt.rn.f64.u16 %fd{{[0-9]+}}, %rs{{[0-9]+}}; ; CHECK: ret; %a = uitofp i16 %x to double ret double %a } define double @cvt_f64_i32(i32 %x) { -; CHECK: cvt.rn.f64.u32 %fl{{[0-9]+}}, %r{{[0-9]+}}; +; CHECK: cvt.rn.f64.u32 %fd{{[0-9]+}}, %r{{[0-9]+}}; ; CHECK: ret; %a = uitofp i32 %x to double ret double %a } define double @cvt_f64_i64(i64 %x) { -; CHECK: cvt.rn.f64.u64 %fl{{[0-9]+}}, %rl{{[0-9]+}}; +; CHECK: cvt.rn.f64.u64 %fd{{[0-9]+}}, %rd{{[0-9]+}}; ; CHECK: ret; %a = uitofp i64 %x to double ret double %a } define double @cvt_f64_f32(float %x) { -; CHECK: cvt.f64.f32 %fl{{[0-9]+}}, %f{{[0-9]+}}; +; CHECK: cvt.f64.f32 %fd{{[0-9]+}}, %f{{[0-9]+}}; ; CHECK: ret; %a = fpext float %x to double ret double %a } define double @cvt_f64_s16(i16 %x) { -; CHECK: cvt.rn.f64.s16 %fl{{[0-9]+}}, %rs{{[0-9]+}} +; CHECK: cvt.rn.f64.s16 %fd{{[0-9]+}}, %rs{{[0-9]+}} ; CHECK: ret %a = sitofp i16 %x to double ret double %a } define double @cvt_f64_s32(i32 %x) { -; CHECK: cvt.rn.f64.s32 %fl{{[0-9]+}}, %r{{[0-9]+}} +; CHECK: cvt.rn.f64.s32 %fd{{[0-9]+}}, %r{{[0-9]+}} ; CHECK: ret %a = sitofp i32 %x to double ret double %a } define double @cvt_f64_s64(i64 %x) { -; CHECK: cvt.rn.f64.s64 %fl{{[0-9]+}}, %rl{{[0-9]+}} +; CHECK: cvt.rn.f64.s64 %fd{{[0-9]+}}, %rd{{[0-9]+}} ; CHECK: ret %a = sitofp i64 %x to double ret double %a diff --git a/test/CodeGen/NVPTX/convert-int-sm20.ll b/test/CodeGen/NVPTX/convert-int-sm20.ll index 227cd31e11b3..57a231629e00 100644 --- a/test/CodeGen/NVPTX/convert-int-sm20.ll +++ b/test/CodeGen/NVPTX/convert-int-sm20.ll @@ -48,16 +48,16 @@ define i32 @cvt_i32_i64(i64 %x) { ; i64 define i64 @cvt_i64_i16(i16 %x) { -; CHECK: ld.param.u16 %rl[[R0:[0-9]+]], [cvt_i64_i16_param_{{[0-9]+}}] -; CHECK: st.param.b64 [func_retval{{[0-9]+}}+0], %rl[[R0]] +; CHECK: ld.param.u16 %rd[[R0:[0-9]+]], [cvt_i64_i16_param_{{[0-9]+}}] +; CHECK: st.param.b64 [func_retval{{[0-9]+}}+0], %rd[[R0]] ; CHECK: ret %a = zext i16 %x to i64 ret i64 %a } define i64 @cvt_i64_i32(i32 %x) { -; CHECK: ld.param.u32 %rl[[R0:[0-9]+]], [cvt_i64_i32_param_{{[0-9]+}}] -; CHECK: st.param.b64 [func_retval{{[0-9]+}}+0], %rl[[R0]] +; CHECK: ld.param.u32 %rd[[R0:[0-9]+]], [cvt_i64_i32_param_{{[0-9]+}}] +; CHECK: st.param.b64 [func_retval{{[0-9]+}}+0], %rd[[R0]] ; CHECK: ret %a = zext i32 %x to i64 ret i64 %a diff --git a/test/CodeGen/NVPTX/div-ri.ll b/test/CodeGen/NVPTX/div-ri.ll new file mode 100644 index 000000000000..7f796e0239fc --- /dev/null +++ b/test/CodeGen/NVPTX/div-ri.ll @@ -0,0 +1,8 @@ +; RUN: llc < %s -march=nvptx -mcpu=sm_20 -nvptx-prec-divf32=0 | FileCheck %s + +define float @foo(float %a) { +; CHECK: div.approx.f32 + %div = fdiv float %a, 13.0 + ret float %div +} + diff --git a/test/CodeGen/NVPTX/envreg.ll b/test/CodeGen/NVPTX/envreg.ll new file mode 100644 index 000000000000..a341b49ecdf3 --- /dev/null +++ b/test/CodeGen/NVPTX/envreg.ll @@ -0,0 +1,139 @@ +; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s + + +declare i32 @llvm.nvvm.read.ptx.sreg.envreg0() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg1() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg2() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg3() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg4() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg5() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg6() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg7() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg8() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg9() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg10() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg11() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg12() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg13() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg14() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg15() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg16() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg17() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg18() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg19() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg20() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg21() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg22() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg23() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg24() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg25() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg26() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg27() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg28() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg29() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg30() +declare i32 @llvm.nvvm.read.ptx.sreg.envreg31() + + +; CHECK: foo +define i32 @foo() { +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg0 + %val0 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg0() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg1 + %val1 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg1() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg2 + %val2 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg2() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg3 + %val3 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg3() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg4 + %val4 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg4() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg5 + %val5 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg5() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg6 + %val6 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg6() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg7 + %val7 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg7() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg8 + %val8 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg8() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg9 + %val9 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg9() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg10 + %val10 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg10() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg11 + %val11 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg11() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg12 + %val12 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg12() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg13 + %val13 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg13() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg14 + %val14 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg14() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg15 + %val15 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg15() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg16 + %val16 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg16() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg17 + %val17 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg17() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg18 + %val18 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg18() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg19 + %val19 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg19() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg20 + %val20 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg20() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg21 + %val21 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg21() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg22 + %val22 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg22() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg23 + %val23 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg23() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg24 + %val24 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg24() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg25 + %val25 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg25() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg26 + %val26 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg26() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg27 + %val27 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg27() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg28 + %val28 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg28() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg29 + %val29 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg29() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg30 + %val30 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg30() +; CHECK: mov.b32 %r{{[0-9]+}}, %envreg31 + %val31 = tail call i32 @llvm.nvvm.read.ptx.sreg.envreg31() + + + %ret0 = add i32 %val0, %val1 + %ret1 = add i32 %ret0, %val2 + %ret2 = add i32 %ret1, %val3 + %ret3 = add i32 %ret2, %val4 + %ret4 = add i32 %ret3, %val5 + %ret5 = add i32 %ret4, %val6 + %ret6 = add i32 %ret5, %val7 + %ret7 = add i32 %ret6, %val8 + %ret8 = add i32 %ret7, %val9 + %ret9 = add i32 %ret8, %val10 + %ret10 = add i32 %ret9, %val11 + %ret11 = add i32 %ret10, %val12 + %ret12 = add i32 %ret11, %val13 + %ret13 = add i32 %ret12, %val14 + %ret14 = add i32 %ret13, %val15 + %ret15 = add i32 %ret14, %val16 + %ret16 = add i32 %ret15, %val17 + %ret17 = add i32 %ret16, %val18 + %ret18 = add i32 %ret17, %val19 + %ret19 = add i32 %ret18, %val20 + %ret20 = add i32 %ret19, %val21 + %ret21 = add i32 %ret20, %val22 + %ret22 = add i32 %ret21, %val23 + %ret23 = add i32 %ret22, %val24 + %ret24 = add i32 %ret23, %val25 + %ret25 = add i32 %ret24, %val26 + %ret26 = add i32 %ret25, %val27 + %ret27 = add i32 %ret26, %val28 + %ret28 = add i32 %ret27, %val29 + %ret29 = add i32 %ret28, %val30 + %ret30 = add i32 %ret29, %val31 + + ret i32 %ret30 +} diff --git a/test/CodeGen/NVPTX/fma.ll b/test/CodeGen/NVPTX/fma.ll index 4ef1a9a4cefb..14b5c45b87d8 100644 --- a/test/CodeGen/NVPTX/fma.ll +++ b/test/CodeGen/NVPTX/fma.ll @@ -1,4 +1,4 @@ -; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s +; RUN: llc < %s -march=nvptx -mcpu=sm_20 -fp-contract=fast | FileCheck %s define ptx_device float @t1_f32(float %x, float %y, float %z) { ; CHECK: fma.rn.f32 %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}; @@ -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: fma.rn.f64 %fl{{[0-9]+}}, %fl{{[0-9]+}}, %fl{{[0-9]+}}, %fl{{[0-9]+}}; +; CHECK: fma.rn.f64 %fd{{[0-9]+}}, %fd{{[0-9]+}}, %fd{{[0-9]+}}, %fd{{[0-9]+}}; ; CHECK: ret; %a = fmul double %x, %y %b = fadd double %a, %z diff --git a/test/CodeGen/NVPTX/fp-contract.ll b/test/CodeGen/NVPTX/fp-contract.ll new file mode 100644 index 000000000000..3f68b188ba75 --- /dev/null +++ b/test/CodeGen/NVPTX/fp-contract.ll @@ -0,0 +1,33 @@ +; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -fp-contract=fast | FileCheck %s --check-prefix=FAST +; RUN: llc < %s -march=nvptx64 -mcpu=sm_30 | FileCheck %s --check-prefix=DEFAULT + +target triple = "nvptx64-unknown-cuda" + +;; Make sure we are generating proper instruction sequences for fused ops +;; If fusion is allowed, we try to form fma.rn at the PTX level, and emit +;; add.f32 otherwise. Without an explicit rounding mode on add.f32, ptxas +;; is free to fuse with a multiply if it is able. If fusion is not allowed, +;; we do not form fma.rn at the PTX level and explicitly generate add.rn +;; for all adds to prevent ptxas from fusion the ops. + +;; FAST-LABEL: @t0 +;; DEFAULT-LABEL: @t0 +define float @t0(float %a, float %b, float %c) { +;; FAST: fma.rn.f32 +;; DEFAULT: mul.rn.f32 +;; DEFAULT: add.rn.f32 + %v0 = fmul float %a, %b + %v1 = fadd float %v0, %c + ret float %v1 +} + +;; FAST-LABEL: @t1 +;; DEFAULT-LABEL: @t1 +define float @t1(float %a, float %b) { +;; We cannot form an fma here, but make sure we explicitly emit add.rn.f32 +;; to prevent ptxas from fusing this with anything else. +;; FAST: add.f32 +;; DEFAULT: add.rn.f32 + %v1 = fadd float %a, %b + ret float %v1 +} diff --git a/test/CodeGen/NVPTX/fp-literals.ll b/test/CodeGen/NVPTX/fp-literals.ll index 0cc2413e009f..755e0f9250a1 100644 --- a/test/CodeGen/NVPTX/fp-literals.ll +++ b/test/CodeGen/NVPTX/fp-literals.ll @@ -1,4 +1,7 @@ -; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s +; RUN: llc < %s -march=nvptx -mcpu=sm_20 -fp-contract=fast | FileCheck %s + +target triple = "nvptx64-unknown-cuda" +target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64" ; Make sure we can properly differentiate between single-precision and ; double-precision FP literals. @@ -11,7 +14,7 @@ define float @myaddf(float %a) { } ; CHECK: myaddd -; CHECK: add.f64 %fl{{[0-9]+}}, %fl{{[0-9]+}}, 0d3FF0000000000000 +; CHECK: add.f64 %fd{{[0-9]+}}, %fd{{[0-9]+}}, 0d3FF0000000000000 define double @myaddd(double %a) { %ret = fadd double %a, 1.0 ret double %ret diff --git a/test/CodeGen/NVPTX/fp16.ll b/test/CodeGen/NVPTX/fp16.ll new file mode 100644 index 000000000000..8770399f2ec9 --- /dev/null +++ b/test/CodeGen/NVPTX/fp16.ll @@ -0,0 +1,45 @@ +; RUN: llc -march=nvptx -verify-machineinstrs < %s | FileCheck %s + +declare float @llvm.convert.from.fp16.f32(i16) nounwind readnone +declare double @llvm.convert.from.fp16.f64(i16) nounwind readnone +declare i16 @llvm.convert.to.fp16.f32(float) nounwind readnone +declare i16 @llvm.convert.to.fp16.f64(double) nounwind readnone + +; CHECK-LABEL: @test_convert_fp16_to_fp32 +; CHECK: cvt.f32.f16 +define void @test_convert_fp16_to_fp32(float addrspace(1)* noalias %out, i16 addrspace(1)* noalias %in) nounwind { + %val = load i16 addrspace(1)* %in, align 2 + %cvt = call float @llvm.convert.from.fp16.f32(i16 %val) nounwind readnone + store float %cvt, float addrspace(1)* %out, align 4 + ret void +} + + +; CHECK-LABEL: @test_convert_fp16_to_fp64 +; CHECK: cvt.f64.f16 +define void @test_convert_fp16_to_fp64(double addrspace(1)* noalias %out, i16 addrspace(1)* noalias %in) nounwind { + %val = load i16 addrspace(1)* %in, align 2 + %cvt = call double @llvm.convert.from.fp16.f64(i16 %val) nounwind readnone + store double %cvt, double addrspace(1)* %out, align 4 + ret void +} + + +; CHECK-LABEL: @test_convert_fp32_to_fp16 +; CHECK: cvt.rn.f16.f32 +define void @test_convert_fp32_to_fp16(i16 addrspace(1)* noalias %out, float addrspace(1)* noalias %in) nounwind { + %val = load float addrspace(1)* %in, align 2 + %cvt = call i16 @llvm.convert.to.fp16.f32(float %val) nounwind readnone + store i16 %cvt, i16 addrspace(1)* %out, align 4 + ret void +} + + +; CHECK-LABEL: @test_convert_fp64_to_fp16 +; CHECK: cvt.rn.f16.f64 +define void @test_convert_fp64_to_fp16(i16 addrspace(1)* noalias %out, double addrspace(1)* noalias %in) nounwind { + %val = load double addrspace(1)* %in, align 2 + %cvt = call i16 @llvm.convert.to.fp16.f64(double %val) nounwind readnone + store i16 %cvt, i16 addrspace(1)* %out, align 4 + ret void +} diff --git a/test/CodeGen/NVPTX/gvar-init.ll b/test/CodeGen/NVPTX/gvar-init.ll new file mode 100644 index 000000000000..8c959422e66a --- /dev/null +++ b/test/CodeGen/NVPTX/gvar-init.ll @@ -0,0 +1,5 @@ +; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s + +; Error out if initializer is given for address spaces that do not support initializers +; XFAIL: * +@g0 = addrspace(3) global i32 42 diff --git a/test/CodeGen/NVPTX/half.ll b/test/CodeGen/NVPTX/half.ll new file mode 100644 index 000000000000..aa08cc78e91a --- /dev/null +++ b/test/CodeGen/NVPTX/half.ll @@ -0,0 +1,70 @@ +; RUN: llc < %s -march=nvptx | FileCheck %s + +define void @test_load_store(half addrspace(1)* %in, half addrspace(1)* %out) { +; CHECK-LABEL: @test_load_store +; CHECK: ld.global.u16 [[TMP:%rs[0-9]+]], [{{%r[0-9]+}}] +; CHECK: st.global.u16 [{{%r[0-9]+}}], [[TMP]] + %val = load half addrspace(1)* %in + store half %val, half addrspace(1) * %out + ret void +} + +define void @test_bitcast_from_half(half addrspace(1)* %in, i16 addrspace(1)* %out) { +; CHECK-LABEL: @test_bitcast_from_half +; CHECK: ld.global.u16 [[TMP:%rs[0-9]+]], [{{%r[0-9]+}}] +; CHECK: st.global.u16 [{{%r[0-9]+}}], [[TMP]] + %val = load half addrspace(1) * %in + %val_int = bitcast half %val to i16 + store i16 %val_int, i16 addrspace(1)* %out + ret void +} + +define void @test_bitcast_to_half(half addrspace(1)* %out, i16 addrspace(1)* %in) { +; CHECK-LABEL: @test_bitcast_to_half +; CHECK: ld.global.u16 [[TMP:%rs[0-9]+]], [{{%r[0-9]+}}] +; CHECK: st.global.u16 [{{%r[0-9]+}}], [[TMP]] + %val = load i16 addrspace(1)* %in + %val_fp = bitcast i16 %val to half + store half %val_fp, half addrspace(1)* %out + ret void +} + +define void @test_extend32(half addrspace(1)* %in, float addrspace(1)* %out) { +; CHECK-LABEL: @test_extend32 +; CHECK: cvt.f32.f16 + + %val16 = load half addrspace(1)* %in + %val32 = fpext half %val16 to float + store float %val32, float addrspace(1)* %out + ret void +} + +define void @test_extend64(half addrspace(1)* %in, double addrspace(1)* %out) { +; CHECK-LABEL: @test_extend64 +; CHECK: cvt.f64.f16 + + %val16 = load half addrspace(1)* %in + %val64 = fpext half %val16 to double + store double %val64, double addrspace(1)* %out + ret void +} + +define void @test_trunc32(float addrspace(1)* %in, half addrspace(1)* %out) { +; CHECK-LABEL: test_trunc32 +; CHECK: cvt.rn.f16.f32 + + %val32 = load float addrspace(1)* %in + %val16 = fptrunc float %val32 to half + store half %val16, half addrspace(1)* %out + ret void +} + +define void @test_trunc64(double addrspace(1)* %in, half addrspace(1)* %out) { +; CHECK-LABEL: @test_trunc64 +; CHECK: cvt.rn.f16.f64 + + %val32 = load double addrspace(1)* %in + %val16 = fptrunc double %val32 to half + store half %val16, half addrspace(1)* %out + ret void +} diff --git a/test/CodeGen/NVPTX/imad.ll b/test/CodeGen/NVPTX/imad.ll new file mode 100644 index 000000000000..67421c7cac4b --- /dev/null +++ b/test/CodeGen/NVPTX/imad.ll @@ -0,0 +1,9 @@ +; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s + +; CHECK: imad +define i32 @imad(i32 %a, i32 %b, i32 %c) { +; CHECK: mad.lo.s32 + %val0 = mul i32 %a, %b + %val1 = add i32 %val0, %c + ret i32 %val1 +} diff --git a/test/CodeGen/NVPTX/implicit-def.ll b/test/CodeGen/NVPTX/implicit-def.ll index 06d3d562046e..2d2c6e527f6d 100644 --- a/test/CodeGen/NVPTX/implicit-def.ll +++ b/test/CodeGen/NVPTX/implicit-def.ll @@ -1,7 +1,7 @@ ; RUN: llc < %s -O0 -march=nvptx -mcpu=sm_20 -asm-verbose=1 | FileCheck %s ; CHECK: // implicit-def: %f[[F0:[0-9]+]] -; CHECK: add.f32 %f{{[0-9]+}}, %f{{[0-9]+}}, %f[[F0]]; +; CHECK: add.rn.f32 %f{{[0-9]+}}, %f{{[0-9]+}}, %f[[F0]]; define float @foo(float %a) { %ret = fadd float %a, undef ret float %ret diff --git a/test/CodeGen/NVPTX/inline-asm.ll b/test/CodeGen/NVPTX/inline-asm.ll index d76eb4239ee3..6f0578d4cff4 100644 --- a/test/CodeGen/NVPTX/inline-asm.ll +++ b/test/CodeGen/NVPTX/inline-asm.ll @@ -7,3 +7,10 @@ entry: %0 = call float asm "ex2.approx.ftz.f32 $0, $1;", "=f,f"(float %x) ret float %0 } + +define i32 @foo(i1 signext %cond, i32 %a, i32 %b) #0 { +entry: +; CHECK: selp.b32 %r{{[0-9]+}}, %r{{[0-9]+}}, %r{{[0-9]+}}, %p{{[0-9]+}} + %0 = tail call i32 asm "selp.b32 $0, $1, $2, $3;", "=r,r,r,b"(i32 %a, i32 %b, i1 %cond) + ret i32 %0 +} diff --git a/test/CodeGen/NVPTX/intrinsic-old.ll b/test/CodeGen/NVPTX/intrinsic-old.ll index af91bb442412..3c51776c0ec9 100644 --- a/test/CodeGen/NVPTX/intrinsic-old.ll +++ b/test/CodeGen/NVPTX/intrinsic-old.ll @@ -198,7 +198,7 @@ define ptx_device i32 @test_clock() { } define ptx_device i64 @test_clock64() { -; CHECK: mov.u64 %rl{{[0-9]+}}, %clock64; +; CHECK: mov.u64 %rd{{[0-9]+}}, %clock64; ; CHECK: ret; %x = call i64 @llvm.ptx.read.clock64() ret i64 %x diff --git a/test/CodeGen/NVPTX/intrinsics.ll b/test/CodeGen/NVPTX/intrinsics.ll index 78e1e7789014..34b671d70e94 100644 --- a/test/CodeGen/NVPTX/intrinsics.ll +++ b/test/CodeGen/NVPTX/intrinsics.ll @@ -9,7 +9,7 @@ define ptx_device float @test_fabsf(float %f) { } define ptx_device double @test_fabs(double %d) { -; CHECK: abs.f64 %fl{{[0-9]+}}, %fl{{[0-9]+}}; +; CHECK: abs.f64 %fd{{[0-9]+}}, %fd{{[0-9]+}}; ; CHECK: ret; %x = call double @llvm.fabs.f64(double %d) ret double %x diff --git a/test/CodeGen/NVPTX/isspacep.ll b/test/CodeGen/NVPTX/isspacep.ll new file mode 100644 index 000000000000..47fa7a6714df --- /dev/null +++ b/test/CodeGen/NVPTX/isspacep.ll @@ -0,0 +1,35 @@ +; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s + +declare i1 @llvm.nvvm.isspacep.const(i8*) readnone noinline +declare i1 @llvm.nvvm.isspacep.global(i8*) readnone noinline +declare i1 @llvm.nvvm.isspacep.local(i8*) readnone noinline +declare i1 @llvm.nvvm.isspacep.shared(i8*) readnone noinline + +; CHECK: is_const +define i1 @is_const(i8* %addr) { +; CHECK: isspacep.const + %v = tail call i1 @llvm.nvvm.isspacep.const(i8* %addr) + ret i1 %v +} + +; CHECK: is_global +define i1 @is_global(i8* %addr) { +; CHECK: isspacep.global + %v = tail call i1 @llvm.nvvm.isspacep.global(i8* %addr) + ret i1 %v +} + +; CHECK: is_local +define i1 @is_local(i8* %addr) { +; CHECK: isspacep.local + %v = tail call i1 @llvm.nvvm.isspacep.local(i8* %addr) + ret i1 %v +} + +; CHECK: is_shared +define i1 @is_shared(i8* %addr) { +; CHECK: isspacep.shared + %v = tail call i1 @llvm.nvvm.isspacep.shared(i8* %addr) + ret i1 %v +} + diff --git a/test/CodeGen/NVPTX/ld-addrspace.ll b/test/CodeGen/NVPTX/ld-addrspace.ll index 133ef09afdb2..f33659c92e84 100644 --- a/test/CodeGen/NVPTX/ld-addrspace.ll +++ b/test/CodeGen/NVPTX/ld-addrspace.ll @@ -6,7 +6,7 @@ define i8 @ld_global_i8(i8 addrspace(1)* %ptr) { ; PTX32: ld.global.u8 %r{{[0-9]+}}, [%r{{[0-9]+}}] ; PTX32: ret -; PTX64: ld.global.u8 %r{{[0-9]+}}, [%rl{{[0-9]+}}] +; PTX64: ld.global.u8 %r{{[0-9]+}}, [%rd{{[0-9]+}}] ; PTX64: ret %a = load i8 addrspace(1)* %ptr ret i8 %a @@ -15,7 +15,7 @@ define i8 @ld_global_i8(i8 addrspace(1)* %ptr) { define i8 @ld_shared_i8(i8 addrspace(3)* %ptr) { ; PTX32: ld.shared.u8 %r{{[0-9]+}}, [%r{{[0-9]+}}] ; PTX32: ret -; PTX64: ld.shared.u8 %r{{[0-9]+}}, [%rl{{[0-9]+}}] +; PTX64: ld.shared.u8 %r{{[0-9]+}}, [%rd{{[0-9]+}}] ; PTX64: ret %a = load i8 addrspace(3)* %ptr ret i8 %a @@ -24,7 +24,7 @@ define i8 @ld_shared_i8(i8 addrspace(3)* %ptr) { define i8 @ld_local_i8(i8 addrspace(5)* %ptr) { ; PTX32: ld.local.u8 %r{{[0-9]+}}, [%r{{[0-9]+}}] ; PTX32: ret -; PTX64: ld.local.u8 %r{{[0-9]+}}, [%rl{{[0-9]+}}] +; PTX64: ld.local.u8 %r{{[0-9]+}}, [%rd{{[0-9]+}}] ; PTX64: ret %a = load i8 addrspace(5)* %ptr ret i8 %a @@ -34,7 +34,7 @@ define i8 @ld_local_i8(i8 addrspace(5)* %ptr) { define i16 @ld_global_i16(i16 addrspace(1)* %ptr) { ; PTX32: ld.global.u16 %r{{[0-9]+}}, [%r{{[0-9]+}}] ; PTX32: ret -; PTX64: ld.global.u16 %r{{[0-9]+}}, [%rl{{[0-9]+}}] +; PTX64: ld.global.u16 %r{{[0-9]+}}, [%rd{{[0-9]+}}] ; PTX64: ret %a = load i16 addrspace(1)* %ptr ret i16 %a @@ -43,7 +43,7 @@ define i16 @ld_global_i16(i16 addrspace(1)* %ptr) { define i16 @ld_shared_i16(i16 addrspace(3)* %ptr) { ; PTX32: ld.shared.u16 %r{{[0-9]+}}, [%r{{[0-9]+}}] ; PTX32: ret -; PTX64: ld.shared.u16 %r{{[0-9]+}}, [%rl{{[0-9]+}}] +; PTX64: ld.shared.u16 %r{{[0-9]+}}, [%rd{{[0-9]+}}] ; PTX64: ret %a = load i16 addrspace(3)* %ptr ret i16 %a @@ -52,7 +52,7 @@ define i16 @ld_shared_i16(i16 addrspace(3)* %ptr) { define i16 @ld_local_i16(i16 addrspace(5)* %ptr) { ; PTX32: ld.local.u16 %r{{[0-9]+}}, [%r{{[0-9]+}}] ; PTX32: ret -; PTX64: ld.local.u16 %r{{[0-9]+}}, [%rl{{[0-9]+}}] +; PTX64: ld.local.u16 %r{{[0-9]+}}, [%rd{{[0-9]+}}] ; PTX64: ret %a = load i16 addrspace(5)* %ptr ret i16 %a @@ -62,7 +62,7 @@ define i16 @ld_local_i16(i16 addrspace(5)* %ptr) { define i32 @ld_global_i32(i32 addrspace(1)* %ptr) { ; PTX32: ld.global.u32 %r{{[0-9]+}}, [%r{{[0-9]+}}] ; PTX32: ret -; PTX64: ld.global.u32 %r{{[0-9]+}}, [%rl{{[0-9]+}}] +; PTX64: ld.global.u32 %r{{[0-9]+}}, [%rd{{[0-9]+}}] ; PTX64: ret %a = load i32 addrspace(1)* %ptr ret i32 %a @@ -71,7 +71,7 @@ define i32 @ld_global_i32(i32 addrspace(1)* %ptr) { define i32 @ld_shared_i32(i32 addrspace(3)* %ptr) { ; PTX32: ld.shared.u32 %r{{[0-9]+}}, [%r{{[0-9]+}}] ; PTX32: ret -; PTX64: ld.shared.u32 %r{{[0-9]+}}, [%rl{{[0-9]+}}] +; PTX64: ld.shared.u32 %r{{[0-9]+}}, [%rd{{[0-9]+}}] ; PTX64: ret %a = load i32 addrspace(3)* %ptr ret i32 %a @@ -80,7 +80,7 @@ define i32 @ld_shared_i32(i32 addrspace(3)* %ptr) { define i32 @ld_local_i32(i32 addrspace(5)* %ptr) { ; PTX32: ld.local.u32 %r{{[0-9]+}}, [%r{{[0-9]+}}] ; PTX32: ret -; PTX64: ld.local.u32 %r{{[0-9]+}}, [%rl{{[0-9]+}}] +; PTX64: ld.local.u32 %r{{[0-9]+}}, [%rd{{[0-9]+}}] ; PTX64: ret %a = load i32 addrspace(5)* %ptr ret i32 %a @@ -88,27 +88,27 @@ define i32 @ld_local_i32(i32 addrspace(5)* %ptr) { ;; i64 define i64 @ld_global_i64(i64 addrspace(1)* %ptr) { -; PTX32: ld.global.u64 %rl{{[0-9]+}}, [%r{{[0-9]+}}] +; PTX32: ld.global.u64 %rd{{[0-9]+}}, [%r{{[0-9]+}}] ; PTX32: ret -; PTX64: ld.global.u64 %rl{{[0-9]+}}, [%rl{{[0-9]+}}] +; PTX64: ld.global.u64 %rd{{[0-9]+}}, [%rd{{[0-9]+}}] ; PTX64: ret %a = load i64 addrspace(1)* %ptr ret i64 %a } define i64 @ld_shared_i64(i64 addrspace(3)* %ptr) { -; PTX32: ld.shared.u64 %rl{{[0-9]+}}, [%r{{[0-9]+}}] +; PTX32: ld.shared.u64 %rd{{[0-9]+}}, [%r{{[0-9]+}}] ; PTX32: ret -; PTX64: ld.shared.u64 %rl{{[0-9]+}}, [%rl{{[0-9]+}}] +; PTX64: ld.shared.u64 %rd{{[0-9]+}}, [%rd{{[0-9]+}}] ; PTX64: ret %a = load i64 addrspace(3)* %ptr ret i64 %a } define i64 @ld_local_i64(i64 addrspace(5)* %ptr) { -; PTX32: ld.local.u64 %rl{{[0-9]+}}, [%r{{[0-9]+}}] +; PTX32: ld.local.u64 %rd{{[0-9]+}}, [%r{{[0-9]+}}] ; PTX32: ret -; PTX64: ld.local.u64 %rl{{[0-9]+}}, [%rl{{[0-9]+}}] +; PTX64: ld.local.u64 %rd{{[0-9]+}}, [%rd{{[0-9]+}}] ; PTX64: ret %a = load i64 addrspace(5)* %ptr ret i64 %a @@ -118,7 +118,7 @@ define i64 @ld_local_i64(i64 addrspace(5)* %ptr) { define float @ld_global_f32(float addrspace(1)* %ptr) { ; PTX32: ld.global.f32 %f{{[0-9]+}}, [%r{{[0-9]+}}] ; PTX32: ret -; PTX64: ld.global.f32 %f{{[0-9]+}}, [%rl{{[0-9]+}}] +; PTX64: ld.global.f32 %f{{[0-9]+}}, [%rd{{[0-9]+}}] ; PTX64: ret %a = load float addrspace(1)* %ptr ret float %a @@ -127,7 +127,7 @@ define float @ld_global_f32(float addrspace(1)* %ptr) { define float @ld_shared_f32(float addrspace(3)* %ptr) { ; PTX32: ld.shared.f32 %f{{[0-9]+}}, [%r{{[0-9]+}}] ; PTX32: ret -; PTX64: ld.shared.f32 %f{{[0-9]+}}, [%rl{{[0-9]+}}] +; PTX64: ld.shared.f32 %f{{[0-9]+}}, [%rd{{[0-9]+}}] ; PTX64: ret %a = load float addrspace(3)* %ptr ret float %a @@ -136,7 +136,7 @@ define float @ld_shared_f32(float addrspace(3)* %ptr) { define float @ld_local_f32(float addrspace(5)* %ptr) { ; PTX32: ld.local.f32 %f{{[0-9]+}}, [%r{{[0-9]+}}] ; PTX32: ret -; PTX64: ld.local.f32 %f{{[0-9]+}}, [%rl{{[0-9]+}}] +; PTX64: ld.local.f32 %f{{[0-9]+}}, [%rd{{[0-9]+}}] ; PTX64: ret %a = load float addrspace(5)* %ptr ret float %a @@ -144,27 +144,27 @@ define float @ld_local_f32(float addrspace(5)* %ptr) { ;; f64 define double @ld_global_f64(double addrspace(1)* %ptr) { -; PTX32: ld.global.f64 %fl{{[0-9]+}}, [%r{{[0-9]+}}] +; PTX32: ld.global.f64 %fd{{[0-9]+}}, [%r{{[0-9]+}}] ; PTX32: ret -; PTX64: ld.global.f64 %fl{{[0-9]+}}, [%rl{{[0-9]+}}] +; PTX64: ld.global.f64 %fd{{[0-9]+}}, [%rd{{[0-9]+}}] ; PTX64: ret %a = load double addrspace(1)* %ptr ret double %a } define double @ld_shared_f64(double addrspace(3)* %ptr) { -; PTX32: ld.shared.f64 %fl{{[0-9]+}}, [%r{{[0-9]+}}] +; PTX32: ld.shared.f64 %fd{{[0-9]+}}, [%r{{[0-9]+}}] ; PTX32: ret -; PTX64: ld.shared.f64 %fl{{[0-9]+}}, [%rl{{[0-9]+}}] +; PTX64: ld.shared.f64 %fd{{[0-9]+}}, [%rd{{[0-9]+}}] ; PTX64: ret %a = load double addrspace(3)* %ptr ret double %a } define double @ld_local_f64(double addrspace(5)* %ptr) { -; PTX32: ld.local.f64 %fl{{[0-9]+}}, [%r{{[0-9]+}}] +; PTX32: ld.local.f64 %fd{{[0-9]+}}, [%r{{[0-9]+}}] ; PTX32: ret -; PTX64: ld.local.f64 %fl{{[0-9]+}}, [%rl{{[0-9]+}}] +; PTX64: ld.local.f64 %fd{{[0-9]+}}, [%rd{{[0-9]+}}] ; PTX64: ret %a = load double addrspace(5)* %ptr ret double %a diff --git a/test/CodeGen/NVPTX/ld-generic.ll b/test/CodeGen/NVPTX/ld-generic.ll index 3728268c24d5..d629e0ecc647 100644 --- a/test/CodeGen/NVPTX/ld-generic.ll +++ b/test/CodeGen/NVPTX/ld-generic.ll @@ -6,7 +6,7 @@ define i8 @ld_global_i8(i8 addrspace(0)* %ptr) { ; PTX32: ld.u8 %r{{[0-9]+}}, [%r{{[0-9]+}}] ; PTX32: ret -; PTX64: ld.u8 %r{{[0-9]+}}, [%rl{{[0-9]+}}] +; PTX64: ld.u8 %r{{[0-9]+}}, [%rd{{[0-9]+}}] ; PTX64: ret %a = load i8 addrspace(0)* %ptr ret i8 %a @@ -16,7 +16,7 @@ define i8 @ld_global_i8(i8 addrspace(0)* %ptr) { define i16 @ld_global_i16(i16 addrspace(0)* %ptr) { ; PTX32: ld.u16 %r{{[0-9]+}}, [%r{{[0-9]+}}] ; PTX32: ret -; PTX64: ld.u16 %r{{[0-9]+}}, [%rl{{[0-9]+}}] +; PTX64: ld.u16 %r{{[0-9]+}}, [%rd{{[0-9]+}}] ; PTX64: ret %a = load i16 addrspace(0)* %ptr ret i16 %a @@ -26,7 +26,7 @@ define i16 @ld_global_i16(i16 addrspace(0)* %ptr) { define i32 @ld_global_i32(i32 addrspace(0)* %ptr) { ; PTX32: ld.u32 %r{{[0-9]+}}, [%r{{[0-9]+}}] ; PTX32: ret -; PTX64: ld.u32 %r{{[0-9]+}}, [%rl{{[0-9]+}}] +; PTX64: ld.u32 %r{{[0-9]+}}, [%rd{{[0-9]+}}] ; PTX64: ret %a = load i32 addrspace(0)* %ptr ret i32 %a @@ -34,9 +34,9 @@ define i32 @ld_global_i32(i32 addrspace(0)* %ptr) { ;; i64 define i64 @ld_global_i64(i64 addrspace(0)* %ptr) { -; PTX32: ld.u64 %rl{{[0-9]+}}, [%r{{[0-9]+}}] +; PTX32: ld.u64 %rd{{[0-9]+}}, [%r{{[0-9]+}}] ; PTX32: ret -; PTX64: ld.u64 %rl{{[0-9]+}}, [%rl{{[0-9]+}}] +; PTX64: ld.u64 %rd{{[0-9]+}}, [%rd{{[0-9]+}}] ; PTX64: ret %a = load i64 addrspace(0)* %ptr ret i64 %a @@ -46,7 +46,7 @@ define i64 @ld_global_i64(i64 addrspace(0)* %ptr) { define float @ld_global_f32(float addrspace(0)* %ptr) { ; PTX32: ld.f32 %f{{[0-9]+}}, [%r{{[0-9]+}}] ; PTX32: ret -; PTX64: ld.f32 %f{{[0-9]+}}, [%rl{{[0-9]+}}] +; PTX64: ld.f32 %f{{[0-9]+}}, [%rd{{[0-9]+}}] ; PTX64: ret %a = load float addrspace(0)* %ptr ret float %a @@ -54,9 +54,9 @@ define float @ld_global_f32(float addrspace(0)* %ptr) { ;; f64 define double @ld_global_f64(double addrspace(0)* %ptr) { -; PTX32: ld.f64 %fl{{[0-9]+}}, [%r{{[0-9]+}}] +; PTX32: ld.f64 %fd{{[0-9]+}}, [%r{{[0-9]+}}] ; PTX32: ret -; PTX64: ld.f64 %fl{{[0-9]+}}, [%rl{{[0-9]+}}] +; PTX64: ld.f64 %fd{{[0-9]+}}, [%rd{{[0-9]+}}] ; PTX64: ret %a = load double addrspace(0)* %ptr ret double %a diff --git a/test/CodeGen/NVPTX/ldparam-v4.ll b/test/CodeGen/NVPTX/ldparam-v4.ll new file mode 100644 index 000000000000..ec306aafe854 --- /dev/null +++ b/test/CodeGen/NVPTX/ldparam-v4.ll @@ -0,0 +1,10 @@ +; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s + +declare <4 x float> @bar() + +define void @foo(<4 x float>* %ptr) { +; CHECK: ld.param.v4.f32 + %val = tail call <4 x float> @bar() + store <4 x float> %val, <4 x float>* %ptr + ret void +} diff --git a/test/CodeGen/NVPTX/ldu-i8.ll b/test/CodeGen/NVPTX/ldu-i8.ll index 81a82b2c38b5..9cc667557906 100644 --- a/test/CodeGen/NVPTX/ldu-i8.ll +++ b/test/CodeGen/NVPTX/ldu-i8.ll @@ -2,13 +2,15 @@ target datalayout = "e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64" -declare i8 @llvm.nvvm.ldu.global.i.i8(i8*) +declare i8 @llvm.nvvm.ldu.global.i.i8.p0i8(i8*) define i8 @foo(i8* %a) { ; Ensure we properly truncate off the high-order 24 bits ; CHECK: ldu.global.u8 ; CHECK: cvt.u32.u16 ; CHECK: and.b32 %r{{[0-9]+}}, %r{{[0-9]+}}, 255 - %val = tail call i8 @llvm.nvvm.ldu.global.i.i8(i8* %a) + %val = tail call i8 @llvm.nvvm.ldu.global.i.i8.p0i8(i8* %a), !align !0 ret i8 %val } + +!0 = metadata !{i32 4} diff --git a/test/CodeGen/NVPTX/ldu-ldg.ll b/test/CodeGen/NVPTX/ldu-ldg.ll new file mode 100644 index 000000000000..3b0619ff5175 --- /dev/null +++ b/test/CodeGen/NVPTX/ldu-ldg.ll @@ -0,0 +1,40 @@ +; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s + + +declare i8 @llvm.nvvm.ldu.global.i.i8.p1i8(i8 addrspace(1)* %ptr) +declare i32 @llvm.nvvm.ldu.global.i.i32.p1i32(i32 addrspace(1)* %ptr) +declare i8 @llvm.nvvm.ldg.global.i.i8.p1i8(i8 addrspace(1)* %ptr) +declare i32 @llvm.nvvm.ldg.global.i.i32.p1i32(i32 addrspace(1)* %ptr) + + +; CHECK: func0 +define i8 @func0(i8 addrspace(1)* %ptr) { +; ldu.global.u8 + %val = tail call i8 @llvm.nvvm.ldu.global.i.i8.p1i8(i8 addrspace(1)* %ptr), !align !0 + ret i8 %val +} + +; CHECK: func1 +define i32 @func1(i32 addrspace(1)* %ptr) { +; ldu.global.u32 + %val = tail call i32 @llvm.nvvm.ldu.global.i.i32.p1i32(i32 addrspace(1)* %ptr), !align !0 + ret i32 %val +} + +; CHECK: func2 +define i8 @func2(i8 addrspace(1)* %ptr) { +; ld.global.nc.u8 + %val = tail call i8 @llvm.nvvm.ldg.global.i.i8.p1i8(i8 addrspace(1)* %ptr), !align !0 + ret i8 %val +} + +; CHECK: func3 +define i32 @func3(i32 addrspace(1)* %ptr) { +; ld.global.nc.u32 + %val = tail call i32 @llvm.nvvm.ldg.global.i.i32.p1i32(i32 addrspace(1)* %ptr), !align !0 + ret i32 %val +} + + + +!0 = metadata !{i32 4} diff --git a/test/CodeGen/NVPTX/ldu-reg-plus-offset.ll b/test/CodeGen/NVPTX/ldu-reg-plus-offset.ll index 26cadc401b79..55707ea85106 100644 --- a/test/CodeGen/NVPTX/ldu-reg-plus-offset.ll +++ b/test/CodeGen/NVPTX/ldu-reg-plus-offset.ll @@ -7,9 +7,9 @@ define void @reg_plus_offset(i32* %a) { ; CHECK: ldu.global.u32 %r{{[0-9]+}}, [%r{{[0-9]+}}+32]; ; CHECK: ldu.global.u32 %r{{[0-9]+}}, [%r{{[0-9]+}}+36]; %p2 = getelementptr i32* %a, i32 8 - %t1 = call i32 @llvm.nvvm.ldu.global.i.i32(i32* %p2), !align !1 + %t1 = call i32 @llvm.nvvm.ldu.global.i.i32.p0i32(i32* %p2), !align !1 %p3 = getelementptr i32* %a, i32 9 - %t2 = call i32 @llvm.nvvm.ldu.global.i.i32(i32* %p3), !align !1 + %t2 = call i32 @llvm.nvvm.ldu.global.i.i32.p0i32(i32* %p3), !align !1 %t3 = mul i32 %t1, %t2 store i32 %t3, i32* %a ret void @@ -17,5 +17,5 @@ define void @reg_plus_offset(i32* %a) { !1 = metadata !{ i32 4 } -declare i32 @llvm.nvvm.ldu.global.i.i32(i32*) +declare i32 @llvm.nvvm.ldu.global.i.i32.p0i32(i32*) declare i32 @llvm.nvvm.read.ptx.sreg.tid.x() diff --git a/test/CodeGen/NVPTX/lit.local.cfg b/test/CodeGen/NVPTX/lit.local.cfg index 85cf8c2c8c07..2cb98eb371b2 100644 --- a/test/CodeGen/NVPTX/lit.local.cfg +++ b/test/CodeGen/NVPTX/lit.local.cfg @@ -1,3 +1,2 @@ -targets = set(config.root.targets_to_build.split()) -if not 'NVPTX' in targets: +if not 'NVPTX' in config.root.targets: config.unsupported = True diff --git a/test/CodeGen/NVPTX/local-stack-frame.ll b/test/CodeGen/NVPTX/local-stack-frame.ll index 178dff1a5d3f..377eee9170e6 100644 --- a/test/CodeGen/NVPTX/local-stack-frame.ll +++ b/test/CodeGen/NVPTX/local-stack-frame.ll @@ -3,16 +3,16 @@ ; Ensure we access the local stack properly -; PTX32: mov.u32 %r{{[0-9]+}}, __local_depot{{[0-9]+}}; -; PTX32: cvta.local.u32 %SP, %r{{[0-9]+}}; -; PTX32: ld.param.u32 %r{{[0-9]+}}, [foo_param_0]; -; PTX32: st.u32 [%SP+0], %r{{[0-9]+}}; -; PTX64: mov.u64 %rl{{[0-9]+}}, __local_depot{{[0-9]+}}; -; PTX64: cvta.local.u64 %SP, %rl{{[0-9]+}}; -; PTX64: ld.param.u32 %r{{[0-9]+}}, [foo_param_0]; -; PTX64: st.u32 [%SP+0], %r{{[0-9]+}}; +; PTX32: mov.u32 %r{{[0-9]+}}, __local_depot{{[0-9]+}}; +; PTX32: cvta.local.u32 %SP, %r{{[0-9]+}}; +; PTX32: ld.param.u32 %r{{[0-9]+}}, [foo_param_0]; +; PTX32: st.volatile.u32 [%SP+0], %r{{[0-9]+}}; +; PTX64: mov.u64 %rd{{[0-9]+}}, __local_depot{{[0-9]+}}; +; PTX64: cvta.local.u64 %SP, %rd{{[0-9]+}}; +; PTX64: ld.param.u32 %r{{[0-9]+}}, [foo_param_0]; +; PTX64: st.volatile.u32 [%SP+0], %r{{[0-9]+}}; define void @foo(i32 %a) { %local = alloca i32, align 4 - store i32 %a, i32* %local + store volatile i32 %a, i32* %local ret void } diff --git a/test/CodeGen/NVPTX/managed.ll b/test/CodeGen/NVPTX/managed.ll new file mode 100644 index 000000000000..4d7e7817f77b --- /dev/null +++ b/test/CodeGen/NVPTX/managed.ll @@ -0,0 +1,11 @@ +; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s + + +; CHECK: .visible .global .align 4 .u32 device_g; +@device_g = addrspace(1) global i32 zeroinitializer +; CHECK: .visible .global .attribute(.managed) .align 4 .u32 managed_g; +@managed_g = addrspace(1) global i32 zeroinitializer + + +!nvvm.annotations = !{!0} +!0 = metadata !{i32 addrspace(1)* @managed_g, metadata !"managed", i32 1} diff --git a/test/CodeGen/NVPTX/misaligned-vector-ldst.ll b/test/CodeGen/NVPTX/misaligned-vector-ldst.ll new file mode 100644 index 000000000000..90c9c4306de7 --- /dev/null +++ b/test/CodeGen/NVPTX/misaligned-vector-ldst.ll @@ -0,0 +1,77 @@ +; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s + +target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64" +target triple = "nvptx64-nvidia-cuda" + +; CHECK-LABEL: t1 +define <4 x float> @t1(i8* %p1) { +; CHECK-NOT: ld.v4 +; CHECK-NOT: ld.v2 +; CHECK-NOT: ld.f32 +; CHECK: ld.u8 + %cast = bitcast i8* %p1 to <4 x float>* + %r = load <4 x float>* %cast, align 1 + ret <4 x float> %r +} + +; CHECK-LABEL: t2 +define <4 x float> @t2(i8* %p1) { +; CHECK-NOT: ld.v4 +; CHECK-NOT: ld.v2 +; CHECK: ld.f32 + %cast = bitcast i8* %p1 to <4 x float>* + %r = load <4 x float>* %cast, align 4 + ret <4 x float> %r +} + +; CHECK-LABEL: t3 +define <4 x float> @t3(i8* %p1) { +; CHECK-NOT: ld.v4 +; CHECK: ld.v2 + %cast = bitcast i8* %p1 to <4 x float>* + %r = load <4 x float>* %cast, align 8 + ret <4 x float> %r +} + +; CHECK-LABEL: t4 +define <4 x float> @t4(i8* %p1) { +; CHECK: ld.v4 + %cast = bitcast i8* %p1 to <4 x float>* + %r = load <4 x float>* %cast, align 16 + ret <4 x float> %r +} + + +; CHECK-LABEL: s1 +define void @s1(<4 x float>* %p1, <4 x float> %v) { +; CHECK-NOT: st.v4 +; CHECK-NOT: st.v2 +; CHECK-NOT: st.f32 +; CHECK: st.u8 + store <4 x float> %v, <4 x float>* %p1, align 1 + ret void +} + +; CHECK-LABEL: s2 +define void @s2(<4 x float>* %p1, <4 x float> %v) { +; CHECK-NOT: st.v4 +; CHECK-NOT: st.v2 +; CHECK: st.f32 + store <4 x float> %v, <4 x float>* %p1, align 4 + ret void +} + +; CHECK-LABEL: s3 +define void @s3(<4 x float>* %p1, <4 x float> %v) { +; CHECK-NOT: st.v4 + store <4 x float> %v, <4 x float>* %p1, align 8 + ret void +} + +; CHECK-LABEL: s4 +define void @s4(<4 x float>* %p1, <4 x float> %v) { +; CHECK: st.v4 + store <4 x float> %v, <4 x float>* %p1, align 16 + ret void +} + diff --git a/test/CodeGen/NVPTX/mulwide.ll b/test/CodeGen/NVPTX/mulwide.ll new file mode 100644 index 000000000000..43bb63098f67 --- /dev/null +++ b/test/CodeGen/NVPTX/mulwide.ll @@ -0,0 +1,46 @@ +; RUN: llc < %s -march=nvptx -mcpu=sm_20 -O3 | FileCheck %s --check-prefix=OPT +; RUN: llc < %s -march=nvptx -mcpu=sm_20 -O0 | FileCheck %s --check-prefix=NOOPT + +; OPT-LABEL: @mulwide16 +; NOOPT-LABEL: @mulwide16 +define i32 @mulwide16(i16 %a, i16 %b) { +; OPT: mul.wide.s16 +; NOOPT: mul.lo.s32 + %val0 = sext i16 %a to i32 + %val1 = sext i16 %b to i32 + %val2 = mul i32 %val0, %val1 + ret i32 %val2 +} + +; OPT-LABEL: @mulwideu16 +; NOOPT-LABEL: @mulwideu16 +define i32 @mulwideu16(i16 %a, i16 %b) { +; OPT: mul.wide.u16 +; NOOPT: mul.lo.s32 + %val0 = zext i16 %a to i32 + %val1 = zext i16 %b to i32 + %val2 = mul i32 %val0, %val1 + ret i32 %val2 +} + +; OPT-LABEL: @mulwide32 +; NOOPT-LABEL: @mulwide32 +define i64 @mulwide32(i32 %a, i32 %b) { +; OPT: mul.wide.s32 +; NOOPT: mul.lo.s64 + %val0 = sext i32 %a to i64 + %val1 = sext i32 %b to i64 + %val2 = mul i64 %val0, %val1 + ret i64 %val2 +} + +; OPT-LABEL: @mulwideu32 +; NOOPT-LABEL: @mulwideu32 +define i64 @mulwideu32(i32 %a, i32 %b) { +; OPT: mul.wide.u32 +; NOOPT: mul.lo.s64 + %val0 = zext i32 %a to i64 + %val1 = zext i32 %b to i64 + %val2 = mul i64 %val0, %val1 + ret i64 %val2 +} diff --git a/test/CodeGen/NVPTX/noduplicate-syncthreads.ll b/test/CodeGen/NVPTX/noduplicate-syncthreads.ll new file mode 100644 index 000000000000..64745fcba3ba --- /dev/null +++ b/test/CodeGen/NVPTX/noduplicate-syncthreads.ll @@ -0,0 +1,74 @@ +; RUN: opt < %s -O3 -S | FileCheck %s + +; Make sure the call to syncthreads is not duplicate here by the LLVM +; optimizations, because it has the noduplicate attribute set. + +; CHECK: call void @llvm.cuda.syncthreads +; CHECK-NOT: call void @llvm.cuda.syncthreads + +; Function Attrs: nounwind +define void @foo(float* %output) #1 { +entry: + %output.addr = alloca float*, align 8 + store float* %output, float** %output.addr, align 8 + %0 = load float** %output.addr, align 8 + %arrayidx = getelementptr inbounds float* %0, i64 0 + %1 = load float* %arrayidx, align 4 + %conv = fpext float %1 to double + %cmp = fcmp olt double %conv, 1.000000e+01 + br i1 %cmp, label %if.then, label %if.else + +if.then: ; preds = %entry + %2 = load float** %output.addr, align 8 + %3 = load float* %2, align 4 + %conv1 = fpext float %3 to double + %add = fadd double %conv1, 1.000000e+00 + %conv2 = fptrunc double %add to float + store float %conv2, float* %2, align 4 + br label %if.end + +if.else: ; preds = %entry + %4 = load float** %output.addr, align 8 + %5 = load float* %4, align 4 + %conv3 = fpext float %5 to double + %add4 = fadd double %conv3, 2.000000e+00 + %conv5 = fptrunc double %add4 to float + store float %conv5, float* %4, align 4 + br label %if.end + +if.end: ; preds = %if.else, %if.then + call void @llvm.cuda.syncthreads() + %6 = load float** %output.addr, align 8 + %arrayidx6 = getelementptr inbounds float* %6, i64 0 + %7 = load float* %arrayidx6, align 4 + %conv7 = fpext float %7 to double + %cmp8 = fcmp olt double %conv7, 1.000000e+01 + br i1 %cmp8, label %if.then9, label %if.else13 + +if.then9: ; preds = %if.end + %8 = load float** %output.addr, align 8 + %9 = load float* %8, align 4 + %conv10 = fpext float %9 to double + %add11 = fadd double %conv10, 3.000000e+00 + %conv12 = fptrunc double %add11 to float + store float %conv12, float* %8, align 4 + br label %if.end17 + +if.else13: ; preds = %if.end + %10 = load float** %output.addr, align 8 + %11 = load float* %10, align 4 + %conv14 = fpext float %11 to double + %add15 = fadd double %conv14, 4.000000e+00 + %conv16 = fptrunc double %add15 to float + store float %conv16, float* %10, align 4 + br label %if.end17 + +if.end17: ; preds = %if.else13, %if.then9 + ret void +} + +; Function Attrs: noduplicate nounwind +declare void @llvm.cuda.syncthreads() #2 + +!0 = metadata !{void (float*)* @foo, metadata !"kernel", i32 1} +!1 = metadata !{null, metadata !"align", i32 8} diff --git a/test/CodeGen/NVPTX/nvvm-reflect.ll b/test/CodeGen/NVPTX/nvvm-reflect.ll index 0d02194651e3..21e9c69e657a 100644 --- a/test/CodeGen/NVPTX/nvvm-reflect.ll +++ b/test/CodeGen/NVPTX/nvvm-reflect.ll @@ -1,7 +1,7 @@ ; RUN: opt < %s -S -nvvm-reflect -nvvm-reflect-list USE_MUL=0 -O2 | FileCheck %s --check-prefix=USE_MUL_0 ; RUN: opt < %s -S -nvvm-reflect -nvvm-reflect-list USE_MUL=1 -O2 | FileCheck %s --check-prefix=USE_MUL_1 -@str = private addrspace(4) unnamed_addr constant [8 x i8] c"USE_MUL\00" +@str = private unnamed_addr addrspace(4) constant [8 x i8] c"USE_MUL\00" declare i32 @__nvvm_reflect(i8*) declare i8* @llvm.nvvm.ptr.constant.to.gen.p0i8.p4i8(i8 addrspace(4)*) @@ -32,3 +32,17 @@ exit: %ret = phi float [%ret1, %use_mul], [%ret2, %use_add] ret float %ret } + +declare i32 @llvm.nvvm.reflect.p0i8(i8*) + +; USE_MUL_0: define i32 @intrinsic +; USE_MUL_1: define i32 @intrinsic +define i32 @intrinsic() { +; USE_MUL_0-NOT: call i32 @llvm.nvvm.reflect +; USE_MUL_0: ret i32 0 +; USE_MUL_1-NOT: call i32 @llvm.nvvm.reflect +; USE_MUL_1: ret i32 1 + %ptr = tail call i8* @llvm.nvvm.ptr.constant.to.gen.p0i8.p4i8(i8 addrspace(4)* getelementptr inbounds ([8 x i8] addrspace(4)* @str, i32 0, i32 0)) + %reflect = tail call i32 @llvm.nvvm.reflect.p0i8(i8* %ptr) + ret i32 %reflect +} diff --git a/test/CodeGen/NVPTX/pr13291-i1-store.ll b/test/CodeGen/NVPTX/pr13291-i1-store.ll index e7a81be01b14..cc67a6fff8e4 100644 --- a/test/CodeGen/NVPTX/pr13291-i1-store.ll +++ b/test/CodeGen/NVPTX/pr13291-i1-store.ll @@ -5,7 +5,7 @@ define ptx_kernel void @t1(i1* %a) { ; PTX32: mov.u16 %rs{{[0-9]+}}, 0; ; PTX32-NEXT: st.u8 [%r{{[0-9]+}}], %rs{{[0-9]+}}; ; PTX64: mov.u16 %rs{{[0-9]+}}, 0; -; PTX64-NEXT: st.u8 [%rl{{[0-9]+}}], %rs{{[0-9]+}}; +; PTX64-NEXT: st.u8 [%rd{{[0-9]+}}], %rs{{[0-9]+}}; store i1 false, i1* %a ret void } @@ -15,7 +15,7 @@ define ptx_kernel void @t2(i1* %a, i8* %b) { ; PTX32: ld.u8 %rs{{[0-9]+}}, [%r{{[0-9]+}}] ; PTX32: and.b16 %rs{{[0-9]+}}, %rs{{[0-9]+}}, 1; ; PTX32: setp.eq.b16 %p{{[0-9]+}}, %rs{{[0-9]+}}, 1; -; PTX64: ld.u8 %rs{{[0-9]+}}, [%rl{{[0-9]+}}] +; PTX64: ld.u8 %rs{{[0-9]+}}, [%rd{{[0-9]+}}] ; PTX64: and.b16 %rs{{[0-9]+}}, %rs{{[0-9]+}}, 1; ; PTX64: setp.eq.b16 %p{{[0-9]+}}, %rs{{[0-9]+}}, 1; diff --git a/test/CodeGen/NVPTX/rotate.ll b/test/CodeGen/NVPTX/rotate.ll new file mode 100644 index 000000000000..dfc8b4fd5fcb --- /dev/null +++ b/test/CodeGen/NVPTX/rotate.ll @@ -0,0 +1,58 @@ +; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck --check-prefix=SM20 %s +; RUN: llc < %s -march=nvptx -mcpu=sm_35 | FileCheck --check-prefix=SM35 %s + + +declare i32 @llvm.nvvm.rotate.b32(i32, i32) +declare i64 @llvm.nvvm.rotate.b64(i64, i32) +declare i64 @llvm.nvvm.rotate.right.b64(i64, i32) + +; SM20: rotate32 +; SM35: rotate32 +define i32 @rotate32(i32 %a, i32 %b) { +; SM20: shl.b32 +; SM20: sub.s32 +; SM20: shr.b32 +; SM20: add.u32 +; SM35: shf.l.wrap.b32 + %val = tail call i32 @llvm.nvvm.rotate.b32(i32 %a, i32 %b) + ret i32 %val +} + +; SM20: rotate64 +; SM35: rotate64 +define i64 @rotate64(i64 %a, i32 %b) { +; SM20: shl.b64 +; SM20: sub.u32 +; SM20: shr.b64 +; SM20: add.u64 +; SM35: shf.l.wrap.b32 +; SM35: shf.l.wrap.b32 + %val = tail call i64 @llvm.nvvm.rotate.b64(i64 %a, i32 %b) + ret i64 %val +} + +; SM20: rotateright64 +; SM35: rotateright64 +define i64 @rotateright64(i64 %a, i32 %b) { +; SM20: shr.b64 +; SM20: sub.u32 +; SM20: shl.b64 +; SM20: add.u64 +; SM35: shf.r.wrap.b32 +; SM35: shf.r.wrap.b32 + %val = tail call i64 @llvm.nvvm.rotate.right.b64(i64 %a, i32 %b) + ret i64 %val +} + +; SM20: rotl0 +; SM35: rotl0 +define i32 @rotl0(i32 %x) { +; SM20: shl.b32 +; SM20: shr.b32 +; SM20: add.u32 +; SM35: shf.l.wrap.b32 + %t0 = shl i32 %x, 8 + %t1 = lshr i32 %x, 24 + %t2 = or i32 %t0, %t1 + ret i32 %t2 +} diff --git a/test/CodeGen/NVPTX/shift-parts.ll b/test/CodeGen/NVPTX/shift-parts.ll new file mode 100644 index 000000000000..748297caf339 --- /dev/null +++ b/test/CodeGen/NVPTX/shift-parts.ll @@ -0,0 +1,38 @@ +; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s + +; CHECK: shift_parts_left_128 +define void @shift_parts_left_128(i128* %val, i128* %amtptr) { +; CHECK: shl.b64 +; CHECK: mov.u32 +; CHECK: sub.s32 +; CHECK: shr.u64 +; CHECK: or.b64 +; CHECK: add.s32 +; CHECK: shl.b64 +; CHECK: setp.gt.s32 +; CHECK: selp.b64 +; CHECK: shl.b64 + %amt = load i128* %amtptr + %a = load i128* %val + %val0 = shl i128 %a, %amt + store i128 %val0, i128* %val + ret void +} + +; CHECK: shift_parts_right_128 +define void @shift_parts_right_128(i128* %val, i128* %amtptr) { +; CHECK: shr.u64 +; CHECK: sub.s32 +; CHECK: shl.b64 +; CHECK: or.b64 +; CHECK: add.s32 +; CHECK: shr.s64 +; CHECK: setp.gt.s32 +; CHECK: selp.b64 +; CHECK: shr.s64 + %amt = load i128* %amtptr + %a = load i128* %val + %val0 = ashr i128 %a, %amt + store i128 %val0, i128* %val + ret void +} diff --git a/test/CodeGen/NVPTX/st-addrspace.ll b/test/CodeGen/NVPTX/st-addrspace.ll index 68c09fe065bc..34a83f343324 100644 --- a/test/CodeGen/NVPTX/st-addrspace.ll +++ b/test/CodeGen/NVPTX/st-addrspace.ll @@ -7,7 +7,7 @@ define void @st_global_i8(i8 addrspace(1)* %ptr, i8 %a) { ; PTX32: st.global.u8 [%r{{[0-9]+}}], %rs{{[0-9]+}} ; PTX32: ret -; PTX64: st.global.u8 [%rl{{[0-9]+}}], %rs{{[0-9]+}} +; PTX64: st.global.u8 [%rd{{[0-9]+}}], %rs{{[0-9]+}} ; PTX64: ret store i8 %a, i8 addrspace(1)* %ptr ret void @@ -16,7 +16,7 @@ define void @st_global_i8(i8 addrspace(1)* %ptr, i8 %a) { define void @st_shared_i8(i8 addrspace(3)* %ptr, i8 %a) { ; PTX32: st.shared.u8 [%r{{[0-9]+}}], %rs{{[0-9]+}} ; PTX32: ret -; PTX64: st.shared.u8 [%rl{{[0-9]+}}], %rs{{[0-9]+}} +; PTX64: st.shared.u8 [%rd{{[0-9]+}}], %rs{{[0-9]+}} ; PTX64: ret store i8 %a, i8 addrspace(3)* %ptr ret void @@ -25,7 +25,7 @@ define void @st_shared_i8(i8 addrspace(3)* %ptr, i8 %a) { define void @st_local_i8(i8 addrspace(5)* %ptr, i8 %a) { ; PTX32: st.local.u8 [%r{{[0-9]+}}], %rs{{[0-9]+}} ; PTX32: ret -; PTX64: st.local.u8 [%rl{{[0-9]+}}], %rs{{[0-9]+}} +; PTX64: st.local.u8 [%rd{{[0-9]+}}], %rs{{[0-9]+}} ; PTX64: ret store i8 %a, i8 addrspace(5)* %ptr ret void @@ -36,7 +36,7 @@ define void @st_local_i8(i8 addrspace(5)* %ptr, i8 %a) { define void @st_global_i16(i16 addrspace(1)* %ptr, i16 %a) { ; PTX32: st.global.u16 [%r{{[0-9]+}}], %rs{{[0-9]+}} ; PTX32: ret -; PTX64: st.global.u16 [%rl{{[0-9]+}}], %rs{{[0-9]+}} +; PTX64: st.global.u16 [%rd{{[0-9]+}}], %rs{{[0-9]+}} ; PTX64: ret store i16 %a, i16 addrspace(1)* %ptr ret void @@ -45,7 +45,7 @@ define void @st_global_i16(i16 addrspace(1)* %ptr, i16 %a) { define void @st_shared_i16(i16 addrspace(3)* %ptr, i16 %a) { ; PTX32: st.shared.u16 [%r{{[0-9]+}}], %rs{{[0-9]+}} ; PTX32: ret -; PTX64: st.shared.u16 [%rl{{[0-9]+}}], %rs{{[0-9]+}} +; PTX64: st.shared.u16 [%rd{{[0-9]+}}], %rs{{[0-9]+}} ; PTX64: ret store i16 %a, i16 addrspace(3)* %ptr ret void @@ -54,7 +54,7 @@ define void @st_shared_i16(i16 addrspace(3)* %ptr, i16 %a) { define void @st_local_i16(i16 addrspace(5)* %ptr, i16 %a) { ; PTX32: st.local.u16 [%r{{[0-9]+}}], %rs{{[0-9]+}} ; PTX32: ret -; PTX64: st.local.u16 [%rl{{[0-9]+}}], %rs{{[0-9]+}} +; PTX64: st.local.u16 [%rd{{[0-9]+}}], %rs{{[0-9]+}} ; PTX64: ret store i16 %a, i16 addrspace(5)* %ptr ret void @@ -65,7 +65,7 @@ define void @st_local_i16(i16 addrspace(5)* %ptr, i16 %a) { define void @st_global_i32(i32 addrspace(1)* %ptr, i32 %a) { ; PTX32: st.global.u32 [%r{{[0-9]+}}], %r{{[0-9]+}} ; PTX32: ret -; PTX64: st.global.u32 [%rl{{[0-9]+}}], %r{{[0-9]+}} +; PTX64: st.global.u32 [%rd{{[0-9]+}}], %r{{[0-9]+}} ; PTX64: ret store i32 %a, i32 addrspace(1)* %ptr ret void @@ -74,7 +74,7 @@ define void @st_global_i32(i32 addrspace(1)* %ptr, i32 %a) { define void @st_shared_i32(i32 addrspace(3)* %ptr, i32 %a) { ; PTX32: st.shared.u32 [%r{{[0-9]+}}], %r{{[0-9]+}} ; PTX32: ret -; PTX64: st.shared.u32 [%rl{{[0-9]+}}], %r{{[0-9]+}} +; PTX64: st.shared.u32 [%rd{{[0-9]+}}], %r{{[0-9]+}} ; PTX64: ret store i32 %a, i32 addrspace(3)* %ptr ret void @@ -83,7 +83,7 @@ define void @st_shared_i32(i32 addrspace(3)* %ptr, i32 %a) { define void @st_local_i32(i32 addrspace(5)* %ptr, i32 %a) { ; PTX32: st.local.u32 [%r{{[0-9]+}}], %r{{[0-9]+}} ; PTX32: ret -; PTX64: st.local.u32 [%rl{{[0-9]+}}], %r{{[0-9]+}} +; PTX64: st.local.u32 [%rd{{[0-9]+}}], %r{{[0-9]+}} ; PTX64: ret store i32 %a, i32 addrspace(5)* %ptr ret void @@ -92,27 +92,27 @@ define void @st_local_i32(i32 addrspace(5)* %ptr, i32 %a) { ;; i64 define void @st_global_i64(i64 addrspace(1)* %ptr, i64 %a) { -; PTX32: st.global.u64 [%r{{[0-9]+}}], %rl{{[0-9]+}} +; PTX32: st.global.u64 [%r{{[0-9]+}}], %rd{{[0-9]+}} ; PTX32: ret -; PTX64: st.global.u64 [%rl{{[0-9]+}}], %rl{{[0-9]+}} +; PTX64: st.global.u64 [%rd{{[0-9]+}}], %rd{{[0-9]+}} ; PTX64: ret store i64 %a, i64 addrspace(1)* %ptr ret void } define void @st_shared_i64(i64 addrspace(3)* %ptr, i64 %a) { -; PTX32: st.shared.u64 [%r{{[0-9]+}}], %rl{{[0-9]+}} +; PTX32: st.shared.u64 [%r{{[0-9]+}}], %rd{{[0-9]+}} ; PTX32: ret -; PTX64: st.shared.u64 [%rl{{[0-9]+}}], %rl{{[0-9]+}} +; PTX64: st.shared.u64 [%rd{{[0-9]+}}], %rd{{[0-9]+}} ; PTX64: ret store i64 %a, i64 addrspace(3)* %ptr ret void } define void @st_local_i64(i64 addrspace(5)* %ptr, i64 %a) { -; PTX32: st.local.u64 [%r{{[0-9]+}}], %rl{{[0-9]+}} +; PTX32: st.local.u64 [%r{{[0-9]+}}], %rd{{[0-9]+}} ; PTX32: ret -; PTX64: st.local.u64 [%rl{{[0-9]+}}], %rl{{[0-9]+}} +; PTX64: st.local.u64 [%rd{{[0-9]+}}], %rd{{[0-9]+}} ; PTX64: ret store i64 %a, i64 addrspace(5)* %ptr ret void @@ -123,7 +123,7 @@ define void @st_local_i64(i64 addrspace(5)* %ptr, i64 %a) { define void @st_global_f32(float addrspace(1)* %ptr, float %a) { ; PTX32: st.global.f32 [%r{{[0-9]+}}], %f{{[0-9]+}} ; PTX32: ret -; PTX64: st.global.f32 [%rl{{[0-9]+}}], %f{{[0-9]+}} +; PTX64: st.global.f32 [%rd{{[0-9]+}}], %f{{[0-9]+}} ; PTX64: ret store float %a, float addrspace(1)* %ptr ret void @@ -132,7 +132,7 @@ define void @st_global_f32(float addrspace(1)* %ptr, float %a) { define void @st_shared_f32(float addrspace(3)* %ptr, float %a) { ; PTX32: st.shared.f32 [%r{{[0-9]+}}], %f{{[0-9]+}} ; PTX32: ret -; PTX64: st.shared.f32 [%rl{{[0-9]+}}], %f{{[0-9]+}} +; PTX64: st.shared.f32 [%rd{{[0-9]+}}], %f{{[0-9]+}} ; PTX64: ret store float %a, float addrspace(3)* %ptr ret void @@ -141,7 +141,7 @@ define void @st_shared_f32(float addrspace(3)* %ptr, float %a) { define void @st_local_f32(float addrspace(5)* %ptr, float %a) { ; PTX32: st.local.f32 [%r{{[0-9]+}}], %f{{[0-9]+}} ; PTX32: ret -; PTX64: st.local.f32 [%rl{{[0-9]+}}], %f{{[0-9]+}} +; PTX64: st.local.f32 [%rd{{[0-9]+}}], %f{{[0-9]+}} ; PTX64: ret store float %a, float addrspace(5)* %ptr ret void @@ -150,27 +150,27 @@ define void @st_local_f32(float addrspace(5)* %ptr, float %a) { ;; f64 define void @st_global_f64(double addrspace(1)* %ptr, double %a) { -; PTX32: st.global.f64 [%r{{[0-9]+}}], %fl{{[0-9]+}} +; PTX32: st.global.f64 [%r{{[0-9]+}}], %fd{{[0-9]+}} ; PTX32: ret -; PTX64: st.global.f64 [%rl{{[0-9]+}}], %fl{{[0-9]+}} +; PTX64: st.global.f64 [%rd{{[0-9]+}}], %fd{{[0-9]+}} ; PTX64: ret store double %a, double addrspace(1)* %ptr ret void } define void @st_shared_f64(double addrspace(3)* %ptr, double %a) { -; PTX32: st.shared.f64 [%r{{[0-9]+}}], %fl{{[0-9]+}} +; PTX32: st.shared.f64 [%r{{[0-9]+}}], %fd{{[0-9]+}} ; PTX32: ret -; PTX64: st.shared.f64 [%rl{{[0-9]+}}], %fl{{[0-9]+}} +; PTX64: st.shared.f64 [%rd{{[0-9]+}}], %fd{{[0-9]+}} ; PTX64: ret store double %a, double addrspace(3)* %ptr ret void } define void @st_local_f64(double addrspace(5)* %ptr, double %a) { -; PTX32: st.local.f64 [%r{{[0-9]+}}], %fl{{[0-9]+}} +; PTX32: st.local.f64 [%r{{[0-9]+}}], %fd{{[0-9]+}} ; PTX32: ret -; PTX64: st.local.f64 [%rl{{[0-9]+}}], %fl{{[0-9]+}} +; PTX64: st.local.f64 [%rd{{[0-9]+}}], %fd{{[0-9]+}} ; PTX64: ret store double %a, double addrspace(5)* %ptr ret void diff --git a/test/CodeGen/NVPTX/st-generic.ll b/test/CodeGen/NVPTX/st-generic.ll index b9c616fbd19e..022f7ab214ca 100644 --- a/test/CodeGen/NVPTX/st-generic.ll +++ b/test/CodeGen/NVPTX/st-generic.ll @@ -7,7 +7,7 @@ define void @st_global_i8(i8 addrspace(0)* %ptr, i8 %a) { ; PTX32: st.u8 [%r{{[0-9]+}}], %rs{{[0-9]+}} ; PTX32: ret -; PTX64: st.u8 [%rl{{[0-9]+}}], %rs{{[0-9]+}} +; PTX64: st.u8 [%rd{{[0-9]+}}], %rs{{[0-9]+}} ; PTX64: ret store i8 %a, i8 addrspace(0)* %ptr ret void @@ -18,7 +18,7 @@ define void @st_global_i8(i8 addrspace(0)* %ptr, i8 %a) { define void @st_global_i16(i16 addrspace(0)* %ptr, i16 %a) { ; PTX32: st.u16 [%r{{[0-9]+}}], %rs{{[0-9]+}} ; PTX32: ret -; PTX64: st.u16 [%rl{{[0-9]+}}], %rs{{[0-9]+}} +; PTX64: st.u16 [%rd{{[0-9]+}}], %rs{{[0-9]+}} ; PTX64: ret store i16 %a, i16 addrspace(0)* %ptr ret void @@ -29,7 +29,7 @@ define void @st_global_i16(i16 addrspace(0)* %ptr, i16 %a) { define void @st_global_i32(i32 addrspace(0)* %ptr, i32 %a) { ; PTX32: st.u32 [%r{{[0-9]+}}], %r{{[0-9]+}} ; PTX32: ret -; PTX64: st.u32 [%rl{{[0-9]+}}], %r{{[0-9]+}} +; PTX64: st.u32 [%rd{{[0-9]+}}], %r{{[0-9]+}} ; PTX64: ret store i32 %a, i32 addrspace(0)* %ptr ret void @@ -38,9 +38,9 @@ define void @st_global_i32(i32 addrspace(0)* %ptr, i32 %a) { ;; i64 define void @st_global_i64(i64 addrspace(0)* %ptr, i64 %a) { -; PTX32: st.u64 [%r{{[0-9]+}}], %rl{{[0-9]+}} +; PTX32: st.u64 [%r{{[0-9]+}}], %rd{{[0-9]+}} ; PTX32: ret -; PTX64: st.u64 [%rl{{[0-9]+}}], %rl{{[0-9]+}} +; PTX64: st.u64 [%rd{{[0-9]+}}], %rd{{[0-9]+}} ; PTX64: ret store i64 %a, i64 addrspace(0)* %ptr ret void @@ -51,7 +51,7 @@ define void @st_global_i64(i64 addrspace(0)* %ptr, i64 %a) { define void @st_global_f32(float addrspace(0)* %ptr, float %a) { ; PTX32: st.f32 [%r{{[0-9]+}}], %f{{[0-9]+}} ; PTX32: ret -; PTX64: st.f32 [%rl{{[0-9]+}}], %f{{[0-9]+}} +; PTX64: st.f32 [%rd{{[0-9]+}}], %f{{[0-9]+}} ; PTX64: ret store float %a, float addrspace(0)* %ptr ret void @@ -60,9 +60,9 @@ define void @st_global_f32(float addrspace(0)* %ptr, float %a) { ;; f64 define void @st_global_f64(double addrspace(0)* %ptr, double %a) { -; PTX32: st.f64 [%r{{[0-9]+}}], %fl{{[0-9]+}} +; PTX32: st.f64 [%r{{[0-9]+}}], %fd{{[0-9]+}} ; PTX32: ret -; PTX64: st.f64 [%rl{{[0-9]+}}], %fl{{[0-9]+}} +; PTX64: st.f64 [%rd{{[0-9]+}}], %fd{{[0-9]+}} ; PTX64: ret store double %a, double addrspace(0)* %ptr ret void diff --git a/test/CodeGen/NVPTX/surf-read-cuda.ll b/test/CodeGen/NVPTX/surf-read-cuda.ll new file mode 100644 index 000000000000..10a1ecc4c473 --- /dev/null +++ b/test/CodeGen/NVPTX/surf-read-cuda.ll @@ -0,0 +1,53 @@ +; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix=SM20 +; RUN: llc < %s -march=nvptx -mcpu=sm_30 | FileCheck %s --check-prefix=SM30 + +target triple = "nvptx-unknown-cuda" + +declare i32 @llvm.nvvm.suld.1d.i32.trap(i64, i32) +declare i64 @llvm.nvvm.texsurf.handle.internal.p1i64(i64 addrspace(1)*) + + +; SM20-LABEL: .entry foo +; SM30-LABEL: .entry foo +define void @foo(i64 %img, float* %red, i32 %idx) { +; SM20: ld.param.u64 %rd[[SURFREG:[0-9]+]], [foo_param_0]; +; SM20: suld.b.1d.b32.trap {%r[[RED:[0-9]+]]}, [%rd[[SURFREG]], {%r{{[0-9]+}}}] +; SM30: ld.param.u64 %rd[[SURFREG:[0-9]+]], [foo_param_0]; +; SM30: suld.b.1d.b32.trap {%r[[RED:[0-9]+]]}, [%rd[[SURFREG]], {%r{{[0-9]+}}}] + %val = tail call i32 @llvm.nvvm.suld.1d.i32.trap(i64 %img, i32 %idx) +; SM20: cvt.rn.f32.s32 %f[[REDF:[0-9]+]], %r[[RED]] +; SM30: cvt.rn.f32.s32 %f[[REDF:[0-9]+]], %r[[RED]] + %ret = sitofp i32 %val to float +; SM20: st.f32 [%r{{[0-9]+}}], %f[[REDF]] +; SM30: st.f32 [%r{{[0-9]+}}], %f[[REDF]] + store float %ret, float* %red + ret void +} + +@surf0 = internal addrspace(1) global i64 0, align 8 + +; SM20-LABEL: .entry bar +; SM30-LABEL: .entry bar +define void @bar(float* %red, i32 %idx) { +; SM30: mov.u64 %rd[[SURFHANDLE:[0-9]+]], surf0 + %surfHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1i64(i64 addrspace(1)* @surf0) +; SM20: suld.b.1d.b32.trap {%r[[RED:[0-9]+]]}, [surf0, {%r{{[0-9]+}}}] +; SM30: suld.b.1d.b32.trap {%r[[RED:[0-9]+]]}, [%rd[[SURFHANDLE]], {%r{{[0-9]+}}}] + %val = tail call i32 @llvm.nvvm.suld.1d.i32.trap(i64 %surfHandle, i32 %idx) +; SM20: cvt.rn.f32.s32 %f[[REDF:[0-9]+]], %r[[RED]] +; SM30: cvt.rn.f32.s32 %f[[REDF:[0-9]+]], %r[[RED]] + %ret = sitofp i32 %val to float +; SM20: st.f32 [%r{{[0-9]+}}], %f[[REDF]] +; SM30: st.f32 [%r{{[0-9]+}}], %f[[REDF]] + store float %ret, float* %red + ret void +} + + + + +!nvvm.annotations = !{!1, !2, !3} +!1 = metadata !{void (i64, float*, i32)* @foo, metadata !"kernel", i32 1} +!2 = metadata !{void (float*, i32)* @bar, metadata !"kernel", i32 1} +!3 = metadata !{i64 addrspace(1)* @surf0, metadata !"surface", i32 1} + diff --git a/test/CodeGen/NVPTX/surf-read.ll b/test/CodeGen/NVPTX/surf-read.ll new file mode 100644 index 000000000000..a69d03efe0d2 --- /dev/null +++ b/test/CodeGen/NVPTX/surf-read.ll @@ -0,0 +1,20 @@ +; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s + +target triple = "nvptx-unknown-nvcl" + +declare i32 @llvm.nvvm.suld.1d.i32.trap(i64, i32) + +; CHECK: .entry foo +define void @foo(i64 %img, float* %red, i32 %idx) { +; CHECK: suld.b.1d.b32.trap {%r[[RED:[0-9]+]]}, [foo_param_0, {%r{{[0-9]+}}}] + %val = tail call i32 @llvm.nvvm.suld.1d.i32.trap(i64 %img, i32 %idx) +; CHECK: cvt.rn.f32.s32 %f[[REDF:[0-9]+]], %r[[RED]] + %ret = sitofp i32 %val to float +; CHECK: st.f32 [%r{{[0-9]+}}], %f[[REDF]] + store float %ret, float* %red + ret void +} + +!nvvm.annotations = !{!1, !2} +!1 = metadata !{void (i64, float*, i32)* @foo, metadata !"kernel", i32 1} +!2 = metadata !{void (i64, float*, i32)* @foo, metadata !"rdwrimage", i32 0} diff --git a/test/CodeGen/NVPTX/surf-write-cuda.ll b/test/CodeGen/NVPTX/surf-write-cuda.ll new file mode 100644 index 000000000000..654c47f46957 --- /dev/null +++ b/test/CodeGen/NVPTX/surf-write-cuda.ll @@ -0,0 +1,42 @@ +; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix=SM20 +; RUN: llc < %s -march=nvptx -mcpu=sm_30 | FileCheck %s --check-prefix=SM30 + +target triple = "nvptx-unknown-cuda" + +declare void @llvm.nvvm.sust.b.1d.i32.trap(i64, i32, i32) +declare i64 @llvm.nvvm.texsurf.handle.internal.p1i64(i64 addrspace(1)*) + + +; SM20-LABEL: .entry foo +; SM30-LABEL: .entry foo +define void @foo(i64 %img, i32 %val, i32 %idx) { +; SM20: ld.param.u64 %rd[[SURFREG:[0-9]+]], [foo_param_0]; +; SM20: sust.b.1d.b32.trap [%rd[[SURFREG]], {%r{{[0-9]+}}}], {%r{{[0-9]+}}} +; SM30: ld.param.u64 %rd[[SURFREG:[0-9]+]], [foo_param_0]; +; SM30: sust.b.1d.b32.trap [%rd[[SURFREG]], {%r{{[0-9]+}}}], {%r{{[0-9]+}}} + tail call void @llvm.nvvm.sust.b.1d.i32.trap(i64 %img, i32 %idx, i32 %val) + ret void +} + + +@surf0 = internal addrspace(1) global i64 0, align 8 + + + +; SM20-LABEL: .entry bar +; SM30-LABEL: .entry bar +define void @bar(i32 %val, i32 %idx) { +; SM30: mov.u64 %rd[[SURFHANDLE:[0-9]+]], surf0 + %surfHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1i64(i64 addrspace(1)* @surf0) +; SM20: sust.b.1d.b32.trap [surf0, {%r{{[0-9]+}}}], {%r{{[0-9]+}}} +; SM30: sust.b.1d.b32.trap [%rd[[SURFREG]], {%r{{[0-9]+}}}], {%r{{[0-9]+}}} + tail call void @llvm.nvvm.sust.b.1d.i32.trap(i64 %surfHandle, i32 %idx, i32 %val) + ret void +} + + +!nvvm.annotations = !{!1, !2, !3} +!1 = metadata !{void (i64, i32, i32)* @foo, metadata !"kernel", i32 1} +!2 = metadata !{void (i32, i32)* @bar, metadata !"kernel", i32 1} +!3 = metadata !{i64 addrspace(1)* @surf0, metadata !"surface", i32 1} + diff --git a/test/CodeGen/NVPTX/surf-write.ll b/test/CodeGen/NVPTX/surf-write.ll new file mode 100644 index 000000000000..880231f96599 --- /dev/null +++ b/test/CodeGen/NVPTX/surf-write.ll @@ -0,0 +1,16 @@ +; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s + +target triple = "nvptx-unknown-nvcl" + +declare void @llvm.nvvm.sust.b.1d.i32.trap(i64, i32, i32) + +; CHECK: .entry foo +define void @foo(i64 %img, i32 %val, i32 %idx) { +; CHECK: sust.b.1d.b32.trap [foo_param_0, {%r{{[0-9]+}}}], {%r{{[0-9]+}}} + tail call void @llvm.nvvm.sust.b.1d.i32.trap(i64 %img, i32 %idx, i32 %val) + ret void +} + +!nvvm.annotations = !{!1, !2} +!1 = metadata !{void (i64, i32, i32)* @foo, metadata !"kernel", i32 1} +!2 = metadata !{void (i64, i32, i32)* @foo, metadata !"wroimage", i32 0} diff --git a/test/CodeGen/NVPTX/symbol-naming.ll b/test/CodeGen/NVPTX/symbol-naming.ll new file mode 100644 index 000000000000..bd1333f1c4e6 --- /dev/null +++ b/test/CodeGen/NVPTX/symbol-naming.ll @@ -0,0 +1,31 @@ +; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix=PTX32 +; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefix=PTX64 + +; Verify that the NVPTX target removes invalid symbol names prior to emitting +; PTX. + +; PTX32-NOT: .str +; PTX64-NOT: .str + +; PTX32-DAG: _$_str1 +; PTX32-DAG: _$_str + +; PTX64-DAG: _$_str1 +; PTX64-DAG: _$_str + +target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64" +target triple = "nvptx64-unknown-unknown" + + +@.str = private unnamed_addr constant [13 x i8] c"%d %f %c %d\0A\00", align 1 +@_$_str = private unnamed_addr constant [13 x i8] c"%d %f %c %d\0A\00", align 1 + + +; Function Attrs: nounwind +define void @foo(i32 %a, float %b, i8 signext %c, i32 %e) { +entry: + %call = call i32 (i8*, ...)* @printf(i8* getelementptr inbounds ([13 x i8]* @.str, i32 0, i32 0)) + ret void +} + +declare i32 @printf(i8*, ...) diff --git a/test/CodeGen/NVPTX/tex-read-cuda.ll b/test/CodeGen/NVPTX/tex-read-cuda.ll new file mode 100644 index 000000000000..ee0cefa919b1 --- /dev/null +++ b/test/CodeGen/NVPTX/tex-read-cuda.ll @@ -0,0 +1,46 @@ +; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix=SM20 +; RUN: llc < %s -march=nvptx -mcpu=sm_30 | FileCheck %s --check-prefix=SM30 + + +target triple = "nvptx-unknown-cuda" + +declare { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64, i32) +declare i64 @llvm.nvvm.texsurf.handle.internal.p1i64(i64 addrspace(1)*) + +; SM20-LABEL: .entry foo +; SM30-LABEL: .entry foo +define void @foo(i64 %img, float* %red, i32 %idx) { +; SM20: ld.param.u64 %rd[[TEXREG:[0-9]+]], [foo_param_0]; +; SM20: tex.1d.v4.f32.s32 {%f[[RED:[0-9]+]], %f[[GREEN:[0-9]+]], %f[[BLUE:[0-9]+]], %f[[ALPHA:[0-9]+]]}, [%rd[[TEXREG]], {%r{{[0-9]+}}}] +; SM30: ld.param.u64 %rd[[TEXREG:[0-9]+]], [foo_param_0]; +; SM30: tex.1d.v4.f32.s32 {%f[[RED:[0-9]+]], %f[[GREEN:[0-9]+]], %f[[BLUE:[0-9]+]], %f[[ALPHA:[0-9]+]]}, [%rd[[TEXREG]], {%r{{[0-9]+}}}] + %val = tail call { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64 %img, i32 %idx) + %ret = extractvalue { float, float, float, float } %val, 0 +; SM20: st.f32 [%r{{[0-9]+}}], %f[[RED]] +; SM30: st.f32 [%r{{[0-9]+}}], %f[[RED]] + store float %ret, float* %red + ret void +} + + +@tex0 = internal addrspace(1) global i64 0, align 8 + +; SM20-LABEL: .entry bar +; SM30-LABEL: .entry bar +define void @bar(float* %red, i32 %idx) { +; SM30: mov.u64 %rd[[TEXHANDLE:[0-9]+]], tex0 + %texHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1i64(i64 addrspace(1)* @tex0) +; SM20: tex.1d.v4.f32.s32 {%f[[RED:[0-9]+]], %f[[GREEN:[0-9]+]], %f[[BLUE:[0-9]+]], %f[[ALPHA:[0-9]+]]}, [tex0, {%r{{[0-9]+}}}] +; SM30: tex.1d.v4.f32.s32 {%f[[RED:[0-9]+]], %f[[GREEN:[0-9]+]], %f[[BLUE:[0-9]+]], %f[[ALPHA:[0-9]+]]}, [%rd[[TEXHANDLE]], {%r{{[0-9]+}}}] + %val = tail call { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64 %texHandle, i32 %idx) + %ret = extractvalue { float, float, float, float } %val, 0 +; SM20: st.f32 [%r{{[0-9]+}}], %f[[RED]] +; SM30: st.f32 [%r{{[0-9]+}}], %f[[RED]] + store float %ret, float* %red + ret void +} + +!nvvm.annotations = !{!1, !2, !3} +!1 = metadata !{void (i64, float*, i32)* @foo, metadata !"kernel", i32 1} +!2 = metadata !{void (float*, i32)* @bar, metadata !"kernel", i32 1} +!3 = metadata !{i64 addrspace(1)* @tex0, metadata !"texture", i32 1} diff --git a/test/CodeGen/NVPTX/tex-read.ll b/test/CodeGen/NVPTX/tex-read.ll new file mode 100644 index 000000000000..55e4bfc9e453 --- /dev/null +++ b/test/CodeGen/NVPTX/tex-read.ll @@ -0,0 +1,20 @@ +; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s + +target triple = "nvptx-unknown-nvcl" + +declare { float, float, float, float } @llvm.nvvm.tex.1d.v4f32.s32(i64, i64, i32) + +; CHECK: .entry foo +define void @foo(i64 %img, i64 %sampler, float* %red, i32 %idx) { +; CHECK: tex.1d.v4.f32.s32 {%f[[RED:[0-9]+]], %f[[GREEN:[0-9]+]], %f[[BLUE:[0-9]+]], %f[[ALPHA:[0-9]+]]}, [foo_param_0, foo_param_1, {%r{{[0-9]+}}}] + %val = tail call { float, float, float, float } @llvm.nvvm.tex.1d.v4f32.s32(i64 %img, i64 %sampler, i32 %idx) + %ret = extractvalue { float, float, float, float } %val, 0 +; CHECK: st.f32 [%r{{[0-9]+}}], %f[[RED]] + store float %ret, float* %red + ret void +} + +!nvvm.annotations = !{!1, !2, !3} +!1 = metadata !{void (i64, i64, float*, i32)* @foo, metadata !"kernel", i32 1} +!2 = metadata !{void (i64, i64, float*, i32)* @foo, metadata !"rdoimage", i32 0} +!3 = metadata !{void (i64, i64, float*, i32)* @foo, metadata !"sampler", i32 1} diff --git a/test/CodeGen/NVPTX/texsurf-queries.ll b/test/CodeGen/NVPTX/texsurf-queries.ll new file mode 100644 index 000000000000..c7637ccff77a --- /dev/null +++ b/test/CodeGen/NVPTX/texsurf-queries.ll @@ -0,0 +1,103 @@ +; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix=SM20 +; RUN: llc < %s -march=nvptx -mcpu=sm_30 | FileCheck %s --check-prefix=SM30 + +target triple = "nvptx-unknown-cuda" + +@tex0 = internal addrspace(1) global i64 0, align 8 +@surf0 = internal addrspace(1) global i64 0, align 8 + +declare i32 @llvm.nvvm.txq.width(i64) +declare i32 @llvm.nvvm.txq.height(i64) +declare i32 @llvm.nvvm.suq.width(i64) +declare i32 @llvm.nvvm.suq.height(i64) +declare i64 @llvm.nvvm.texsurf.handle.internal.p1i64(i64 addrspace(1)*) + + +; SM20-LABEL: @t0 +; SM30-LABEL: @t0 +define i32 @t0(i64 %texHandle) { +; SM20: txq.width.b32 +; SM30: txq.width.b32 + %width = tail call i32 @llvm.nvvm.txq.width(i64 %texHandle) + ret i32 %width +} + +; SM20-LABEL: @t1 +; SM30-LABEL: @t1 +define i32 @t1() { +; SM30: mov.u64 %rd[[HANDLE:[0-9]+]], tex0 + %texHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1i64(i64 addrspace(1)* @tex0) +; SM20: txq.width.b32 %r{{[0-9]+}}, [tex0] +; SM30: txq.width.b32 %r{{[0-9]+}}, [%rd[[HANDLE:[0-9]+]]] + %width = tail call i32 @llvm.nvvm.txq.width(i64 %texHandle) + ret i32 %width +} + + +; SM20-LABEL: @t2 +; SM30-LABEL: @t2 +define i32 @t2(i64 %texHandle) { +; SM20: txq.height.b32 +; SM30: txq.height.b32 + %height = tail call i32 @llvm.nvvm.txq.height(i64 %texHandle) + ret i32 %height +} + +; SM20-LABEL: @t3 +; SM30-LABEL: @t3 +define i32 @t3() { +; SM30: mov.u64 %rd[[HANDLE:[0-9]+]], tex0 + %texHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1i64(i64 addrspace(1)* @tex0) +; SM20: txq.height.b32 %r{{[0-9]+}}, [tex0] +; SM30: txq.height.b32 %r{{[0-9]+}}, [%rd[[HANDLE:[0-9]+]]] + %height = tail call i32 @llvm.nvvm.txq.height(i64 %texHandle) + ret i32 %height +} + + +; SM20-LABEL: @s0 +; SM30-LABEL: @s0 +define i32 @s0(i64 %surfHandle) { +; SM20: suq.width.b32 +; SM30: suq.width.b32 + %width = tail call i32 @llvm.nvvm.suq.width(i64 %surfHandle) + ret i32 %width +} + +; SM20-LABEL: @s1 +; SM30-LABEL: @s1 +define i32 @s1() { +; SM30: mov.u64 %rd[[HANDLE:[0-9]+]], surf0 + %surfHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1i64(i64 addrspace(1)* @surf0) +; SM20: suq.width.b32 %r{{[0-9]+}}, [surf0] +; SM30: suq.width.b32 %r{{[0-9]+}}, [%rd[[HANDLE:[0-9]+]]] + %width = tail call i32 @llvm.nvvm.suq.width(i64 %surfHandle) + ret i32 %width +} + + +; SM20-LABEL: @s2 +; SM30-LABEL: @s2 +define i32 @s2(i64 %surfHandle) { +; SM20: suq.height.b32 +; SM30: suq.height.b32 + %height = tail call i32 @llvm.nvvm.suq.height(i64 %surfHandle) + ret i32 %height +} + +; SM20-LABEL: @s3 +; SM30-LABEL: @s3 +define i32 @s3() { +; SM30: mov.u64 %rd[[HANDLE:[0-9]+]], surf0 + %surfHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1i64(i64 addrspace(1)* @surf0) +; SM20: suq.height.b32 %r{{[0-9]+}}, [surf0] +; SM30: suq.height.b32 %r{{[0-9]+}}, [%rd[[HANDLE:[0-9]+]]] + %height = tail call i32 @llvm.nvvm.suq.height(i64 %surfHandle) + ret i32 %height +} + + + +!nvvm.annotations = !{!1, !2} +!1 = metadata !{i64 addrspace(1)* @tex0, metadata !"texture", i32 1} +!2 = metadata !{i64 addrspace(1)* @surf0, metadata !"surface", i32 1} diff --git a/test/CodeGen/NVPTX/vec-param-load.ll b/test/CodeGen/NVPTX/vec-param-load.ll index a384348a6590..4193ac4085cc 100644 --- a/test/CodeGen/NVPTX/vec-param-load.ll +++ b/test/CodeGen/NVPTX/vec-param-load.ll @@ -5,9 +5,9 @@ target datalayout = "e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f3 define <16 x float> @foo(<16 x float> %a) { ; Make sure we index into vectors properly -; CHECK: ld.param.v4.f32 {%f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}}, [foo_param_0]; -; CHECK: ld.param.v4.f32 {%f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}}, [foo_param_0+16]; -; CHECK: ld.param.v4.f32 {%f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}}, [foo_param_0+32]; ; CHECK: ld.param.v4.f32 {%f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}}, [foo_param_0+48]; +; CHECK: ld.param.v4.f32 {%f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}}, [foo_param_0+32]; +; CHECK: ld.param.v4.f32 {%f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}}, [foo_param_0+16]; +; CHECK: ld.param.v4.f32 {%f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}}, [foo_param_0]; ret <16 x float> %a } diff --git a/test/CodeGen/NVPTX/vector-call.ll b/test/CodeGen/NVPTX/vector-call.ll new file mode 100644 index 000000000000..a03d7fd41914 --- /dev/null +++ b/test/CodeGen/NVPTX/vector-call.ll @@ -0,0 +1,12 @@ +; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s + +target triple = "nvptx-unknown-cuda" + +declare void @bar(<4 x i32>) + +; CHECK-LABEL @foo +define void @foo(<4 x i32> %a) { +; CHECK: st.param.v4.b32 + tail call void @bar(<4 x i32> %a) + ret void +} diff --git a/test/CodeGen/NVPTX/weak-global.ll b/test/CodeGen/NVPTX/weak-global.ll new file mode 100644 index 000000000000..2bef4c5228a9 --- /dev/null +++ b/test/CodeGen/NVPTX/weak-global.ll @@ -0,0 +1,9 @@ +; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s + +; CHECK: .weak .global .align 4 .u32 g +@g = common addrspace(1) global i32 zeroinitializer + +define i32 @func0() { + %val = load i32 addrspace(1)* @g + ret i32 %val +} diff --git a/test/CodeGen/NVPTX/weak-linkage.ll b/test/CodeGen/NVPTX/weak-linkage.ll new file mode 100644 index 000000000000..7a1335783642 --- /dev/null +++ b/test/CodeGen/NVPTX/weak-linkage.ll @@ -0,0 +1,12 @@ +; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s + + +; CHECK: .weak .func foo +define weak void @foo() { + ret void +} + +; CHECK: .visible .func bar +define void @bar() { + ret void +} |