80 lines · plain
1// Check that when we link a bitcode module into a file using2// -mlink-builtin-bitcode, we apply the same attributes to the functions in that3// bitcode module as we apply to functions we generate.4//5// In particular, we check that ftz and unsafe-math are propagated into the6// bitcode library as appropriate.7 8// Build the bitcode library. This is not built in CUDA mode, otherwise it9// might have incompatible attributes. This mirrors how libdevice is built.10// RUN: %clang_cc1 -x c++ -fconvergent-functions -emit-llvm-bc -DLIB \11// RUN: %s -o %t.bc -triple nvptx-unknown-unknown12 13// RUN: %clang_cc1 -x cuda %s -emit-llvm -mlink-builtin-bitcode %t.bc -o - \14// RUN: -fcuda-is-device -triple nvptx-unknown-unknown \15// RUN: | FileCheck %s --check-prefix=CHECK --check-prefix=NOFTZ16 17// RUN: %clang_cc1 -x cuda %s -emit-llvm -mlink-builtin-bitcode %t.bc \18// RUN: -fdenormal-fp-math-f32=preserve-sign -o - \19// RUN: -fcuda-is-device -triple nvptx-unknown-unknown \20// RUN: | FileCheck %s --check-prefix=CHECK --check-prefix=FTZ21 22// RUN: %clang_cc1 -x cuda %s -emit-llvm -mlink-builtin-bitcode %t.bc \23// RUN: -fdenormal-fp-math-f32=preserve-sign -o - \24// RUN: -fcuda-is-device -funsafe-math-optimizations -triple nvptx-unknown-unknown \25// RUN: | FileCheck %s --check-prefix=CHECK26 27#ifndef LIB28#include "Inputs/cuda.h"29#endif30 31// Wrap everything in extern "C" so we don't have to worry about name mangling32// in the IR.33extern "C" {34#ifdef LIB35 36// This function is defined in the library and only declared in the main37// compilation.38void lib_fn() {}39 40#else41 42__device__ void lib_fn();43__global__ void kernel() { lib_fn(); }44 45#endif46}47 48// The kernel and lib function should have the same attributes.49// CHECK: define{{.*}} void @kernel() [[kattr:#[0-9]+]]50// CHECK: define internal void @lib_fn() [[fattr:#[0-9]+]]51 52// FIXME: These -NOT checks do not work as intended and do not check on the same53// line.54 55// Check the attribute list for kernel.56// CHECK: attributes [[kattr]] = {57 58// CHECK-SAME: convergent59// CHECK-SAME: norecurse60 61// FTZ-NOT: "denormal-fp-math"62// FTZ-SAME: "denormal-fp-math-f32"="preserve-sign,preserve-sign"63// NOFTZ-NOT: "denormal-fp-math-f32"64 65// CHECK-SAME: "no-trapping-math"="true"66 67// Check the attribute list for lib_fn.68// CHECK: attributes [[fattr]] = {69 70// CHECK-SAME: convergent71// CHECK-NOT: norecurse72 73// FTZ-NOT: "denormal-fp-math"74// NOFTZ-NOT: "denormal-fp-math"75 76// FTZ-SAME: "denormal-fp-math-f32"="preserve-sign,preserve-sign"77// NOFTZ-NOT: "denormal-fp-math-f32"78 79// CHECK-SAME: "no-trapping-math"="true"80