e84614833e
This commit adds NVPTX support for div.full PTX instruction with test under div.ll. [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#floating-point-instructions-div)
27 lines
1.1 KiB
LLVM
27 lines
1.1 KiB
LLVM
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
|
|
; RUN: llc < %s -march=nvptx64 | FileCheck %s
|
|
; RUN: %if ptxas %{ llc < %s -march=nvptx64 | %ptxas-verify %}
|
|
|
|
define float @div_full(float %a, float %b) {
|
|
; CHECK-LABEL: div_full(
|
|
; CHECK: {
|
|
; CHECK-NEXT: .reg .f32 %f<9>;
|
|
; CHECK-EMPTY:
|
|
; CHECK-NEXT: // %bb.0:
|
|
; CHECK-NEXT: ld.param.f32 %f1, [div_full_param_0];
|
|
; CHECK-NEXT: ld.param.f32 %f2, [div_full_param_1];
|
|
; CHECK-NEXT: div.full.f32 %f3, %f1, %f2;
|
|
; CHECK-NEXT: mov.f32 %f4, 0f40400000;
|
|
; CHECK-NEXT: div.full.f32 %f5, %f3, %f4;
|
|
; CHECK-NEXT: div.full.ftz.f32 %f6, %f5, %f2;
|
|
; CHECK-NEXT: mov.f32 %f7, 0f40800000;
|
|
; CHECK-NEXT: div.full.ftz.f32 %f8, %f6, %f7;
|
|
; CHECK-NEXT: st.param.f32 [func_retval0], %f8;
|
|
; CHECK-NEXT: ret;
|
|
%1 = call float @llvm.nvvm.div.full(float %a, float %b)
|
|
%2 = call float @llvm.nvvm.div.full(float %1, float 3.0)
|
|
%3 = call float @llvm.nvvm.div.full.ftz(float %2, float %b)
|
|
%4 = call float @llvm.nvvm.div.full.ftz(float %3, float 4.0)
|
|
ret float %4
|
|
}
|