Renable -fopenacc and emit warning for amdgpu

Author: ronlieb

clang/include/clang/Basic/DiagnosticDriverKinds.td

@@ -146,6 +146,9 @@ def warn_drv_unsupported_openmp_library : Warning<
   "the library '%0=%1' is not supported, OpenMP will not be enabled">,
   InGroup<OptionIgnored>;
 
+def warn_openacc_experimental : Warning<
+"OpenACC is NOT supported for AMDGPU">;
+
 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<

clang/include/clang/Driver/Options.td

@@ -3737,8 +3737,10 @@ def fno_openmp_target_debug : Flag<["-"], "fno-openmp-target-debug">;
 // FlangOption + FC1 + ClangOption + CC1Option
 //===----------------------------------------------------------------------===//
 // Unsupported on AMD downstream
+let Visibility = [FC1Option, FlangOption, CC1Option, ClangOption] in {
 def fopenacc : Flag<["-"], "fopenacc">, Group<f_Group>, Visibility<[]>,
   HelpText<"Enable OpenACC">;
+}
 
 //===----------------------------------------------------------------------===//
 // Optimisation remark options

clang/lib/Driver/ToolChains/Clang.cpp

@@ -4028,6 +4028,7 @@ static void RenderOpenACCOptions(const Driver &D, const ArgList &Args,
   if (!Args.hasArg(options::OPT_fopenacc))
     return;
 
+  D.Diag(diag::warn_openacc_experimental);
   CmdArgs.push_back("-fopenacc");
 
   if (Arg *A = Args.getLastArg(options::OPT_openacc_macro_override)) {

clang/lib/Driver/ToolChains/Flang.cpp

@@ -132,6 +132,11 @@ void Flang::addOtherOptions(const ArgList &Args, ArgStringList &CmdArgs) const {
                    options::OPT_fno_offload_global_filtering,
                    options::OPT_funsigned, options::OPT_fno_unsigned});
 
+  if (Args.hasArg(options::OPT_fopenacc)) {
+     const Driver &D = getToolChain().getDriver();
+     D.Diag(diag::warn_openacc_experimental);
+  }
+
   llvm::codegenoptions::DebugInfoKind DebugInfoKind;
   if (Args.hasArg(options::OPT_gN_Group)) {
     Arg *gNArg = Args.getLastArg(options::OPT_gN_Group);

clang/test/AST/ast-print-openacc-atomic-construct.cpp

@@ -1,5 +1,4 @@
 // RUN: %clang_cc1 -fopenacc -ast-print %s -o - | FileCheck %s
-// XFAIL: *
 
 void foo(int v, int x) {
 // CHECK: #pragma acc atomic read

clang/test/AST/ast-print-openacc-cache-construct.cpp

@@ -1,5 +1,4 @@
 // RUN: %clang_cc1 -fopenacc -ast-print %s -o - | FileCheck %s
-// XFAIL: *
 void foo() {
   int Array[5];
   // CHECK: #pragma acc cache(readonly: Array[1], Array[1:2])

clang/test/AST/ast-print-openacc-declare-construct.cpp

@@ -1,6 +1,5 @@
 // RUN: %clang_cc1 -fopenacc -ast-print %s -o - | FileCheck %s
 // flang decided that we would NOT accept -fopenacc downsteam
-// XFAIL: *
 
 int *Global, *Global2;
 int GlobalArray[5];

clang/test/Misc/warning-flags.c

@@ -18,7 +18,7 @@ This test serves two purposes:
 
 The list of warnings below should NEVER grow.  It should gradually shrink to 0.
 
-CHECK: Warnings without flags (58):
+CHECK: Warnings without flags (59):
 
 CHECK-NEXT:   ext_expected_semi_decl_list
 CHECK-NEXT:   ext_missing_whitespace_after_macro_name
@@ -63,6 +63,7 @@ CHECK-NEXT:   warn_not_compound_assign
 CHECK-NEXT:   warn_objc_property_copy_missing_on_block
 CHECK-NEXT:   warn_objc_protocol_qualifier_missing_id
 CHECK-NEXT:   warn_on_superclass_use
+CHECK-NEXT:   warn_openacc_experimental
 CHECK-NEXT:   warn_pp_convert_to_positive
 CHECK-NEXT:   warn_pp_expr_overflow
 CHECK-NEXT:   warn_pp_line_decimal

flang/test/Semantics/cuf18.cuf

@@ -1,5 +1,4 @@
 ! RUN: %python %S/test_errors.py %s %flang_fc1 -fopenacc
-! XFAIL: *
 subroutine sub1()
   real, allocatable, device :: a(:)
   integer :: i