[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:

Author: jhuber6

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<

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);
     }

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

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