[clang] 9d2378b - [OpenMP] Add Error Handling for Conflicting Pointer Sizes for Target Offload
via cfe-commits
cfe-commits at lists.llvm.org
Wed Sep 30 10:59:51 PDT 2020
Author: Joseph Huber
Date: 2020-09-30T13:58:24-04:00
New Revision: 9d2378b59150f6f1cb5c9cf42ea06b0bb57029a1
URL: https://github.com/llvm/llvm-project/commit/9d2378b59150f6f1cb5c9cf42ea06b0bb57029a1
DIFF: https://github.com/llvm/llvm-project/commit/9d2378b59150f6f1cb5c9cf42ea06b0bb57029a1.diff
LOG: [OpenMP] Add Error Handling for Conflicting Pointer Sizes for Target Offload
Summary:
This patch adds an error to Clang that detects if OpenMP offloading is used
between two architectures with incompatible pointer sizes. This ensures that
the data mapping can be done correctly and solves an issue in code generation
generating the wrong size pointer.
Reviewer: jdoerfert
Subscribers:
Tags: #OpenMP #Clang
Differential Revision:
Added:
clang/test/OpenMP/target_incompatible_architecture_messages.cpp
Modified:
clang/include/clang/Basic/DiagnosticDriverKinds.td
clang/lib/Frontend/CompilerInvocation.cpp
clang/test/OpenMP/nvptx_target_parallel_reduction_codegen_tbaa_PR46146.cpp
Removed:
################################################################################
diff --git a/clang/include/clang/Basic/DiagnosticDriverKinds.td b/clang/include/clang/Basic/DiagnosticDriverKinds.td
index 3bf1bb19b7ae..29bc19e5a84e 100644
--- a/clang/include/clang/Basic/DiagnosticDriverKinds.td
+++ b/clang/include/clang/Basic/DiagnosticDriverKinds.td
@@ -253,6 +253,7 @@ def err_drv_optimization_remark_format : Error<
"unknown remark serializer format: '%0'">;
def err_drv_no_neon_modifier : Error<"[no]neon is not accepted as modifier, please use [no]simd instead">;
def err_drv_invalid_omp_target : Error<"OpenMP target is invalid: '%0'">;
+def err_drv_incompatible_omp_arch : Error<"OpenMP target architecture '%0' pointer size is incompatible with host '%1'">;
def err_drv_omp_host_ir_file_not_found : Error<
"The provided host compiler IR file '%0' is required to generate code for OpenMP target regions but cannot be found.">;
def err_drv_omp_host_target_not_supported : Error<
diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp
index b402f53cc765..bbdf0e3be7ae 100644
--- a/clang/lib/Frontend/CompilerInvocation.cpp
+++ b/clang/lib/Frontend/CompilerInvocation.cpp
@@ -3206,6 +3206,14 @@ static void ParseLangArgs(LangOptions &Opts, ArgList &Args, InputKind IK,
TT.getArch() == llvm::Triple::x86 ||
TT.getArch() == llvm::Triple::x86_64))
Diags.Report(diag::err_drv_invalid_omp_target) << A->getValue(i);
+ else if ((T.isArch64Bit() && TT.isArch32Bit()) ||
+ (T.isArch64Bit() && TT.isArch16Bit()) ||
+ (T.isArch32Bit() && TT.isArch64Bit()) ||
+ (T.isArch32Bit() && TT.isArch16Bit()) ||
+ (T.isArch16Bit() && TT.isArch32Bit()) ||
+ (T.isArch16Bit() && TT.isArch64Bit()))
+ Diags.Report(diag::err_drv_incompatible_omp_arch)
+ << A->getValue(i) << T.str();
else
Opts.OMPTargetTriples.push_back(TT);
}
diff --git a/clang/test/OpenMP/nvptx_target_parallel_reduction_codegen_tbaa_PR46146.cpp b/clang/test/OpenMP/nvptx_target_parallel_reduction_codegen_tbaa_PR46146.cpp
index aefe00f1cadf..031c7b6c778e 100644
--- a/clang/test/OpenMP/nvptx_target_parallel_reduction_codegen_tbaa_PR46146.cpp
+++ b/clang/test/OpenMP/nvptx_target_parallel_reduction_codegen_tbaa_PR46146.cpp
@@ -1,8 +1,8 @@
// RUN: %clang_cc1 -x c++ -O1 -disable-llvm-optzns -verify -fopenmp -internal-isystem %S/../Headers/Inputs/include -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
// RUN: %clang_cc1 -x c++ -O1 -disable-llvm-optzns -verify -fopenmp -internal-isystem %S/../Headers/Inputs/include -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s
// RUN: %clang_cc1 -x c++ -O1 -disable-llvm-optzns -verify -fopenmp -internal-isystem %S/../Headers/Inputs/include -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
-// RUN: %clang_cc1 -x c++ -O1 -disable-llvm-optzns -verify -fopenmp -internal-isystem %S/../Headers/Inputs/include -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -triple nvptx-unknown-unknown -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s
-// RUN: %clang_cc1 -x c++ -O1 -disable-llvm-optzns -verify -fopenmp -internal-isystem %S/../Headers/Inputs/include -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -fexceptions -fcxx-exceptions -aux-triple powerpc64le-unknown-unknown -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s
+// RUN: %clang_cc1 -x c++ -O1 -disable-llvm-optzns -verify -fopenmp -internal-isystem %S/../Headers/Inputs/include -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s
+// RUN: %clang_cc1 -x c++ -O1 -disable-llvm-optzns -verify -fopenmp -internal-isystem %S/../Headers/Inputs/include -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -fexceptions -fcxx-exceptions -aux-triple powerpc64le-unknown-unknown -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s
// expected-no-diagnostics
#ifndef HEADER
#define HEADER
diff --git a/clang/test/OpenMP/target_incompatible_architecture_messages.cpp b/clang/test/OpenMP/target_incompatible_architecture_messages.cpp
new file mode 100644
index 000000000000..f0f9d236d764
--- /dev/null
+++ b/clang/test/OpenMP/target_incompatible_architecture_messages.cpp
@@ -0,0 +1,14 @@
+// RUN: not %clang_cc1 -x c++ -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -o - %s 2>&1 | FileCheck %s
+// RUN: not %clang_cc1 -x c++ -fopenmp -triple i386-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -o - %s 2>&1 | FileCheck %s
+// RUN: not %clang_cc1 -x c++ -fopenmp -triple x86_64-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -o - %s 2>&1 | FileCheck %s
+// RUN: not %clang_cc1 -x c++ -fopenmp -triple x86_64-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -o - %s 2>&1 | FileCheck %s
+// CHECK: error: OpenMP target architecture '{{.+}}' pointer size is incompatible with host '{{.+}}'
+#ifndef HEADER
+#define HEADER
+
+void test() {
+#pragma omp target
+ {}
+}
+
+#endif
More information about the cfe-commits
mailing list