Alex MacLean 76c9bfefa4
[NVPTX] Remove Float register classes (#140487)
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.
2025-05-21 11:33:57 -07:00

59 lines
1.7 KiB
LLVM

; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc < %s -mtriple=nvptx64 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 | %ptxas-verify %}
target triple = "nvptx64-nvidia-cuda"
;; Check if fneg (fdiv 1, X) lowers to fneg (rcp.rn X).
define double @test1(double %in) {
; CHECK-LABEL: test1(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b64 %rd1, [test1_param_0];
; CHECK-NEXT: rcp.rn.f64 %rd2, %rd1;
; CHECK-NEXT: neg.f64 %rd3, %rd2;
; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
; CHECK-NEXT: ret;
%div = fdiv double 1.000000e+00, %in
%neg = fsub double -0.000000e+00, %div
ret double %neg
}
;; Check if fdiv -1, X lowers to fneg (rcp.rn X).
define double @test2(double %in) {
; CHECK-LABEL: test2(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b64 %rd1, [test2_param_0];
; CHECK-NEXT: rcp.rn.f64 %rd2, %rd1;
; CHECK-NEXT: neg.f64 %rd3, %rd2;
; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
; CHECK-NEXT: ret;
%div = fdiv double -1.000000e+00, %in
ret double %div
}
;; Check if fdiv 1, (fneg X) lowers to fneg (rcp.rn X).
define double @test3(double %in) {
; CHECK-LABEL: test3(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b64 %rd1, [test3_param_0];
; CHECK-NEXT: rcp.rn.f64 %rd2, %rd1;
; CHECK-NEXT: neg.f64 %rd3, %rd2;
; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
; CHECK-NEXT: ret;
%neg = fsub double -0.000000e+00, %in
%div = fdiv double 1.000000e+00, %neg
ret double %div
}