From 5ad52587ec182f03636649e2cb66a0a4d9ffeab2 Mon Sep 17 00:00:00 2001
From: Alexey Bataev <a.bataev@hotmail.com>
Date: Thu, 12 Dec 2019 15:33:18 -0500
Subject: [PATCH] [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.
---
 clang/lib/CodeGen/CGOpenMPRuntime.cpp             |  2 +-
 clang/test/OpenMP/declare_variant_mixed_codegen.c | 49 +++++++++++++++++++++++
 2 files changed, 50 insertions(+), 1 deletion(-)
 create mode 100644 clang/test/OpenMP/declare_variant_mixed_codegen.c

diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index e3e9bd4..1aae18b 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 0000000..6345709
--- /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
-- 
2.7.4