blob: 6f4eb222e0b380b60b456c06488a86be6eaac60f (
plain)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
|
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc < %s -mcpu=sm_75 -mattr=+ptx70 | FileCheck %s
; RUN: %if ptxas-11.0 %{ llc < %s -mcpu=sm_75 -mattr=+ptx70 | %ptxas-verify -arch=sm_75 %}
target triple = "nvptx64-nvidia-cuda"
define float @test1(float %in) local_unnamed_addr {
; CHECK-LABEL: test1(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [test1_param_0];
; CHECK-NEXT: tanh.approx.f32 %r2, %r1;
; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
; CHECK-NEXT: ret;
%call = call afn float @llvm.tanh.f32(float %in)
ret float %call
}
define half @test2(half %in) local_unnamed_addr {
; CHECK-LABEL: test2(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<3>;
; CHECK-NEXT: .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [test2_param_0];
; CHECK-NEXT: cvt.f32.f16 %r1, %rs1;
; CHECK-NEXT: tanh.approx.f32 %r2, %r1;
; CHECK-NEXT: cvt.rn.f16.f32 %rs2, %r2;
; CHECK-NEXT: st.param.b16 [func_retval0], %rs2;
; CHECK-NEXT: ret;
%call = call afn half @llvm.tanh.f16(half %in)
ret half %call
}
declare float @llvm.tanh.f32(float)
declare half @llvm.tanh.f16(half)
|
