[PATCH] D87761: [clang][codegen] Skip adding default function attributes on intrinsics.

[Michael Liao via Phabricator via cfe-commits]

hliao created this revision.
hliao added reviewers: rjmccall, tra, yaxunl.
Herald added a project: clang.
Herald added a subscriber: cfe-commits.
hliao requested review of this revision.

- 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.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D87761

Files:
  clang/lib/CodeGen/CodeGenAction.cpp
  clang/test/CodeGenCUDA/Inputs/device-lib-code.ll
  clang/test/CodeGenCUDA/dft-func-attr-skip-intrinsic.hip


Index: clang/test/CodeGenCUDA/dft-func-attr-skip-intrinsic.hip
===================================================================
--- /dev/null
+++ 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: }
Index: clang/test/CodeGenCUDA/Inputs/device-lib-code.ll
===================================================================
--- /dev/null
+++ 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)
Index: clang/lib/CodeGen/CodeGenAction.cpp
===================================================================
--- clang/lib/CodeGen/CodeGenAction.cpp
+++ 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();
 



Index: clang/test/CodeGenCUDA/dft-func-attr-skip-intrinsic.hip
===================================================================
--- /dev/null
+++ 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: }
Index: clang/test/CodeGenCUDA/Inputs/device-lib-code.ll
===================================================================
--- /dev/null
+++ 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)
Index: clang/lib/CodeGen/CodeGenAction.cpp
===================================================================
--- clang/lib/CodeGen/CodeGenAction.cpp
+++ 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();
 

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits