132 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100a -mattr=+ptx86 | FileCheck %s3; RUN: %if ptxas-sm_100a && ptxas-isa-8.6 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100a -mattr=+ptx86 | %ptxas-verify -arch=sm_100a %}4 5declare float @llvm.nvvm.redux.sync.fmin(float, i32)6define float @redux_sync_fmin(float %src, i32 %mask) {7; CHECK-LABEL: redux_sync_fmin(8; CHECK: {9; CHECK-NEXT: .reg .b32 %r<4>;10; CHECK-EMPTY:11; CHECK-NEXT: // %bb.0:12; CHECK-NEXT: ld.param.b32 %r1, [redux_sync_fmin_param_0];13; CHECK-NEXT: ld.param.b32 %r2, [redux_sync_fmin_param_1];14; CHECK-NEXT: redux.sync.min.f32 %r3, %r1, %r2;15; CHECK-NEXT: st.param.b32 [func_retval0], %r3;16; CHECK-NEXT: ret;17 %val = call float @llvm.nvvm.redux.sync.fmin(float %src, i32 %mask)18 ret float %val19}20 21declare float @llvm.nvvm.redux.sync.fmin.abs(float, i32)22define float @redux_sync_fmin_abs(float %src, i32 %mask) {23; CHECK-LABEL: redux_sync_fmin_abs(24; CHECK: {25; CHECK-NEXT: .reg .b32 %r<4>;26; CHECK-EMPTY:27; CHECK-NEXT: // %bb.0:28; CHECK-NEXT: ld.param.b32 %r1, [redux_sync_fmin_abs_param_0];29; CHECK-NEXT: ld.param.b32 %r2, [redux_sync_fmin_abs_param_1];30; CHECK-NEXT: redux.sync.min.abs.f32 %r3, %r1, %r2;31; CHECK-NEXT: st.param.b32 [func_retval0], %r3;32; CHECK-NEXT: ret;33 %val = call float @llvm.nvvm.redux.sync.fmin.abs(float %src, i32 %mask)34 ret float %val35}36 37declare float @llvm.nvvm.redux.sync.fmin.NaN(float, i32)38define float @redux_sync_fmin_NaN(float %src, i32 %mask) {39; CHECK-LABEL: redux_sync_fmin_NaN(40; CHECK: {41; CHECK-NEXT: .reg .b32 %r<4>;42; CHECK-EMPTY:43; CHECK-NEXT: // %bb.0:44; CHECK-NEXT: ld.param.b32 %r1, [redux_sync_fmin_NaN_param_0];45; CHECK-NEXT: ld.param.b32 %r2, [redux_sync_fmin_NaN_param_1];46; CHECK-NEXT: redux.sync.min.NaN.f32 %r3, %r1, %r2;47; CHECK-NEXT: st.param.b32 [func_retval0], %r3;48; CHECK-NEXT: ret;49 %val = call float @llvm.nvvm.redux.sync.fmin.NaN(float %src, i32 %mask)50 ret float %val51}52 53declare float @llvm.nvvm.redux.sync.fmin.abs.NaN(float, i32)54define float @redux_sync_fmin_abs_NaN(float %src, i32 %mask) {55; CHECK-LABEL: redux_sync_fmin_abs_NaN(56; CHECK: {57; CHECK-NEXT: .reg .b32 %r<4>;58; CHECK-EMPTY:59; CHECK-NEXT: // %bb.0:60; CHECK-NEXT: ld.param.b32 %r1, [redux_sync_fmin_abs_NaN_param_0];61; CHECK-NEXT: ld.param.b32 %r2, [redux_sync_fmin_abs_NaN_param_1];62; CHECK-NEXT: redux.sync.min.abs.NaN.f32 %r3, %r1, %r2;63; CHECK-NEXT: st.param.b32 [func_retval0], %r3;64; CHECK-NEXT: ret;65 %val = call float @llvm.nvvm.redux.sync.fmin.abs.NaN(float %src, i32 %mask)66 ret float %val67}68 69declare float @llvm.nvvm.redux.sync.fmax(float, i32)70define float @redux_sync_fmax(float %src, i32 %mask) {71; CHECK-LABEL: redux_sync_fmax(72; CHECK: {73; CHECK-NEXT: .reg .b32 %r<4>;74; CHECK-EMPTY:75; CHECK-NEXT: // %bb.0:76; CHECK-NEXT: ld.param.b32 %r1, [redux_sync_fmax_param_0];77; CHECK-NEXT: ld.param.b32 %r2, [redux_sync_fmax_param_1];78; CHECK-NEXT: redux.sync.max.f32 %r3, %r1, %r2;79; CHECK-NEXT: st.param.b32 [func_retval0], %r3;80; CHECK-NEXT: ret;81 %val = call float @llvm.nvvm.redux.sync.fmax(float %src, i32 %mask)82 ret float %val83}84 85declare float @llvm.nvvm.redux.sync.fmax.abs(float, i32)86define float @redux_sync_fmax_abs(float %src, i32 %mask) {87; CHECK-LABEL: redux_sync_fmax_abs(88; CHECK: {89; CHECK-NEXT: .reg .b32 %r<4>;90; CHECK-EMPTY:91; CHECK-NEXT: // %bb.0:92; CHECK-NEXT: ld.param.b32 %r1, [redux_sync_fmax_abs_param_0];93; CHECK-NEXT: ld.param.b32 %r2, [redux_sync_fmax_abs_param_1];94; CHECK-NEXT: redux.sync.max.abs.f32 %r3, %r1, %r2;95; CHECK-NEXT: st.param.b32 [func_retval0], %r3;96; CHECK-NEXT: ret;97 %val = call float @llvm.nvvm.redux.sync.fmax.abs(float %src, i32 %mask)98 ret float %val99}100 101declare float @llvm.nvvm.redux.sync.fmax.NaN(float, i32)102define float @redux_sync_fmax_NaN(float %src, i32 %mask) {103; CHECK-LABEL: redux_sync_fmax_NaN(104; CHECK: {105; CHECK-NEXT: .reg .b32 %r<4>;106; CHECK-EMPTY:107; CHECK-NEXT: // %bb.0:108; CHECK-NEXT: ld.param.b32 %r1, [redux_sync_fmax_NaN_param_0];109; CHECK-NEXT: ld.param.b32 %r2, [redux_sync_fmax_NaN_param_1];110; CHECK-NEXT: redux.sync.max.NaN.f32 %r3, %r1, %r2;111; CHECK-NEXT: st.param.b32 [func_retval0], %r3;112; CHECK-NEXT: ret;113 %val = call float @llvm.nvvm.redux.sync.fmax.NaN(float %src, i32 %mask)114 ret float %val115}116 117declare float @llvm.nvvm.redux.sync.fmax.abs.NaN(float, i32)118define float @redux_sync_fmax_abs_NaN(float %src, i32 %mask) {119; CHECK-LABEL: redux_sync_fmax_abs_NaN(120; CHECK: {121; CHECK-NEXT: .reg .b32 %r<4>;122; CHECK-EMPTY:123; CHECK-NEXT: // %bb.0:124; CHECK-NEXT: ld.param.b32 %r1, [redux_sync_fmax_abs_NaN_param_0];125; CHECK-NEXT: ld.param.b32 %r2, [redux_sync_fmax_abs_NaN_param_1];126; CHECK-NEXT: redux.sync.max.abs.NaN.f32 %r3, %r1, %r2;127; CHECK-NEXT: st.param.b32 [func_retval0], %r3;128; CHECK-NEXT: ret;129 %val = call float @llvm.nvvm.redux.sync.fmax.abs.NaN(float %src, i32 %mask)130 ret float %val131}132