| Artem Belevich | 31c3bad | 2016-05-19 18:44:45 +0000 | [diff] [blame] | 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 | |
| Adam Nemet | 049a31d | 2017-03-29 21:54:24 +0000 | [diff] [blame] | 13 | // Explicit -ffp-contract=on -- fusing by front-end. |
| Artem Belevich | 31c3bad | 2016-05-19 18:44:45 +0000 | [diff] [blame] | 14 | // RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \ |
| 15 | // RUN: -ffp-contract=on -disable-llvm-passes -o - %s \ |
| Adam Nemet | 049a31d | 2017-03-29 21:54:24 +0000 | [diff] [blame] | 16 | // RUN: | FileCheck -check-prefix ENABLED %s |
| Artem Belevich | 31c3bad | 2016-05-19 18:44:45 +0000 | [diff] [blame] | 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 | |
| 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 |