diff --git a/clang/lib/CodeGen/CodeGenAction.cpp b/clang/lib/CodeGen/CodeGenAction.cpp --- a/clang/lib/CodeGen/CodeGenAction.cpp +++ b/clang/lib/CodeGen/CodeGenAction.cpp @@ -245,8 +245,13 @@ bool LinkInModules() { for (auto &LM : LinkModules) { if (LM.PropagateAttrs) - for (Function &F : *LM.Module) + for (Function &F : *LM.Module) { + // Skip intrinsics. Keep consistent with how intrinsics are created + // in LLVM IR. + if (F.isIntrinsic()) + continue; Gen->CGM().addDefaultFunctionDefinitionAttributes(F); + } CurLinkModule = LM.Module.get(); diff --git a/clang/test/CodeGenCUDA/Inputs/device-lib-code.ll b/clang/test/CodeGenCUDA/Inputs/device-lib-code.ll new file mode 100644 --- /dev/null +++ b/clang/test/CodeGenCUDA/Inputs/device-lib-code.ll @@ -0,0 +1,5 @@ +define linkonce_odr protected float @__ocml_fma_f32(float %0, float %1, float %2) local_unnamed_addr { + %4 = tail call float @llvm.fma.f32(float %0, float %1, float %2) + ret float %4 +} +declare float @llvm.fma.f32(float, float, float) diff --git a/clang/test/CodeGenCUDA/dft-func-attr-skip-intrinsic.hip b/clang/test/CodeGenCUDA/dft-func-attr-skip-intrinsic.hip new file mode 100644 --- /dev/null +++ b/clang/test/CodeGenCUDA/dft-func-attr-skip-intrinsic.hip @@ -0,0 +1,18 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -x ir -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm-bc -disable-llvm-passes -o %t.bc %S/Inputs/device-lib-code.ll +// RUN: %clang_cc1 -x hip -fcuda-is-device -triple amdgcn-amd-amdhsa -mlink-builtin-bitcode %t.bc -emit-llvm -disable-llvm-passes -o - %s | FileCheck %s + +#include "Inputs/cuda.h" + +extern "C" __device__ float __ocml_fma_f32(float x, float y, float z); + +__device__ float foo(float x) { + return __ocml_fma_f32(x, x, x); +} + +// CHECK: {{^}}define{{.*}} @__ocml_fma_f32{{.*}} [[ATTR1:#[0-9]+]] +// CHECK: {{^}}declare{{.*}} @llvm.fma.f32{{.*}} [[ATTR2:#[0-9]+]] +// CHECK: attributes [[ATTR1]] = { convergent +// CHECK: attributes [[ATTR2]] = { +// CHECK-NOT: convergent +// CHECK: }