299 lines · plain
1// REQUIRES: x86-registered-target, nvptx-registered-target, amdgpu-registered-target2 3// By default CUDA uses -ffp-contract=fast, HIP uses -ffp-contract=fast-honor-pragmas.4// we should fuse multiply/add into fma instruction.5// In IR, fmul/fadd instructions with contract flag are emitted.6// In backend7// nvptx - assumes fast fp fuse option, which fuses8// mult/add insts disregarding contract flag and9// llvm.fmuladd intrinsics.10// amdgcn - assumes standard fp fuse option, which only11// fuses mult/add insts with contract flag and12// llvm.fmuladd intrinsics.13 14// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \15// RUN: -disable-llvm-passes -o - %s \16// RUN: | FileCheck -check-prefixes=COMMON,NV-ON %s17// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -S \18// RUN: -target-cpu gfx906 -disable-llvm-passes -o - -x hip %s \19// RUN: | FileCheck -check-prefixes=COMMON,AMD-ON %s20// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \21// RUN: -O3 -o - %s \22// RUN: | FileCheck -check-prefixes=COMMON,NV-OPT-FAST %s23// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -S \24// RUN: -O3 -target-cpu gfx906 -o - -x hip %s \25// RUN: | FileCheck -check-prefixes=COMMON,AMD-OPT-FASTSTD %s26 27// Check separate compile/backend steps corresponding to -save-temps.28 29// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \30// RUN: -O3 -disable-llvm-passes -target-cpu gfx906 -o %t.ll -x hip %s31// RUN: cat %t.ll | FileCheck -check-prefixes=COMMON,AMD-OPT-FAST-IR %s32// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -S \33// RUN: -O3 -target-cpu gfx906 -o - -x ir %t.ll \34// RUN: | FileCheck -check-prefixes=COMMON,AMD-OPT-FASTSTD %s35 36// Explicit -ffp-contract=fast37// In IR, fmul/fadd instructions with contract flag are emitted.38// In backend39// nvptx/amdgcn - assumes fast fp fuse option, which fuses40// mult/add insts disregarding contract flag and41// llvm.fmuladd intrinsics.42 43// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \44// RUN: -ffp-contract=fast -disable-llvm-passes -o - %s \45// RUN: | FileCheck -check-prefixes=COMMON,NV-ON %s46// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -S \47// RUN: -target-cpu gfx906 -disable-llvm-passes -o - -x hip %s \48// RUN: -ffp-contract=fast \49// RUN: | FileCheck -check-prefixes=COMMON,AMD-ON %s50// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \51// RUN: -O3 -o - %s \52// RUN: -ffp-contract=fast \53// RUN: | FileCheck -check-prefixes=COMMON,NV-OPT-FAST %s54// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -S \55// RUN: -O3 -target-cpu gfx906 -o - -x hip %s \56// RUN: -ffp-contract=fast \57// RUN: | FileCheck -check-prefixes=COMMON,AMD-OPT-FAST %s58 59// Check separate compile/backend steps corresponding to -save-temps.60// When input is IR, -ffp-contract has no effect. Backend uses default61// default FP fuse option.62 63// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \64// RUN: -ffp-contract=fast \65// RUN: -O3 -disable-llvm-passes -target-cpu gfx906 -o %t.ll -x hip %s66// RUN: cat %t.ll | FileCheck -check-prefixes=COMMON,AMD-OPT-FAST-IR %s67// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -S \68// RUN: -O3 -target-cpu gfx906 -o - -x ir %t.ll \69// RUN: | FileCheck -check-prefixes=COMMON,AMD-OPT-FASTSTD %s70 71// Explicit -ffp-contract=fast-honor-pragmas72// In IR, fmul/fadd instructions with contract flag are emitted.73// In backend74// nvptx/amdgcn - assumes standard fp fuse option, which only75// fuses mult/add insts with contract flag or76// llvm.fmuladd intrinsics.77 78// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \79// RUN: -ffp-contract=fast-honor-pragmas -disable-llvm-passes -o - %s \80// RUN: | FileCheck -check-prefixes=COMMON,NV-ON %s81// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -S \82// RUN: -target-cpu gfx906 -disable-llvm-passes -o - -x hip %s \83// RUN: -ffp-contract=fast-honor-pragmas \84// RUN: | FileCheck -check-prefixes=COMMON,AMD-ON %s85// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \86// RUN: -O3 -o - %s \87// RUN: -ffp-contract=fast-honor-pragmas \88// RUN: | FileCheck -check-prefixes=COMMON,NV-OPT-FASTSTD %s89// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -S \90// RUN: -O3 -target-cpu gfx906 -o - -x hip %s \91// RUN: -ffp-contract=fast-honor-pragmas \92// RUN: | FileCheck -check-prefixes=COMMON,AMD-OPT-FASTSTD %s93 94// Check separate compile/backend steps corresponding to -save-temps.95// When input is IR, -ffp-contract has no effect. Backend uses default96// default FP fuse option.97 98// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \99// RUN: -ffp-contract=fast-honor-pragmas \100// RUN: -O3 -disable-llvm-passes -target-cpu gfx906 -o %t.ll -x hip %s101// RUN: cat %t.ll | FileCheck -check-prefixes=COMMON,AMD-OPT-FAST-IR %s102// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -S \103// RUN: -O3 -target-cpu gfx906 -o - -x ir %t.ll \104// RUN: | FileCheck -check-prefixes=COMMON,AMD-OPT-FASTSTD %s105 106// Explicit -ffp-contract=on -- fusing by front-end.107// In IR,108// mult/add in the same statement - llvm.fmuladd intrinsic emitted109// mult/add in different statement - fmul/fadd instructions without110// contract flag are emitted.111// In backend112// nvptx/amdgcn - assumes standard fp fuse option, which only113// fuses mult/add insts with contract flag or114// llvm.fmuladd intrinsics.115 116// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \117// RUN: -ffp-contract=on -disable-llvm-passes -o - %s \118// RUN: | FileCheck -check-prefixes=COMMON,NV-ON %s119// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -S \120// RUN: -target-cpu gfx906 -disable-llvm-passes -o - -x hip %s \121// RUN: -ffp-contract=on \122// RUN: | FileCheck -check-prefixes=COMMON,AMD-ON %s123// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \124// RUN: -O3 -o - %s \125// RUN: -ffp-contract=on \126// RUN: | FileCheck -check-prefixes=COMMON,NV-OPT-ON %s127// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -S \128// RUN: -O3 -target-cpu gfx906 -o - -x hip %s \129// RUN: -ffp-contract=on \130// RUN: | FileCheck -check-prefixes=COMMON,AMD-OPT-ON %s131 132// Check separate compile/backend steps corresponding to -save-temps.133 134// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \135// RUN: -ffp-contract=on \136// RUN: -O3 -disable-llvm-passes -target-cpu gfx906 -o %t.ll -x hip %s137// RUN: cat %t.ll | FileCheck -check-prefixes=COMMON,AMD-OPT-ON-IR %s138// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -S \139// RUN: -O3 -target-cpu gfx906 -o - -x ir %t.ll \140// RUN: | FileCheck -check-prefixes=COMMON,AMD-OPT-ON %s141 142// Explicit -ffp-contract=off should disable instruction fusing.143// In IR, fmul/fadd instructions without contract flag are emitted.144// In backend145// nvptx/amdgcn - assumes standard fp fuse option, which only146// fuses mult/add insts with contract flag or147// llvm.fmuladd intrinsics.148 149// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \150// RUN: -ffp-contract=off -disable-llvm-passes -o - %s \151// RUN: | FileCheck -check-prefixes=COMMON,NV-OFF %s152// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -S \153// RUN: -target-cpu gfx906 -disable-llvm-passes -o - -x hip %s \154// RUN: -ffp-contract=off \155// RUN: | FileCheck -check-prefixes=COMMON,AMD-OFF %s156// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \157// RUN: -O3 -o - %s \158// RUN: -ffp-contract=off \159// RUN: | FileCheck -check-prefixes=COMMON,NV-OPT-OFF %s160// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -S \161// RUN: -O3 -target-cpu gfx906 -o - -x hip %s \162// RUN: -ffp-contract=off \163// RUN: | FileCheck -check-prefixes=COMMON,AMD-OPT-OFF %s164 165// Check separate compile/backend steps corresponding to -save-temps.166 167// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \168// RUN: -ffp-contract=off \169// RUN: -O3 -disable-llvm-passes -target-cpu gfx906 -o %t.ll -x hip %s170// RUN: cat %t.ll | FileCheck -check-prefixes=COMMON,AMD-OPT-OFF-IR %s171// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -S \172// RUN: -O3 -target-cpu gfx906 -o - -x ir %t.ll \173// RUN: | FileCheck -check-prefixes=COMMON,AMD-OPT-OFF %s174 175#include "Inputs/cuda.h"176 177// Test multiply/add in the same statement, which can be emitted as FMA when178// fp-contract is on or fast.179__host__ __device__ float func(float a, float b, float c) { return a + b * c; }180// COMMON-LABEL: _Z4funcfff181// NV-ON: fma.rn.f32182// NV-ON-NEXT: st.param.b32183// AMD-ON: v_fmac_f32_e64184// AMD-ON-NEXT: s_setpc_b64185 186// NV-OFF: mul.rn.f32187// NV-OFF-NEXT: add.rn.f32188// NV-OFF-NEXT: st.param.b32189// AMD-OFF: v_mul_f32_e64190// AMD-OFF-NEXT: v_add_f32_e64191// AMD-OFF-NEXT: s_setpc_b64192 193// NV-OPT-FAST: fma.rn.f32194// NV-OPT-FAST-NEXT: st.param.b32195// NV-OPT-FASTSTD: fma.rn.f32196// NV-OPT-FASTSTD-NEXT: st.param.b32197// NV-OPT-ON: fma.rn.f32198// NV-OPT-ON-NEXT: st.param.b32199// NV-OPT-OFF: mul.rn.f32200// NV-OPT-OFF-NEXT: add.rn.f32201// NV-OPT-OFF-NEXT: st.param.b32202 203// AMD-OPT-FAST-IR: fmul contract float204// AMD-OPT-FAST-IR: fadd contract float205// AMD-OPT-ON-IR: @llvm.fmuladd.f32206// AMD-OPT-OFF-IR: fmul float207// AMD-OPT-OFF-IR: fadd float208 209// AMD-OPT-FAST: v_fmac_f32_e32210// AMD-OPT-FAST-NEXT: s_setpc_b64211// AMD-OPT-FASTSTD: v_fmac_f32_e32212// AMD-OPT-FASTSTD-NEXT: s_setpc_b64213// AMD-OPT-ON: v_fmac_f32_e32214// AMD-OPT-ON-NEXT: s_setpc_b64215// AMD-OPT-OFF: v_mul_f32_e32216// AMD-OPT-OFF-NEXT: v_add_f32_e32217// AMD-OPT-OFF-NEXT: s_setpc_b64218 219// Test multiply/add in the different statements, which can be emitted as220// FMA when fp-contract is fast but not on.221__host__ __device__ float func2(float a, float b, float c) {222 float t = b * c;223 return t + a;224}225// COMMON-LABEL: _Z5func2fff226// NV-OPT-FAST: fma.rn.f32227// NV-OPT-FAST-NEXT: st.param.b32228// NV-OPT-FASTSTD: fma.rn.f32229// NV-OPT-FASTSTD-NEXT: st.param.b32230// NV-OPT-ON: mul.rn.f32231// NV-OPT-ON: add.rn.f32232// NV-OPT-ON-NEXT: st.param.b32233// NV-OPT-OFF: mul.rn.f32234// NV-OPT-OFF: add.rn.f32235// NV-OPT-OFF-NEXT: st.param.b32236 237// AMD-OPT-FAST-IR: fmul contract float238// AMD-OPT-FAST-IR: fadd contract float239// AMD-OPT-ON-IR: fmul float240// AMD-OPT-ON-IR: fadd float241// AMD-OPT-OFF-IR: fmul float242// AMD-OPT-OFF-IR: fadd float243 244// AMD-OPT-FAST: v_fmac_f32_e32245// AMD-OPT-FAST-NEXT: s_setpc_b64246// AMD-OPT-FASTSTD: v_fmac_f32_e32247// AMD-OPT-FASTSTD-NEXT: s_setpc_b64248// AMD-OPT-ON: v_mul_f32_e32249// AMD-OPT-ON-NEXT: v_add_f32_e32250// AMD-OPT-ON-NEXT: s_setpc_b64251// AMD-OPT-OFF: v_mul_f32_e32252// AMD-OPT-OFF-NEXT: v_add_f32_e32253// AMD-OPT-OFF-NEXT: s_setpc_b64254 255// Test multiply/add in the different statements, which is forced256// to be compiled with fp contract on. fmul/fadd without contract257// flags are emitted in IR. In nvptx, they are emitted as FMA in258// fp-contract is fast but not on, as nvptx backend uses the same259// fp fuse option as front end, whereas fast fp fuse option in260// backend fuses fadd/fmul disregarding contract flag. In amdgcn261// they are not fused as amdgcn always use standard fp fusion262// option which respects contract flag.263 __host__ __device__ float func3(float a, float b, float c) {264#pragma clang fp contract(on)265 float t = b * c;266 return t + a;267}268// COMMON-LABEL: _Z5func3fff269// NV-OPT-FAST: fma.rn.f32270// NV-OPT-FAST-NEXT: st.param.b32271// NV-OPT-FASTSTD: mul.rn.f32272// NV-OPT-FASTSTD: add.rn.f32273// NV-OPT-FASTSTD-NEXT: st.param.b32274// NV-OPT-ON: mul.rn.f32275// NV-OPT-ON: add.rn.f32276// NV-OPT-ON-NEXT: st.param.b32277// NV-OPT-OFF: mul.rn.f32278// NV-OPT-OFF: add.rn.f32279// NV-OPT-OFF-NEXT: st.param.b32280 281// AMD-OPT-FAST-IR: fmul float282// AMD-OPT-FAST-IR: fadd float283// AMD-OPT-ON-IR: fmul float284// AMD-OPT-ON-IR: fadd float285// AMD-OPT-OFF-IR: fmul float286// AMD-OPT-OFF-IR: fadd float287 288// AMD-OPT-FAST: v_fmac_f32_e32289// AMD-OPT-FAST-NEXT: s_setpc_b64290// AMD-OPT-FASTSTD: v_mul_f32_e32291// AMD-OPT-FASTSTD-NEXT: v_add_f32_e32292// AMD-OPT-FASTSTD-NEXT: s_setpc_b64293// AMD-OPT-ON: v_mul_f32_e32294// AMD-OPT-ON-NEXT: v_add_f32_e32295// AMD-OPT-ON-NEXT: s_setpc_b64296// AMD-OPT-OFF: v_mul_f32_e32297// AMD-OPT-OFF-NEXT: v_add_f32_e32298// AMD-OPT-OFF-NEXT: s_setpc_b64299