[HIP] Warn no --offload-arch option

[yaxunl]

Summary

This patch let clang emit a warning when no --offload-arch option is specified, which
usually indicates that users forget to specify this option.

[tra]

While I agree that the default GPU choice is not likely to be correct, or usable, for everyone, but the warning seems to be a half-measure.
If the default is not usable, then it should not be the default. If it's usable, then we don't need a warning.

Having a warning would make sense if it were a part of the plan to transition towards making GPU arch a mandatory option. Is that the case here?
Just a warning is not very useful, IMO. The users would need to specify the GPU in order to silence it, so why not just require it.

[yaxunl]

In D92363#2426401, @tra wrote:

While I agree that the default GPU choice is not likely to be correct, or usable, for everyone, but the warning seems to be a half-measure.
If the default is not usable, then it should not be the default. If it's usable, then we don't need a warning.

Having a warning would make sense if it were a part of the plan to transition towards making GPU arch a mandatory option. Is that the case here?
Just a warning is not very useful, IMO. The users would need to specify the GPU in order to silence it, so why not just require it.

There are still valid use cases for not providing GPU arch. For example, users would like to test syntax of a HIP program for which GPU arch does not matter, or users would like to get PCH file where GPU arch does not matter. Another example is cmake may test whether a compiler can compile HIP program without options. Make it a warning allows such uses cases.

[tra]

In D92363#2431812, @yaxunl wrote:

There are still valid use cases for not providing GPU arch. For example, users would like to test syntax of a HIP program for which GPU arch does not matter, or users would like to get PCH file where GPU arch does not matter. Another example is cmake may test whether a compiler can compile HIP program without options. Make it a warning allows such uses cases.

Issuing no warning allows such use cases even better. :-)

Also, how do I silence that warning, if I'm compiling with -Werror?
If we provide a knob to silence the warning, how would that be better than asking the user to specify a GPU arch?
If we do ask users to specify GPU arch in order to silence the warning => we effectively make the GPU arch a mandatory option.

I believe that no warning or a mandatory GPU arch would be a more principled way. With the warning we effectively flag the default value as 'it's not a very useful default'. If it's not useful by default, then it should not be the default. If it is useful, there should be no warning.

In this case, I think making GPU arch mandatory would have a stronger case, as the default we use right now is indeed quite often wrong. However, we probably have existing users who have valid use cases that work without the explicitly specified GPU arch, so requiring the GPU arch would break them. That constrains us to the no-warning scenario in practice.

Introducing a warning would break existing users, too -- compiling with -Werror is not that uncommon. If we're willing to break users, we may as well break them by requiring a GPU arch.
I just don't see a compelling argument where adding a warning would provide a net benefit over no-warning or a mandatory-GPU-arch scenarios.

[yaxunl]

@tra I got some issue with this patch. There are cases that an expression using a host variable is compile-time constant, e.g.

int x;
__device__ void fun() {
  sizeof(x);
}

Do we want to allow that? Thanks.

[yaxunl]

In D92363#2607072, @yaxunl wrote:

@tra I got some issue with this patch. There are cases that an expression using a host variable is compile-time constant, e.g.

int x;
__device__ void fun() {
  sizeof(x);
}

Do we want to allow that? Thanks.

sorry, wrong patch.

clang/include/clang/Basic/DiagnosticDriverKinds.td

def warn_drv_unknown_cuda_version: Warning<
   InGroup<CudaUnknownVersion>;
 def err_drv_cuda_host_arch : Error<"unsupported architecture '%0' for host compilation.">;
 def err_drv_mix_cuda_hip : Error<"Mixed Cuda and HIP compilation is not supported.">;
 def err_drv_bad_target_id : Error<"Invalid target ID: %0 (A target ID is a processor name "
   "followed by an optional list of predefined features post-fixed by a plus or minus sign deliminated "
   "by colon, e.g. 'gfx908:sram-ecc+:xnack-')">;
 def err_drv_bad_offload_arch_combo : Error<"Invalid offload arch combinations: %0 and %1 (For a specific "
   "processor, a feature should either exist in all offload archs, or not exist in any offload archs)">;
+def warn_drv_no_offload_arch : Warning<"no offload arch is specified. Assuming %0. Use --offload-arch=<arch> to"
+  " specify offload arch">, InGroup<HIPOptionMising>;
 def err_drv_invalid_thread_model_for_target : Error<
   "invalid thread model '%0' in '%1' for this target">;
 def err_drv_invalid_linker_name : Error<
   "invalid linker name in argument '%0'">;
 def err_drv_invalid_pgo_instrumentor : Error<
   "invalid PGO instrumentor in argument '%0'">;
 def err_drv_invalid_rtlib_name : Error<
   "invalid runtime library name in argument '%0'">;

clang/include/clang/Basic/DiagnosticGroups.td

 
 // Warning about unknown CUDA SDK version.
 def CudaUnknownVersion: DiagGroup<"unknown-cuda-version">;
 
 // A warning group for warnings about features supported by HIP but
 // ignored by CUDA.
 def HIPOnly : DiagGroup<"hip-only">;
 
+// A warning group for warnings about missing HIP options.
+def HIPOptionMising : DiagGroup<"hip-option-missing">;
+
 // Warnings which cause linking of the runtime libraries like
 // libc and the CRT to be skipped.
 def AVRRtlibLinkingQuirks : DiagGroup<"avr-rtlib-linking-quirks">;
 
 // A warning group for things that will change semantics in the future.
 def FutureCompat : DiagGroup<"future-compat">;
 
 def InvalidOrNonExistentDirectory : DiagGroup<"invalid-or-nonexistent-directory">;

clang/lib/Driver/Driver.cpp

bool initialize() override {
 
       // Collect list of GPUs remaining in the set.
       for (auto Arch : GpuArchs)
         GpuArchList.push_back(Arch.data());
 
       // Default to sm_20 which is the lowest common denominator for
       // supported GPUs.  sm_20 code should work correctly, if
       // suboptimally, on all newer GPUs.
-      if (GpuArchList.empty())
+      if (GpuArchList.empty()) {
         GpuArchList.push_back(DefaultCudaArch);
+        if (AssociatedOffloadKind == Action::OFK_HIP)
+          C.getDriver().Diag(clang::diag::warn_drv_no_offload_arch)
+              << CudaArchToString(DefaultCudaArch);
+      }
 
       return Error;
     }
   };
 
   /// \brief CUDA action builder. It injects device code in the host backend
   /// action.
   class CudaActionBuilder final : public CudaActionBuilderBase {

clang/test/Driver/hip-no-offload-arch.hip

+// REQUIRES: clang-driver, x86-registered-target, amdgpu-registered-target, nvptx-registered-target
+
+// RUN: %clang -### -target x86_64-linux-gnu -nogpuinc -nogpulib \
+// RUN:   %s 2>&1 | FileCheck -check-prefix=WARN %s
+// WARN: warning: no offload arch is specified. Assuming gfx803. Use --offload-arch=<arch> to specify offload arch [-Whip-option-missing]
+
+// RUN: %clang -### -target x86_64-linux-gnu -nogpuinc -nogpulib \
+// RUN:   -x cuda %s 2>&1 | FileCheck -check-prefix=NOWARN %s
+// NOWARN-NOT: warning