Mercurial > hg > CbC > CbC_llvm

view clang/test/CodeGenCUDA/fp-contract.cu @ 150:1d019706d866

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
LLVM10
author anatofuz
date Thu, 13 Feb 2020 15:10:13 +0900
parents
children 2e18cbf3894f
line wrap: on
line source

// REQUIRES: x86-registered-target
// REQUIRES: nvptx-registered-target

// By default we should fuse multiply/add into fma instruction.
// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \
// RUN:   -disable-llvm-passes -o - %s | FileCheck -check-prefix ENABLED %s

// Explicit -ffp-contract=fast
// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \
// RUN:   -ffp-contract=fast -disable-llvm-passes -o - %s \
// RUN:   | FileCheck -check-prefix ENABLED %s

// Explicit -ffp-contract=on -- fusing by front-end.
// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \
// RUN:   -ffp-contract=on -disable-llvm-passes -o - %s \
// RUN:   | FileCheck -check-prefix ENABLED %s

// Explicit -ffp-contract=off should disable instruction fusing.
// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \
// RUN:   -ffp-contract=off -disable-llvm-passes -o - %s \
// RUN:   | FileCheck -check-prefix DISABLED %s


#include "Inputs/cuda.h"

__host__ __device__ float func(float a, float b, float c) { return a + b * c; }
// ENABLED:       fma.rn.f32
// ENABLED-NEXT:  st.param.f32

// DISABLED:      mul.rn.f32
// DISABLED-NEXT: add.rn.f32
// DISABLED-NEXT: st.param.f32