76c9bfefa4
These classes are redundant, as the untyped "Int" classes can be used for all float operations. This change is intended to be as minimal as possible and leaves the many potential simplifications and refactors this exposes as future work.
44 lines
1.5 KiB
LLVM
44 lines
1.5 KiB
LLVM
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
|
|
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s
|
|
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
|
|
|
|
define float @test(float %x) {
|
|
; CHECK-LABEL: test(
|
|
; CHECK: {
|
|
; CHECK-NEXT: .reg .b32 %r<5>;
|
|
; CHECK-EMPTY:
|
|
; CHECK-NEXT: // %bb.0: // %entry
|
|
; CHECK-NEXT: ld.param.b32 %r2, [test_param_0];
|
|
; CHECK-NEXT: // begin inline asm
|
|
; CHECK-NEXT: ex2.approx.ftz.f32 %r4, %r2;
|
|
; CHECK-NEXT: // end inline asm
|
|
; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
|
|
; CHECK-NEXT: ret;
|
|
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 {
|
|
; CHECK-LABEL: foo(
|
|
; CHECK: {
|
|
; CHECK-NEXT: .reg .pred %p<2>;
|
|
; CHECK-NEXT: .reg .b16 %rs<3>;
|
|
; CHECK-NEXT: .reg .b32 %r<4>;
|
|
; CHECK-EMPTY:
|
|
; CHECK-NEXT: // %bb.0: // %entry
|
|
; CHECK-NEXT: ld.param.b8 %rs1, [foo_param_0];
|
|
; CHECK-NEXT: and.b16 %rs2, %rs1, 1;
|
|
; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0;
|
|
; CHECK-NEXT: ld.param.b32 %r2, [foo_param_1];
|
|
; CHECK-NEXT: ld.param.b32 %r3, [foo_param_2];
|
|
; CHECK-NEXT: // begin inline asm
|
|
; CHECK-NEXT: selp.b32 %r1, %r2, %r3, %p1;
|
|
; CHECK-NEXT: // end inline asm
|
|
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
|
|
; CHECK-NEXT: ret;
|
|
entry:
|
|
%0 = tail call i32 asm "selp.b32 $0, $1, $2, $3;", "=r,r,r,b"(i32 %a, i32 %b, i1 %cond)
|
|
ret i32 %0
|
|
}
|