1 // REQUIRES: x86-registered-target
2 // REQUIRES: nvptx-registered-target
3
4 // By default we should fuse multiply/add into fma instruction.
5 // RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \
6 // RUN: -disable-llvm-passes -o - %s | FileCheck -check-prefix ENABLED %s
7
8 // Explicit -ffp-contract=fast
9 // RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \
10 // RUN: -ffp-contract=fast -disable-llvm-passes -o - %s \
11 // RUN: | FileCheck -check-prefix ENABLED %s
12
13 // Explicit -ffp-contract=on -- fusing by front-end (disabled).
14 // RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \
15 // RUN: -ffp-contract=on -disable-llvm-passes -o - %s \
16 // RUN: | FileCheck -check-prefix DISABLED %s
17
18 // Explicit -ffp-contract=off should disable instruction fusing.
19 // RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \
20 // RUN: -ffp-contract=off -disable-llvm-passes -o - %s \
21 // RUN: | FileCheck -check-prefix DISABLED %s
22
23
24 #include "Inputs/cuda.h"
25
func(float a,float b,float c)26 __host__ __device__ float func(float a, float b, float c) { return a + b * c; }
27 // ENABLED: fma.rn.f32
28 // ENABLED-NEXT: st.param.f32
29
30 // DISABLED: mul.rn.f32
31 // DISABLED-NEXT: add.rn.f32
32 // DISABLED-NEXT: st.param.f32
33