[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

3227

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

[jdoerfert]

clang/lib/AST/Decl.cpp

3227

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

clang/lib/Frontend/CompilerInvocation.cpp

3172

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

700

What's the difference between an AMDGPU and AMDGCN?

[sameerds]

clang/lib/AST/Decl.cpp

3227

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

3172

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

700

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

700

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

[sameerds]

clang/lib/AST/Decl.cpp

3230

This needs a test.

clang/lib/Frontend/CompilerInvocation.cpp

3112–3113

Needs a test.

3147

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

[saiislam]

clang/lib/AST/Decl.cpp

3227

@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.

3227

@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.

3228

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

clang/lib/Frontend/CompilerInvocation.cpp

3172

@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

700

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

3227

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

3110

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.

3112–3113

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.

3166

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.

3171

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

[sameerds]

clang/lib/AST/Decl.cpp

3227

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

3230

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

3227

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

[sameerds]

clang/lib/AST/Decl.cpp

3229

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

3229

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.

3230

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

28

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

28

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

[sameerds]

LGTM!

clang/lib/AST/Decl.cpp

unsigned FunctionDecl::getBuiltinID(bool ConsiderWrapperFunctions) const {
 
   // 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>() &&
       !(BuiltinID == Builtin::BIprintf || BuiltinID == Builtin::BImalloc))
     return 0;
 
+  // As AMDGCN implementation of OpenMP does not have a device-side standard
+  // library, none of the predefined library functions except printf and malloc
+  // should be treated as a builtin i.e. 0 should be returned for them.
+  if (Context.getTargetInfo().getTriple().isAMDGCN() &&
+      Context.getLangOpts().OpenMPIsDevice &&
+      Context.BuiltinInfo.isPredefinedLibFunction(BuiltinID) &&
+      !(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>();

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.isNVPTX() || T.isAMDGCN())) ||
+      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.isNVPTX() || T.isAMDGCN()) &&
                         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.isNVPTX() || T.isAMDGCN()) &&
       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: "--return-at-end"
 // HAS_DEBUG: nvlink
 // HAS_DEBUG-SAME: "-g"
 
 // 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"
+// 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
+// CUDA_MODE: clang{{.*}}"-cc1"{{.*}}"-triple" "{{nvptx64-nvidia-cuda|amdgcn-amd-amdhsa}}"
 // 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
+// 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
 // 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-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"
+// 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
+// FULL_RUNTIME: clang{{.*}}"-cc1"{{.*}}"-triple" "{{nvptx64-nvidia-cuda|amdgcn-amd-amdhsa}}"
 // 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
+// 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
 // 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-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/amdgcn_device_function_call.cpp

+// REQUIRES: amdgpu-registered-target
+
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple amdgcn-amd-amdhsa -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s
+// RUN: llvm-dis < %t-ppc-host.bc | FileCheck %s -check-prefix=HOST
+
+// device side declarations
+#pragma omp declare target
+extern "C" float cosf(float __x);
+#pragma omp end declare target
+
+// host side declaration
+extern "C" float cosf(float __x);
+
+void test_amdgcn_openmp_device(float __x) {
+  // the default case where predefined library functions are treated as
+  // builtins on the host
+  // HOST: call float @llvm.cos.f32(float
+  __x = cosf(__x);
+
+#pragma omp target
+  {
+    // cosf should not be treated as builtin on device
+    // CHECK-NOT: call float @llvm.cos.f32(float
+    __x = cosf(__x);
+  }
+}

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; }
+
   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 {