[clang] 5ad5258 - [OPENMP50]Fix possible conflict when emitting an alias for the functions

[Alexey Bataev via cfe-commits]

Author: Alexey Bataev
Date: 2019-12-12T15:48:33-05:00
New Revision: 5ad52587ec182f03636649e2cb66a0a4d9ffeab2

URL: https://github.com/llvm/llvm-project/commit/5ad52587ec182f03636649e2cb66a0a4d9ffeab2
DIFF: https://github.com/llvm/llvm-project/commit/5ad52587ec182f03636649e2cb66a0a4d9ffeab2.diff

LOG: [OPENMP50]Fix possible conflict when emitting an alias for the functions
in declare variant.

If the types of the fnction are not equal, but match, at the codegen
thei may have different types. This may lead to compiler crash.

Added: 
    clang/test/OpenMP/declare_variant_mixed_codegen.c

Modified: 
    clang/lib/CodeGen/CGOpenMPRuntime.cpp

Removed: 
    


################################################################################
diff  --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index e3e9bd4284c3..1aae18b4504c 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -1280,7 +1280,7 @@ bool CGOpenMPRuntime::tryEmitDeclareVariant(const GlobalDecl &NewGD,
   llvm::GlobalValue *Addr = CGM.GetGlobalValue(NewMangledName);
   if (Addr && !Addr->isDeclaration()) {
     const auto *D = cast<FunctionDecl>(OldGD.getDecl());
-    const CGFunctionInfo &FI = CGM.getTypes().arrangeGlobalDeclaration(OldGD);
+    const CGFunctionInfo &FI = CGM.getTypes().arrangeGlobalDeclaration(NewGD);
     llvm::Type *DeclTy = CGM.getTypes().GetFunctionType(FI);
 
     // Create a reference to the named value.  This ensures that it is emitted

diff  --git a/clang/test/OpenMP/declare_variant_mixed_codegen.c b/clang/test/OpenMP/declare_variant_mixed_codegen.c
new file mode 100644
index 000000000000..63457095b934
--- /dev/null
+++ b/clang/test/OpenMP/declare_variant_mixed_codegen.c
@@ -0,0 +1,49 @@
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple x86_64-unknown-linux -emit-llvm %s -o - | FileCheck %s --check-prefix HOST
+// RUN: %clang_cc1 -fopenmp -x c -triple x86_64-unknown-linux -emit-pch -o %t -fopenmp-version=50 %s
+// RUN: %clang_cc1 -fopenmp -x c -triple x86_64-unknown-linux -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=50 | FileCheck %s --check-prefix HOST
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -fopenmp-version=50
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -fopenmp-version=50 | FileCheck %s --check-prefix GPU
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t -fopenmp-version=50
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -o - -fopenmp-version=50 | FileCheck %s --check-prefix GPU
+// expected-no-diagnostics
+
+// HOST: @base = alias i32 (double), i32 (double)* @hst
+#ifndef HEADER
+#define HEADER
+
+int dev(double i) { return 0; }
+
+int hst(double i) { return 1; }
+
+#pragma omp declare variant(hst) match(device = {kind(host)})
+#pragma omp declare variant(dev) match(device = {kind(gpu)})
+int base();
+
+// HOST-LABEL: define void @foo()
+// HOST: call i32 (double, ...) bitcast (i32 (double)* @base to i32 (double, ...)*)(double -1.000000e+00)
+// HOST: call i32 @hst(double -2.000000e+00)
+// HOST: call void [[OFFL:@.+_foo_l29]]()
+void foo() {
+  base(-1);
+  hst(-2);
+#pragma omp target
+  {
+    base(-3);
+    dev(-4);
+  }
+}
+
+// HOST: define {{.*}}void [[OFFL]]()
+// HOST: call i32 (double, ...) bitcast (i32 (double)* @base to i32 (double, ...)*)(double -3.000000e+00)
+// HOST: call i32 @dev(double -4.000000e+00)
+
+// GPU: define {{.*}}void @__omp_offloading_{{.+}}_foo_l29()
+// GPU: call i32 @base(double -3.000000e+00)
+// GPU: call i32 @dev(double -4.000000e+00)
+
+// GPU: define {{.*}}i32 @base(double
+// GPU: ret i32 0
+// GPU: define {{.*}}i32 @dev(double
+// GPU: ret i32 0
+
+#endif // HEADER