[OpenMP] Customize CUDA-based tool chain selection

[gtbercea]

Summary

This patch provides a generic way of selecting CUDA based tool chains as host-device pairs.

[Hahnfeld]

lib/Frontend/CompilerInvocation.cpp

2181–2190 ↗

(On Diff #87454)

This is unrelated to CUDA, please split that into its own patch.

[gtbercea]

Split patch.

[Hahnfeld]

Is there a way to actually test the changed code? The current test changes seem unrelated...

include/clang/Basic/DiagnosticDriverKinds.td

165–166

This diagnostic seems unrelated to the code.

test/OpenMP/target_messages.cpp

7–8

This test seems unrelated to the changes to the code.

[gtbercea]

Remove tests which belong into a different patch.

[Hahnfeld]

LGTM with one small note

lib/Driver/Driver.cpp

564

The code above uses HostTriple.str(), maybe better align to this?

[ABataev]

Tests?

lib/Driver/Driver.cpp

556–571

Seems to me you're not using this initial value, so it is better to leave this variable uninitialized

[gtbercea]

Rebase on latest master.

[gtbercea]

lib/Driver/Driver.cpp

564

HostTriple is equivalent with HostTC->getTriple(). HostTC->getTriple() is only used once so no need for a local variable unless you want me to refactor the CUDA code above and take out HostTriple from within the if statement and then re0use HostTriple that way. I'd rather not make changes to the CUDA toolchain above.

[gtbercea]

In D29658#732520, @ABataev wrote:

Tests?

Can't write any meaningful tests. This will be tested implicitely by all future patches that perform offloading using OpenMP.

[Hahnfeld]

lib/Driver/Driver.cpp

564

Sorry, I meant the mismatch between .str() above and .normalize() here

[gtbercea]

Use str()

include/clang/Basic/DiagnosticDriverKinds.td

 def err_drv_optimization_remark_pattern : Error<
   "%0 in '%1'">;
 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_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<
   "The target '%0' is not a supported OpenMP host target.">;
+def err_drv_omp_device_target_not_supported : Error<
+  "The target '%0' is not a supported OpenMP device target.">;
 def err_drv_expecting_fopenmp_with_fopenmp_targets : Error<
   "The option -fopenmp-targets must be used in conjunction with a -fopenmp option compatible with offloading, please use -fopenmp=libomp or -fopenmp=libiomp5.">;
 def warn_drv_omp_offload_target_duplicate : Warning<
   "The OpenMP offloading target '%0' is similar to target '%1' already specified - will be ignored.">, 
   InGroup<OpenMPTarget>;
 def err_drv_bitcode_unsupported_on_toolchain : Error<
   "-fembed-bitcode is not supported on versions of iOS prior to 6.0">;
 

lib/Driver/Driver.cpp

if (OpenMPTargets->getNumValues()) {
           // Store the current triple so that we can check for duplicates in the
           // following iterations.
           FoundNormalizedTriples[NormalizedName] = Val;
 
           // If the specified target is invalid, emit a diagnostic.
           if (TT.getArch() == llvm::Triple::UnknownArch)
             Diag(clang::diag::err_drv_invalid_omp_target) << Val;
           else {
-            const ToolChain &TC = getToolChain(C.getInputArgs(), TT);
-            C.addOffloadDeviceToolChain(&TC, Action::OFK_OpenMP);
+            const ToolChain *TC = nullptr;
+            // CUDA toolchains have to be selected differently. They pair host
+            // and device in their implementation.
+            if (TT.isNVPTX()) {
+              const ToolChain *HostTC =
+                  C.getSingleOffloadToolChain<Action::OFK_Host>();
+              assert(HostTC && "Host toolchain should be always defined.");
+              auto &CudaTC =
+                  ToolChains[TT.str() + "/" + HostTC->getTriple().normalize()];
+              if (!CudaTC)
+                CudaTC = llvm::make_unique<toolchains::CudaToolChain>(
+                    *this, TT, *HostTC, C.getInputArgs());
+              TC = CudaTC.get();
+            } else
+              TC = &getToolChain(C.getInputArgs(), TT);
+            C.addOffloadDeviceToolChain(TC, Action::OFK_OpenMP);
           }
         }
       } else
         Diag(clang::diag::err_drv_expecting_fopenmp_with_fopenmp_targets);
     } else
       Diag(clang::diag::warn_drv_empty_joined_argument)
           << OpenMPTargets->getAsString(C.getInputArgs());
   }

test/OpenMP/target_messages.cpp

 // RUN: %clang_cc1 -verify -fopenmp -std=c++11 -o - %s
 // RUN: not %clang_cc1 -fopenmp -std=c++11 -fopenmp-targets=aaa-bbb-ccc-ddd -o - %s 2>&1 | FileCheck %s
 // CHECK: error: OpenMP target is invalid: 'aaa-bbb-ccc-ddd'
 // RUN: not %clang_cc1 -fopenmp -std=c++11 -triple nvptx64-nvidia-cuda -o - %s 2>&1 | FileCheck --check-prefix CHECK-UNSUPPORTED-HOST-TARGET %s
 // RUN: not %clang_cc1 -fopenmp -std=c++11 -triple nvptx-nvidia-cuda -o - %s 2>&1 | FileCheck --check-prefix CHECK-UNSUPPORTED-HOST-TARGET %s
 // CHECK-UNSUPPORTED-HOST-TARGET: error: The target '{{nvptx64-nvidia-cuda|nvptx-nvidia-cuda}}' is not a supported OpenMP host target.
+// RUN: not %clang_cc1 -fopenmp -std=c++11 -fopenmp-targets=hexagon-linux-gnu -o - %s 2>&1 | FileCheck --check-prefix CHECK-UNSUPPORTED-DEVICE-TARGET %s
+// CHECK-UNSUPPORTED-DEVICE-TARGET: The target 'hexagon-linux-gnu' is not a supported OpenMP device target.
 
 void foo() {
 }
 
 #pragma omp target // expected-error {{unexpected OpenMP directive '#pragma omp target'}}
 
 int main(int argc, char **argv) {
   #pragma omp target { // expected-warning {{extra tokens at the end of '#pragma omp target' are ignored}}