[OpenMP][AMDGCN] Support OpenMP offloading for AMDGCN architecture - Part 1

[saiislam]

Summary

Allow AMDGCN as a GPU offloading target for OpenMP during compiler
invocation and allow setting CUDAMode for it.

Originally authored by Greg Rodgers (@gregrodgers).

[saiislam]

Fixed lint errors.

[arsenm]

Test?

clang/lib/AST/Decl.cpp

3225

This is also identical to the cuda handling above, can you merge them

[jdoerfert]

clang/lib/AST/Decl.cpp

3225

Yes, wasn't there an idea to have isGPU()?

clang/lib/Frontend/CompilerInvocation.cpp

3159

Can we please not call it CUDA mode for non-CUDA targets? Or do you expect the compilation to "identify" as CUDA compilation?

llvm/include/llvm/ADT/Triple.h

703

What's the difference between an AMDGPU and AMDGCN?

[sameerds]

clang/lib/AST/Decl.cpp

3225

It's not identical, because CUDA is filtering out host code and its check is target independent.

But then, Saiyed, it seems that the new patch disallows library builtins on all languages that reach this point, both on host and device code. Is that the intention? Does this point in the flow preclude any side-effects outside of "OpenMP on AMDGCN"?

llvm/include/llvm/ADT/Triple.h

696

Why not just isAMDGPU()? I myself don't have an opinion either way, but still curious.

[sameerds]

clang/lib/Frontend/CompilerInvocation.cpp

3159

I suspect it's just a lot of water under the bridge. All over Clang, HIP has traditionally co-opted CUDA code without introducing new identifiers, like "-fcuda-is-device". It won't be easy to redo that now, and definitely looks out of scope for this change. A typical HIP compilation usually does identify itself as a HIP compilation like setting the isHIP flag. This allows the frontend to distinguish between CUDA and HIP when it matters.

[yaxunl]

llvm/include/llvm/ADT/Triple.h

703

AMDGPU inlude r600 and amdgcn. r600 are old AMD GPUs. They do not support flat address space therefore cannot support CUDA or HIP.

[sameerds]

llvm/include/llvm/ADT/Triple.h

703

As a separate topic, does that mean that Clang should reject r600 at an earlier entry point itself?

[sameerds]

clang/lib/AST/Decl.cpp

3228

This needs a test.

clang/lib/Frontend/CompilerInvocation.cpp

3100

Needs a test.

3134

This is probably too fundamental to need a test on its own.

[saiislam]

clang/lib/AST/Decl.cpp

3225

@sameerds this function returns a value indicating whether input function corresponds to a builtin function (returns BuiltinID), or not (returns 0) i.e. all conditions returning 0 are the exceptions where a valid BuiltinID can't be returned. The new condition only applies to non-OpenCL (line 3213) builtin functions which are not printf and malloc, and have been targeted for amdgcn.

3225

@jdoerfert Can you please elaborate on isGPU() idea? I am not aware about it.
Guessing by the name, I have added isGPU() as a wrapper over isNVPTX() and isAMDGCN() in the next revision.

3226

this condition is not required because it has already been checked in line #3200 earlier.

clang/lib/Frontend/CompilerInvocation.cpp

3159

@jdoerfert OpenMP support document of clang defines two data sharing modes for cuda devices: CUDA and Generic mode. NVPTX and AMDGCN both are cuda devices. Doesn't it make sense to not refactor CUDA mode as of now?

[saiislam]

Added test cases. Added a wrapper isGPU() for isNVPTX()/isAMDGCN().

[yaxunl]

llvm/include/llvm/ADT/Triple.h

703

I think HIP checks triple and rejects triple other than amdgcn

[arsenm]

llvm/include/llvm/ADT/Triple.h

699

This is skipping out r600, and probably should leave this in clang?

[saiislam]

Moved isGPU() from llvm's Triple.h to clang's TargetInfo. Renamed it to isOpenMPGPU()
to represent target's compatibility with OpenMP offloading and reduce its scope.

[sameerds]

clang/include/clang/Basic/TargetInfo.h

1261 ↗

(On Diff #264192)

How is "OpenMP-compatible GPU" defined? I think it's too early to start designing predicates about whether a target is a GPU and whether it supports OpenMP.

clang/lib/AST/Decl.cpp

3221

This seems awkward to me. Why mix it up with only CUDA and HIP? The earlier factoring is better, where CUDA/HIP took care of their own business, and the catch-all case of AMDGCN was a separate clause by itself. It doesn't matter that the builtins being checked for AMDGCN on OpenMP are currently identical to CUDA/HIP. When this situation later changes (I am sure OpenMP will support more builtins), we will have to split it out again anyway.

clang/lib/Frontend/CompilerInvocation.cpp

3098

I am not particularly in favour of introducing a variable just because it looks smaller than a call at each appropriate location. If you really want it this way, at least make it a const.

3100

Looking at the comment before this line, the correct predicate would "target supports exceptions with OpenMP". Is it always true that every GPU that supports OpenMP will not support exception handling? I would recommend just checking individual targets for now instead of inventing predicates.

3153

Is there any reason to believe that every future GPU added to this predicate will also want the CUDA mode set? I would recommend using individual targets for now instead of inventing a new predicate.

3158

Same doubt about this use of an artificial predicate as commented earlier.

[sameerds]

clang/lib/AST/Decl.cpp

3221

This clause seems to assume that AMDGCN automatically means OpenMP whenever it is not HIP. That does not sound right. Is there a Context.getLangOpts().OpenMP flag? If not, why not? There also seems to be an Opts.OpenMPIsDevice ... perhaps that could be used here to write a separate OpenMP clause?

[saiislam]

Removed isOpenMPGPU() to avoid defining OpenMP compatibility of an architecture.
Reverting back to explicitly checking NVPTX and AMDGCN architectures. Also, split
handling of NVPTX's and AMDGCN's handling of getBuiltinID. For AMDGCN it now uses
OpenMPIsDevice LangOpt and returns 0 for every device library function, except for
printf and malloc.

[sameerds]

clang/lib/AST/Decl.cpp

3228

This still needs a test. One that shows that specific builtins are allowed and others are not.

clang/test/Driver/openmp-offload-gpu.c

257

Is there a functional reason to move these lines? Changes to existing files should be minimized to show only functional changes. Any NFC rearrangement should be a separate patch.

[sameerds]

clang/lib/AST/Decl.cpp

3225

This should say "OpenMP does not have ..."?

[sameerds]

clang/lib/AST/Decl.cpp

3227

Why is the check for AMDGCN required here? Doesn't the language define the set of supported builtins in a language-independent way? At least that seems to be true for OpenCL and CUDA above.

[saiislam]

Fixed typo. Added test case to check for device side functions.

[saiislam]

Added test case to show treatment of specific functions as builtins or functions on the device

[saiislam]

clang/lib/AST/Decl.cpp

3227

AMDGCN check is required because this block deals with the openmp target when it is amdgcn. A future or existing target device might decide to deal with library calls in its own way.

3228

Added test case to show difference in treatment of printf() and other predefined library functions, on the device.

[saiislam]

Shifted test cases in openmp-offload-gpu.c for better visual segmentation.
Updated device function test case to be more accurate.

[jdoerfert]

I'm generally fine with this, don't wait for my approval.

clang/test/OpenMP/amdgcn_device_function_call.cpp

27 ↗

(On Diff #266171)

Not for this patch:
FWIW, we will need to make math functions work inside target regions. The way aomp does it is afaik different from the way we do it. We can however adopt our way for this target though. Feel free to ping me on this later.

[saiislam]

clang/test/OpenMP/amdgcn_device_function_call.cpp

27 ↗

(On Diff #266171)

Yes, sure. I will ping you as we move towards it. Thanks.

[sameerds]

LGTM!

clang/lib/AST/Decl.cpp

if (!ConsiderWrapperFunctions && getStorageClass() == SC_Static)
     return 0;
 
   // OpenCL v1.2 s6.9.f - The library functions defined in
   // the C99 standard headers are not available.
   if (Context.getLangOpts().OpenCL &&
       Context.BuiltinInfo.isPredefinedLibFunction(BuiltinID))
     return 0;
 
-  // CUDA does not have device-side standard library. printf and malloc are the
-  // only special cases that are supported by device-side runtime.
-  if (Context.getLangOpts().CUDA && hasAttr<CUDADeviceAttr>() &&
-      !hasAttr<CUDAHostAttr>() &&
+  // CUDA/HIP does not have device-side standard library. printf and malloc are
+  // the only special cases that are supported by device-side runtime.
+  if (((Context.getLangOpts().CUDA && hasAttr<CUDADeviceAttr>() &&
+        !hasAttr<CUDAHostAttr>()) ||
+       Context.getTargetInfo().getTriple().isAMDGCN()) &&
       !(BuiltinID == Builtin::BIprintf || BuiltinID == Builtin::BImalloc))
     return 0;
 
   return BuiltinID;
 }
 
 /// getNumParams - Return the number of parameters this function must have
 /// based on its FunctionType.  This is the length of the ParamInfo array
 /// after it has been created.
 unsigned FunctionDecl::getNumParams() const {
   const auto *FPT = getType()->getAs<FunctionProtoType>();
   return FPT ? FPT->getNumParams() : 0;
 }
 
 void FunctionDecl::setParams(ASTContext &C,

clang/lib/Frontend/CompilerInvocation.cpp

if (!Opts.OpenMPIsDevice) {
       case llvm::Triple::nvptx64:
         Diags.Report(diag::err_drv_omp_host_target_not_supported)
             << TargetOpts.Triple;
         break;
       }
     }
   }
 
   // Set the flag to prevent the implementation from emitting device exception
   // handling code for those requiring so.
-  if ((Opts.OpenMPIsDevice && T.isNVPTX()) || Opts.OpenCLCPlusPlus) {
+  if ((Opts.OpenMPIsDevice && T.isGPU()) || Opts.OpenCLCPlusPlus) {
     Opts.Exceptions = 0;
     Opts.CXXExceptions = 0;
   }
   if (Opts.OpenMPIsDevice && T.isNVPTX()) {
     Opts.OpenMPCUDANumSMs =
         getLastArgIntValue(Args, options::OPT_fopenmp_cuda_number_of_sm_EQ,
                            Opts.OpenMPCUDANumSMs, Diags);
     Opts.OpenMPCUDABlocksPerSM =
for (unsigned i = 0; i < A->getNumValues(); ++i) {
 
       if (TT.getArch() == llvm::Triple::UnknownArch ||
           !(TT.getArch() == llvm::Triple::aarch64 ||
             TT.getArch() == llvm::Triple::ppc ||
             TT.getArch() == llvm::Triple::ppc64 ||
             TT.getArch() == llvm::Triple::ppc64le ||
             TT.getArch() == llvm::Triple::nvptx ||
             TT.getArch() == llvm::Triple::nvptx64 ||
+            TT.getArch() == llvm::Triple::amdgcn ||
             TT.getArch() == llvm::Triple::x86 ||
             TT.getArch() == llvm::Triple::x86_64))
         Diags.Report(diag::err_drv_invalid_omp_target) << A->getValue(i);
       else
         Opts.OMPTargetTriples.push_back(TT);
     }
   }
 
   // Get OpenMP host file path if any and report if a non existent file is
   // found
   if (Arg *A = Args.getLastArg(options::OPT_fopenmp_host_ir_file_path)) {
     Opts.OMPHostIRFile = A->getValue();
     if (!llvm::sys::fs::exists(Opts.OMPHostIRFile))
       Diags.Report(diag::err_drv_omp_host_ir_file_not_found)
           << Opts.OMPHostIRFile;
   }
 
-  // Set CUDA mode for OpenMP target NVPTX if specified in options
-  Opts.OpenMPCUDAMode = Opts.OpenMPIsDevice && T.isNVPTX() &&
+  // Set CUDA mode for OpenMP target NVPTX/AMDGCN if specified in options
+  Opts.OpenMPCUDAMode = Opts.OpenMPIsDevice && T.isGPU() &&
                         Args.hasArg(options::OPT_fopenmp_cuda_mode);
 
-  // Set CUDA mode for OpenMP target NVPTX if specified in options
+  // Set CUDA mode for OpenMP target NVPTX/AMDGCN if specified in options
   Opts.OpenMPCUDAForceFullRuntime =
-      Opts.OpenMPIsDevice && T.isNVPTX() &&
+      Opts.OpenMPIsDevice && T.isGPU() &&
       Args.hasArg(options::OPT_fopenmp_cuda_force_full_runtime);
 
   // Record whether the __DEPRECATED define was requested.
   Opts.Deprecated = Args.hasFlag(OPT_fdeprecated_macro,
                                  OPT_fno_deprecated_macro,
                                  Opts.Deprecated);
 
   // FIXME: Eliminate this dependency.
   unsigned Opt = getOptimizationLevel(Args, IK, Diags),

clang/test/Driver/openmp-offload-gpu.c

 ///
 /// Perform several driver tests for OpenMP offloading
 ///
 
 // REQUIRES: clang-driver
 // REQUIRES: x86-registered-target
 // REQUIRES: powerpc-registered-target
 // REQUIRES: nvptx-registered-target
+// REQUIRES: amdgpu-registered-target
 
 /// ###########################################################################
 
 /// Check -Xopenmp-target uses one of the archs provided when several archs are used.
 // RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda \
 // RUN:          -Xopenmp-target -march=sm_35 -Xopenmp-target -march=sm_60 %s 2>&1 \
 // RUN:   | FileCheck -check-prefix=CHK-FOPENMP-TARGET-ARCHS %s
 

 // HAS_DEBUG-SAME: "-dwarf-version=2"
 // HAS_DEBUG-SAME: "-fopenmp-is-device"
 // HAS_DEBUG: ptxas
 // HAS_DEBUG-SAME: "-g"
 // HAS_DEBUG-SAME: "--dont-merge-basicblocks"
 // HAS_DEBUG-SAME: "--return-at-end"
 // HAS_DEBUG: nvlink
 // HAS_DEBUG-SAME: "-g"
+// CUDA_MODE: clang{{.*}}"-cc1"{{.*}}"-triple" "{{nvptx64-nvidia-cuda|amdgcn-amd-amdhsa}}"
+// CUDA_MODE-SAME: "-fopenmp-cuda-mode"
+// NO_CUDA_MODE-NOT: "-{{fno-|f}}openmp-cuda-mode"
 
 // RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_60 %s -fopenmp-cuda-mode 2>&1 \
 // RUN:   | FileCheck -check-prefix=CUDA_MODE %s
 // RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_60 %s -fno-openmp-cuda-mode -fopenmp-cuda-mode 2>&1 \
 // RUN:   | FileCheck -check-prefix=CUDA_MODE %s
-// CUDA_MODE: clang{{.*}}"-cc1"{{.*}}"-triple" "nvptx64-nvidia-cuda"
-// CUDA_MODE-SAME: "-fopenmp-cuda-mode"
 // RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_60 %s -fno-openmp-cuda-mode 2>&1 \
 // RUN:   | FileCheck -check-prefix=NO_CUDA_MODE %s
 // RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_60 %s -fopenmp-cuda-mode -fno-openmp-cuda-mode 2>&1 \
 // RUN:   | FileCheck -check-prefix=NO_CUDA_MODE %s
-// NO_CUDA_MODE-NOT: "-{{fno-|f}}openmp-cuda-mode"
+
+// RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target -march=gfx906 %s -fopenmp-cuda-mode 2>&1 \
+// RUN:   | FileCheck -check-prefix=CUDA_MODE %s
+// RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target -march=gfx906 %s -fno-openmp-cuda-mode -fopenmp-cuda-mode 2>&1 \
+// RUN:   | FileCheck -check-prefix=CUDA_MODE %s
+// RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target -march=gfx906 %s -fno-openmp-cuda-mode 2>&1 \
+// RUN:   | FileCheck -check-prefix=NO_CUDA_MODE %s
+// RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target -march=gfx906 %s -fopenmp-cuda-mode -fno-openmp-cuda-mode 2>&1 \
+// RUN:   | FileCheck -check-prefix=NO_CUDA_MODE %s
+
+// FULL_RUNTIME: clang{{.*}}"-cc1"{{.*}}"-triple" "{{nvptx64-nvidia-cuda|amdgcn-amd-amdhsa}}"
+// FULL_RUNTIME-SAME: "-fopenmp-cuda-force-full-runtime"
+// NO_FULL_RUNTIME-NOT: "-{{fno-|f}}openmp-cuda-force-full-runtime"
 
 // RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_60 %s -fopenmp-cuda-force-full-runtime 2>&1 \
 // RUN:   | FileCheck -check-prefix=FULL_RUNTIME %s
 // RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_60 %s -fno-openmp-cuda-force-full-runtime -fopenmp-cuda-force-full-runtime 2>&1 \
 // RUN:   | FileCheck -check-prefix=FULL_RUNTIME %s
-// FULL_RUNTIME: clang{{.*}}"-cc1"{{.*}}"-triple" "nvptx64-nvidia-cuda"
-// FULL_RUNTIME-SAME: "-fopenmp-cuda-force-full-runtime"
 // RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_60 %s -fno-openmp-cuda-force-full-runtime 2>&1 \
 // RUN:   | FileCheck -check-prefix=NO_FULL_RUNTIME %s
 // RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_60 %s -fopenmp-cuda-force-full-runtime -fno-openmp-cuda-force-full-runtime 2>&1 \
 // RUN:   | FileCheck -check-prefix=NO_FULL_RUNTIME %s
-// NO_FULL_RUNTIME-NOT: "-{{fno-|f}}openmp-cuda-force-full-runtime"
+
+// RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target -march=gfx906 %s -fopenmp-cuda-force-full-runtime 2>&1 \
+// RUN:   | FileCheck -check-prefix=FULL_RUNTIME %s
+// RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target -march=gfx906 %s -fno-openmp-cuda-force-full-runtime -fopenmp-cuda-force-full-runtime 2>&1 \
+// RUN:   | FileCheck -check-prefix=FULL_RUNTIME %s
+// RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target -march=gfx906 %s -fno-openmp-cuda-force-full-runtime 2>&1 \
+// RUN:   | FileCheck -check-prefix=NO_FULL_RUNTIME %s
+// RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target -march=gfx906 %s -fopenmp-cuda-force-full-runtime -fno-openmp-cuda-force-full-runtime 2>&1 \
+// RUN:   | FileCheck -check-prefix=NO_FULL_RUNTIME %s
 
 // RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_60 %s -fopenmp-cuda-teams-reduction-recs-num=2048 2>&1 \
 // RUN:   | FileCheck -check-prefix=CUDA_RED_RECS %s
 // CUDA_RED_RECS: clang{{.*}}"-cc1"{{.*}}"-triple" "nvptx64-nvidia-cuda"
 // CUDA_RED_RECS-SAME: "-fopenmp-cuda-teams-reduction-recs-num=2048"
 
 // RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda %s 2>&1 \
 // RUN:   | FileCheck -check-prefix=OPENMP_NVPTX_WRAPPERS %s
 // OPENMP_NVPTX_WRAPPERS: clang{{.*}}"-cc1"{{.*}}"-triple" "nvptx64-nvidia-cuda"
 // OPENMP_NVPTX_WRAPPERS-SAME: "-internal-isystem" "{{.*}}openmp_wrappers"

clang/test/OpenMP/target_parallel_no_exceptions.cpp

 /// Make sure no exception messages are inclided in the llvm output.
 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple nvptx64-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 --check-prefix CHK-EXCEPTION
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CHK-EXCEPTION
 
 void test_increment() {
 #pragma omp target
   {
     []() { return; }();
   }
 }
 
 int main() {
   test_increment();
   return 0;
 }
 
 //CHK-EXCEPTION-NOT: invoke
 

llvm/include/llvm/ADT/Triple.h

bool isSPIR() const {
     return getArch() == Triple::spir || getArch() == Triple::spir64;
   }
 
   /// Tests whether the target is NVPTX (32- or 64-bit).
   bool isNVPTX() const {
     return getArch() == Triple::nvptx || getArch() == Triple::nvptx64;
   }
 
+  /// Tests whether the target is AMDGCN
+  bool isAMDGCN() const { return getArch() == Triple::amdgcn; }
+
+  /// Tests whether the target is a GPU i.e. NVPTX or AMDGCN
+  bool isGPU() const { return isNVPTX() || isAMDGCN(); }
+
   bool isAMDGPU() const {
     return getArch() == Triple::r600 || getArch() == Triple::amdgcn;
   }
 
   /// Tests whether the target is Thumb (little and big endian).
   bool isThumb() const {
     return getArch() == Triple::thumb || getArch() == Triple::thumbeb;
   }
 
   /// Tests whether the target is ARM (little and big endian).
   bool isARM() const {