1 // Check that when we link a bitcode module into a file using 2 // -mlink-cuda-bitcode, we apply the same attributes to the functions in that 3 // bitcode module as we apply to functions we generate. 4 // 5 // In particular, we check that ftz and unsafe-math are propagated into the 6 // bitcode library as appropriate. 7 // 8 // In addition, we set -ftrapping-math on the bitcode library, but then set 9 // -fno-trapping-math on the main compilations, and ensure that the latter flag 10 // overrides the flag on the bitcode library. 11 12 // Build the bitcode library. This is not built in CUDA mode, otherwise it 13 // might have incompatible attributes. This mirrors how libdevice is built. 14 // RUN: %clang_cc1 -x c++ -emit-llvm-bc -ftrapping-math -DLIB \ 15 // RUN: %s -o %t.bc -triple nvptx-unknown-unknown 16 17 // RUN: %clang_cc1 -x cuda %s -emit-llvm -mlink-cuda-bitcode %t.bc -o - \ 18 // RUN: -fno-trapping-math -fcuda-is-device -triple nvptx-unknown-unknown \ 19 // RUN: | FileCheck %s --check-prefix=CHECK --check-prefix=NOFTZ --check-prefix=NOFAST 20 21 // RUN: %clang_cc1 -x cuda %s -emit-llvm -mlink-cuda-bitcode %t.bc \ 22 // RUN: -fno-trapping-math -fcuda-flush-denormals-to-zero -o - \ 23 // RUN: -fcuda-is-device -triple nvptx-unknown-unknown \ 24 // RUN: | FileCheck %s --check-prefix=CHECK --check-prefix=FTZ \ 25 // RUN: --check-prefix=NOFAST 26 27 // RUN: %clang_cc1 -x cuda %s -emit-llvm -mlink-cuda-bitcode %t.bc \ 28 // RUN: -fno-trapping-math -fcuda-flush-denormals-to-zero -o - \ 29 // RUN: -fcuda-is-device -menable-unsafe-fp-math -triple nvptx-unknown-unknown \ 30 // RUN: | FileCheck %s --check-prefix=CHECK --check-prefix=FAST 31 32 // Wrap everything in extern "C" so we don't ahve to worry about name mangling 33 // in the IR. 34 extern "C" { 35 #ifdef LIB 36 37 // This function is defined in the library and only declared in the main 38 // compilation. 39 void lib_fn() {} 40 41 #else 42 43 #include "Inputs/cuda.h" 44 __device__ void lib_fn(); 45 __global__ void kernel() { lib_fn(); } 46 47 #endif 48 } 49 50 // The kernel and lib function should have the same attributes. 51 // CHECK: define void @kernel() [[attr:#[0-9]+]] 52 // CHECK: define internal void @lib_fn() [[attr]] 53 54 // Check the attribute list. 55 // CHECK: attributes [[attr]] = { 56 // CHECK: "no-trapping-math"="true" 57 58 // FTZ-SAME: "nvptx-f32ftz"="true" 59 // NOFTZ-NOT: "nvptx-f32ftz"="true" 60 61 // FAST-SAME: "unsafe-fp-math"="true" 62 // NOFAST-NOT: "unsafe-fp-math"="true" 63