Author: Michael Liao Date: 2020-09-16T14:10:05-04:00 New Revision: 4d4f0922837de3f1aa9862ae8a8d941b3b6e5f78
URL: https://github.com/llvm/llvm-project/commit/4d4f0922837de3f1aa9862ae8a8d941b3b6e5f78 DIFF: https://github.com/llvm/llvm-project/commit/4d4f0922837de3f1aa9862ae8a8d941b3b6e5f78.diff LOG: [clang][codegen] Skip adding default function attributes on intrinsics. - After loading builtin bitcode for linking, skip adding default function attributes on LLVM intrinsics as their attributes are well-defined and retrieved directly from internal definitions. Adding extra attributes on intrinsics results in inconsistent result when `-save-temps` is present. Also, that makes few optimizations conservative. Differential Revision: https://reviews.llvm.org/D87761 Added: clang/test/CodeGenCUDA/Inputs/device-lib-code.ll clang/test/CodeGenCUDA/dft-func-attr-skip-intrinsic.hip Modified: clang/lib/CodeGen/CodeGenAction.cpp Removed: ################################################################################ diff --git a/clang/lib/CodeGen/CodeGenAction.cpp b/clang/lib/CodeGen/CodeGenAction.cpp index 5a6ce0f5dbd5..eda4beff78b7 100644 --- a/clang/lib/CodeGen/CodeGenAction.cpp +++ b/clang/lib/CodeGen/CodeGenAction.cpp @@ -245,8 +245,13 @@ namespace clang { 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 index 000000000000..43ec911fb02c --- /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 index 000000000000..9e3e436200fc --- /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: } _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
