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.
143 lines
4.3 KiB
LLVM
143 lines
4.3 KiB
LLVM
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
|
|
; RUN: llc < %s -mcpu=sm_80 -mattr=+ptx70 | FileCheck --check-prefixes=CHECK %s
|
|
; RUN: %if ptxas-11.8 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx70 | %ptxas-verify -arch=sm_80 %}
|
|
|
|
target triple = "nvptx-nvidia-cuda"
|
|
|
|
declare float @llvm.nvvm.fabs.f32(float)
|
|
declare float @llvm.nvvm.fabs.ftz.f32(float)
|
|
declare double @llvm.nvvm.fabs.f64(double)
|
|
declare half @llvm.nvvm.fabs.f16(half)
|
|
declare half @llvm.nvvm.fabs.ftz.f16(half)
|
|
declare <2 x half> @llvm.nvvm.fabs.v2f16(<2 x half>)
|
|
declare <2 x half> @llvm.nvvm.fabs.ftz.v2f16(<2 x half>)
|
|
declare bfloat @llvm.nvvm.fabs.bf16(bfloat)
|
|
declare <2 x bfloat> @llvm.nvvm.fabs.v2bf16(<2 x bfloat>)
|
|
|
|
|
|
define float @fabs_float(float %a) {
|
|
; CHECK-LABEL: fabs_float(
|
|
; CHECK: {
|
|
; CHECK-NEXT: .reg .b32 %r<3>;
|
|
; CHECK-EMPTY:
|
|
; CHECK-NEXT: // %bb.0:
|
|
; CHECK-NEXT: ld.param.b32 %r1, [fabs_float_param_0];
|
|
; CHECK-NEXT: abs.f32 %r2, %r1;
|
|
; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
|
|
; CHECK-NEXT: ret;
|
|
%ret = call float @llvm.nvvm.fabs.f32(float %a)
|
|
ret float %ret
|
|
}
|
|
|
|
define float @fabs_float_ftz(float %a) {
|
|
; CHECK-LABEL: fabs_float_ftz(
|
|
; CHECK: {
|
|
; CHECK-NEXT: .reg .b32 %r<3>;
|
|
; CHECK-EMPTY:
|
|
; CHECK-NEXT: // %bb.0:
|
|
; CHECK-NEXT: ld.param.b32 %r1, [fabs_float_ftz_param_0];
|
|
; CHECK-NEXT: abs.ftz.f32 %r2, %r1;
|
|
; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
|
|
; CHECK-NEXT: ret;
|
|
%ret = call float @llvm.nvvm.fabs.ftz.f32(float %a)
|
|
ret float %ret
|
|
}
|
|
|
|
define double @fabs_double(double %a) {
|
|
; CHECK-LABEL: fabs_double(
|
|
; CHECK: {
|
|
; CHECK-NEXT: .reg .b64 %rd<3>;
|
|
; CHECK-EMPTY:
|
|
; CHECK-NEXT: // %bb.0:
|
|
; CHECK-NEXT: ld.param.b64 %rd1, [fabs_double_param_0];
|
|
; CHECK-NEXT: abs.f64 %rd2, %rd1;
|
|
; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
|
|
; CHECK-NEXT: ret;
|
|
%ret = call double @llvm.nvvm.fabs.f64(double %a)
|
|
ret double %ret
|
|
}
|
|
|
|
define half @fabs_half(half %a) {
|
|
; CHECK-LABEL: fabs_half(
|
|
; CHECK: {
|
|
; CHECK-NEXT: .reg .b16 %rs<3>;
|
|
; CHECK-EMPTY:
|
|
; CHECK-NEXT: // %bb.0:
|
|
; CHECK-NEXT: ld.param.b16 %rs1, [fabs_half_param_0];
|
|
; CHECK-NEXT: abs.f16 %rs2, %rs1;
|
|
; CHECK-NEXT: st.param.b16 [func_retval0], %rs2;
|
|
; CHECK-NEXT: ret;
|
|
%ret = call half @llvm.nvvm.fabs.f16(half %a)
|
|
ret half %ret
|
|
}
|
|
|
|
define half @fabs_half_ftz(half %a) {
|
|
; CHECK-LABEL: fabs_half_ftz(
|
|
; CHECK: {
|
|
; CHECK-NEXT: .reg .b16 %rs<3>;
|
|
; CHECK-EMPTY:
|
|
; CHECK-NEXT: // %bb.0:
|
|
; CHECK-NEXT: ld.param.b16 %rs1, [fabs_half_ftz_param_0];
|
|
; CHECK-NEXT: abs.ftz.f16 %rs2, %rs1;
|
|
; CHECK-NEXT: st.param.b16 [func_retval0], %rs2;
|
|
; CHECK-NEXT: ret;
|
|
%ret = call half @llvm.nvvm.fabs.ftz.f16(half %a)
|
|
ret half %ret
|
|
}
|
|
|
|
define <2 x half> @fabs_v2half(<2 x half> %a) {
|
|
; CHECK-LABEL: fabs_v2half(
|
|
; CHECK: {
|
|
; CHECK-NEXT: .reg .b32 %r<3>;
|
|
; CHECK-EMPTY:
|
|
; CHECK-NEXT: // %bb.0:
|
|
; CHECK-NEXT: ld.param.b32 %r1, [fabs_v2half_param_0];
|
|
; CHECK-NEXT: abs.f16x2 %r2, %r1;
|
|
; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
|
|
; CHECK-NEXT: ret;
|
|
%ret = call <2 x half> @llvm.nvvm.fabs.v2f16(<2 x half> %a)
|
|
ret <2 x half> %ret
|
|
}
|
|
|
|
define <2 x half> @fabs_v2half_ftz(<2 x half> %a) {
|
|
; CHECK-LABEL: fabs_v2half_ftz(
|
|
; CHECK: {
|
|
; CHECK-NEXT: .reg .b32 %r<3>;
|
|
; CHECK-EMPTY:
|
|
; CHECK-NEXT: // %bb.0:
|
|
; CHECK-NEXT: ld.param.b32 %r1, [fabs_v2half_ftz_param_0];
|
|
; CHECK-NEXT: abs.ftz.f16x2 %r2, %r1;
|
|
; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
|
|
; CHECK-NEXT: ret;
|
|
%ret = call <2 x half> @llvm.nvvm.fabs.ftz.v2f16(<2 x half> %a)
|
|
ret <2 x half> %ret
|
|
}
|
|
|
|
define bfloat @fabs_bf16(bfloat %a) {
|
|
; CHECK-LABEL: fabs_bf16(
|
|
; CHECK: {
|
|
; CHECK-NEXT: .reg .b16 %rs<3>;
|
|
; CHECK-EMPTY:
|
|
; CHECK-NEXT: // %bb.0:
|
|
; CHECK-NEXT: ld.param.b16 %rs1, [fabs_bf16_param_0];
|
|
; CHECK-NEXT: abs.bf16 %rs2, %rs1;
|
|
; CHECK-NEXT: st.param.b16 [func_retval0], %rs2;
|
|
; CHECK-NEXT: ret;
|
|
%ret = call bfloat @llvm.nvvm.fabs.bf16(bfloat %a)
|
|
ret bfloat %ret
|
|
}
|
|
|
|
define <2 x bfloat> @fabs_v2bf16(<2 x bfloat> %a) {
|
|
; CHECK-LABEL: fabs_v2bf16(
|
|
; CHECK: {
|
|
; CHECK-NEXT: .reg .b32 %r<3>;
|
|
; CHECK-EMPTY:
|
|
; CHECK-NEXT: // %bb.0:
|
|
; CHECK-NEXT: ld.param.b32 %r1, [fabs_v2bf16_param_0];
|
|
; CHECK-NEXT: abs.bf16x2 %r2, %r1;
|
|
; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
|
|
; CHECK-NEXT: ret;
|
|
%ret = call <2 x bfloat> @llvm.nvvm.fabs.v2bf16(<2 x bfloat> %a)
|
|
ret <2 x bfloat> %ret
|
|
}
|