[AMDGPU][OpenMP] Enable Libomptarget runtime tests

[pdhaliwal]

Summary

This enables the runtime tests on amdgpu targets.
10 tests have been marked as XFAIL on amdgcn currently mostly due to
missing printf.

[JonChesterfield]

I guess there's no XFAIL equivalent here? In that case we should probably leave off the RUN line for the tests that can't work, as they'll otherwise break the build once an amdgpu CI machine goes live

openmp/libomptarget/CMakeLists.txt

40

I don't understand this string. Maybe in '-march=$LIBOMPTARGET_AMDGCN_TEST_TARGET'? Seems bad to be hard coding a gfx908 here, presumably it should be whatever is in the system running the tests

openmp/libomptarget/deviceRTLs/CMakeLists.txt

12

would rather enable this unconditionally, which works today except for people who build with the amdgpu target disabled. See D98746

openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt

94

This we can patch separately. Maybe also drop the gfx7? I don't think it's being tested, and if it's broken it would be friendlier to not build it by default.

openmp/libomptarget/plugins/CMakeLists.txt

80

This seems to be a regression, why not continue building the plugin unconditionally?

openmp/libomptarget/test/lit.site.cfg.in

19

this seems very inconvenient for the tests, I think we should handle missing march by looking at the system we're running the tests on

[pdhaliwal]

openmp/libomptarget/CMakeLists.txt

40

Its just a default cmake option value. Since there is no automatic march detection, I thought of having it configurable from command line. I will need dig around for detecting system gpu.

openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt

94

I will create a separate patch for it.

openmp/libomptarget/plugins/CMakeLists.txt

80

This was causing issues when plugin was always getting built but deviceRTL not. When plugin is built, the runtime tests will be executed which are going to fail.

openmp/libomptarget/test/lit.site.cfg.in

19

Agreed. This seems like a problem which could be solved by Driver/ToolChain.

[pdhaliwal]

• Addressed review comments.

[pdhaliwal]

Rebase

[protze.joachim]

In D99656#2661379, @JonChesterfield wrote:

I guess there's no XFAIL equivalent here? In that case we should probably leave off the RUN line for the tests that can't work, as they'll otherwise break the build once an amdgpu CI machine goes live

It's a general problem of the current libomptarget tests. A test fails if any of the run lines fails and will not check the remaining lines.
I think, it would be better to run the tests individually for each available target. Would it be possible to run the test for each supported target in a different work directory?

[pdhaliwal]

Add four more tests.

[pdhaliwal]

Remove whitespace change.

[pdhaliwal]

In D99656#2715982, @protze.joachim wrote:

In D99656#2661379, @JonChesterfield wrote:

I guess there's no XFAIL equivalent here? In that case we should probably leave off the RUN line for the tests that can't work, as they'll otherwise break the build once an amdgpu CI machine goes live

It's a general problem of the current libomptarget tests. A test fails if any of the run lines fails and will not check the remaining lines.
I think, it would be better to run the tests individually for each available target. Would it be possible to run the test for each supported target in a different work directory?

With properly configured cmake and lit, it should be possible to add tests run separately for architectures. For e.g. it should be possible to create targets like check-libomptarget-amdgcn-amd-amdhsa which would run tests only for amdgcn and similar for others. These can be grouped under existing check-libomptarget.

[protze.joachim]

In D99656#2716049, @pdhaliwal wrote:

With properly configured cmake and lit, it should be possible to add tests run separately for architectures. For e.g. it should be possible to create targets like check-libomptarget-amdgcn-amd-amdhsa which would run tests only for amdgcn and similar for others. These can be grouped under existing check-libomptarget.

I posted D101315 and D101326. I suggest that you rebase on these two patches. This should cut down the necessary changes for your patch.

You might want to mark unsupported tests with // XFAIL: amdgcn-amd-amdhsa (currently failing, but will be fix in future) or // UNSUPPORTED: amdgcn-amd-amdhsa (no plan to support the test)

[pdhaliwal]

Rebased.

[JonChesterfield]

Patch is much improved by Joachim's refactor!

Looks good to me, will leave the check box to someone who has edited these tests before. Bringing printf online looks like a useful thing to do soon.

openmp/libomptarget/test/unified_shared_memory/close_enter_exit.c

4 ↗

(On Diff #341080)

Inconsistent with above. I.e. XFAIL on two lines or one XFAIL with comma delimited. No preference either way here, but probably stay with a convention

[protze.joachim]

It might help in the future, if you add comments for each XFAIL you add, saying what feature is missing to make the test case work (like you did with printf).
-> This would naturally result in a separate XFAIL line for each offloading target.

[pdhaliwal]

Add failure reason for the failing tests.

[protze.joachim]

Thanks!
lgtm

[JonChesterfield]

Where's the logic to skip running these tests if ./amdgpu-arch fails?

[protze.joachim]

In D99656#2739517, @JonChesterfield wrote:

Where's the logic to skip running these tests if ./amdgpu-arch fails?

The tests are included to check-libomptarget, if llvm-project/openmp/libomptarget/plugins/amdgpu/CMakeLists.txt reaches the last line, i.e., if the plugin is built.

[JonChesterfield]

In D99656#2739599, @protze.joachim wrote:

In D99656#2739517, @JonChesterfield wrote:

Where's the logic to skip running these tests if ./amdgpu-arch fails?

The tests are included to check-libomptarget, if llvm-project/openmp/libomptarget/plugins/amdgpu/CMakeLists.txt reaches the last line, i.e., if the plugin is built.

Ah, that's not ideal. Picked up the same behaviour as nvptx then.

The plugin existing is necessary but not sufficient for running the tests to succeed. One requires the plugin, but also the hardware available locally.

amdgpu-arch returning 0 will correlate strongly with there being a local gpu that tests can run on. Do we have a good place to put a call to that which would control whether to run the amdgpu tests?

We should be able to use a successful call to one of the binaries in the cuda toolkit to do the same test for nvidia, so that people with cuda installed but no nvidia gpu don't run the nvptx tests.

openmp/libomptarget/CMakeLists.txt

 # Get dependencies for the different components of the project.
 include(LibomptargetGetDependencies)
 
 # LLVM source tree is required at build time for libomptarget
 if (NOT LIBOMPTARGET_LLVM_INCLUDE_DIRS)
   message(FATAL_ERROR "Missing definition for LIBOMPTARGET_LLVM_INCLUDE_DIRS")
 endif()
 
+# options to control amdgcn building and testing
+option(LIBOMPTARGET_ENABLE_EXPERIMENTAL_AMDGCN_TARGET "Enable building plugin and device runtime library for amdgpu" Off)
+option(LIBOMPTARGET_AMDGCN_TEST_TARGET "AMDGPU target to be used in -march=gfx906 when running runtime tests" "gfx908")
+
 # This is a list of all the targets that are supported/tested right now.
 set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} aarch64-unknown-linux-gnu")
+set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} amdgcn-amd-amdhsa")
 set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} powerpc64le-ibm-linux-gnu")
 set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} powerpc64-ibm-linux-gnu")
 set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} x86_64-pc-linux-gnu")
 set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} nvptx64-nvidia-cuda")
 
 # Once the plugins for the different targets are validated, they will be added to
 # the list of supported targets in the current system.
 set (LIBOMPTARGET_SYSTEM_TARGETS "")

 add_subdirectory(${LIBOMPTARGET_SRC_DIR})
 
 # Definitions for testing, for reuse when testing libomptarget-nvptx.
 set(LIBOMPTARGET_OPENMP_HEADER_FOLDER "${LIBOMP_INCLUDE_DIR}" CACHE STRING
   "Path to folder containing omp.h")
 set(LIBOMPTARGET_OPENMP_HOST_RTL_FOLDER "${LIBOMP_LIBRARY_DIR}" CACHE STRING
   "Path to folder containing libomp.so, and libLLVMSupport.so with profiling enabled")
 
-
 # Build offloading plugins and device RTLs if they are available.
 add_subdirectory(plugins)
 add_subdirectory(deviceRTLs)
 
 # Add tests.
 add_subdirectory(test)

openmp/libomptarget/deviceRTLs/CMakeLists.txt

 ##===----------------------------------------------------------------------===##
 #
 # Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 # See https://llvm.org/LICENSE.txt for license information.
 # SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 #
 # ##===----------------------------------------------------------------------===##
 #
 # Build a device RTL for each available machine.
 #
 ##===----------------------------------------------------------------------===##
-
+if(LIBOMPTARGET_ENABLE_EXPERIMENTAL_AMDGCN_TARGET)
+    add_subdirectory(amdgcn)
+endif()
 add_subdirectory(nvptx)

openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt

 # for both in-tree and out-of-tree build
 if (NOT CMAKE_ARCHIVE_OUTPUT_DIRECTORY)
   set(OUTPUTDIR ${CMAKE_CURRENT_BINARY_DIR})
 else()
   set(OUTPUTDIR ${CMAKE_ARCHIVE_OUTPUT_DIRECTORY})
 endif()
 
 # create libraries
-set(mcpus gfx700 gfx701 gfx801 gfx803 gfx900 gfx906)
+set(mcpus gfx700 gfx701 gfx801 gfx803 gfx900 gfx906 gfx908)
 if (DEFINED LIBOMPTARGET_AMDGCN_GFXLIST)
   set(mcpus ${LIBOMPTARGET_AMDGCN_GFXLIST})
 endif()
 
 macro(add_cuda_bc_library)
   set(cu_cmd ${AOMP_BINDIR}/clang++
     -xc++
     -c

openmp/libomptarget/plugins/CMakeLists.txt

else(LIBOMPTARGET_DEP_LIBELF_FOUND)
     libomptarget_say("Not building ${tmachine_name} offloading plugin: libelf dependency not found.")
   endif(LIBOMPTARGET_DEP_LIBELF_FOUND)
 else()
   libomptarget_say("Not building ${tmachine_name} offloading plugin: machine not found in the system.")
 endif()
 endmacro()
 
 add_subdirectory(aarch64)
-add_subdirectory(amdgpu)
 add_subdirectory(cuda)
 add_subdirectory(ppc64)
 add_subdirectory(ppc64le)
 add_subdirectory(ve)
 add_subdirectory(x86_64)
 add_subdirectory(remote)
 
+if(LIBOMPTARGET_ENABLE_EXPERIMENTAL_AMDGCN_TARGET)
+    add_subdirectory(amdgpu)
+endif()
+
+
 # Make sure the parent scope can see the plugins that will be created.
 set(LIBOMPTARGET_SYSTEM_TARGETS "${LIBOMPTARGET_SYSTEM_TARGETS}" PARENT_SCOPE)
 set(LIBOMPTARGET_TESTED_PLUGINS "${LIBOMPTARGET_TESTED_PLUGINS}" PARENT_SCOPE)
 

openmp/libomptarget/test/env/base_ptr_ref_count.c

 // RUN: %libomptarget-compile-aarch64-unknown-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 | %fcheck-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 | %fcheck-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 | %fcheck-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-compile-x86_64-pc-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-x86_64-pc-linux-gnu 2>&1 | %fcheck-x86_64-pc-linux-gnu
 // RUN: %libomptarget-compile-nvptx64-nvidia-cuda && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-nvptx64-nvidia-cuda 2>&1 | %fcheck-nvptx64-nvidia-cuda
+// RUN: %libomptarget-compile-amdgcn-amd-amdhsa && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-amdgcn-amd-amdhsa 2>&1 | %fcheck-amdgcn-amd-amdhsa
+
 // REQUIRES: libomptarget-debug
 
 #include <stdlib.h>
 
 int *allocate(size_t n) {
   int *ptr = malloc(sizeof(int) * n);
 #pragma omp target enter data map(to : ptr[:n])
   return ptr;

openmp/libomptarget/test/env/omp_target_debug.c

 // RUN: %libomptarget-compile-aarch64-unknown-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 | %fcheck-aarch64-unknown-linux-gnu -allow-empty -check-prefix=DEBUG
 // RUN: %libomptarget-compile-aarch64-unknown-linux-gnu && env LIBOMPTARGET_DEBUG=0 %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 | %fcheck-aarch64-unknown-linux-gnu -allow-empty -check-prefix=NDEBUG
 // RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 | %fcheck-powerpc64-ibm-linux-gnu -allow-empty -check-prefix=DEBUG
 // RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=0 %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 | %fcheck-powerpc64-ibm-linux-gnu -allow-empty -check-prefix=NDEBUG
 // RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 | %fcheck-powerpc64le-ibm-linux-gnu -allow-empty -check-prefix=DEBUG
 // RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=0 %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 | %fcheck-powerpc64le-ibm-linux-gnu -allow-empty -check-prefix=NDEBUG
 // RUN: %libomptarget-compile-x86_64-pc-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-x86_64-pc-linux-gnu 2>&1 | %fcheck-x86_64-pc-linux-gnu -allow-empty -check-prefix=DEBUG
 // RUN: %libomptarget-compile-x86_64-pc-linux-gnu && env LIBOMPTARGET_DEBUG=0 %libomptarget-run-x86_64-pc-linux-gnu 2>&1 | %fcheck-x86_64-pc-linux-gnu -allow-empty -check-prefix=NDEBUG
 // RUN: %libomptarget-compile-nvptx64-nvidia-cuda && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-nvptx64-nvidia-cuda 2>&1 | %fcheck-nvptx64-nvidia-cuda -allow-empty -check-prefix=DEBUG
 // RUN: %libomptarget-compile-nvptx64-nvidia-cuda && env LIBOMPTARGET_DEBUG=0 %libomptarget-run-nvptx64-nvidia-cuda 2>&1 | %fcheck-nvptx64-nvidia-cuda -allow-empty -check-prefix=NDEBUG
+// RUN: %libomptarget-compile-amdgcn-amd-amdhsa && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-amdgcn-amd-amdhsa 2>&1 | %fcheck-amdgcn-amd-amdhsa -allow-empty -check-prefix=DEBUG
+// RUN: %libomptarget-compile-amdgcn-amd-amdhsa && env LIBOMPTARGET_DEBUG=0 %libomptarget-run-amdgcn-amd-amdhsa 2>&1 | %fcheck-amdgcn-amd-amdhsa -allow-empty -check-prefix=NDEBUG
 // REQUIRES: libomptarget-debug
 
 int main(void) {
 #pragma omp target
   {}
   return 0;
 }
 
 // DEBUG: Libomptarget
 // NDEBUG-NOT: Libomptarget
 // NDEBUG-NOT: Target
 

openmp/libomptarget/test/lit.cfg

 # - for targets that exist in the system create the actual command.
 # - for valid targets that do not exist in the system, return false, so that the
 #   same test can be used for different targets.
 
 # Scan all the valid targets.
 for libomptarget_target in config.libomptarget_all_targets:
     # Is this target in the current system? If so create a compile, run and test
     # command. Otherwise create command that return false.
+    xopenmp_target_flags = ''
+    if libomptarget_target == 'amdgcn-amd-amdhsa':
+        if config.libomptarget_amdgcn_test_target:
+            xopenmp_target_flags = '-Xopenmp-target=amdgcn-amd-amdhsa -march=' + config.libomptarget_amdgcn_test_target
+
     if libomptarget_target in config.libomptarget_system_targets:
         config.substitutions.append(("%libomptarget-compilexx-run-and-check-" + \
             libomptarget_target, \
             "%libomptarget-compilexx-and-run-" + libomptarget_target + \
             " | " + config.libomptarget_filecheck + " %s"))
         config.substitutions.append(("%libomptarget-compile-run-and-check-" + \
             libomptarget_target, \
             "%libomptarget-compile-and-run-" + libomptarget_target + \
if libomptarget_target in config.libomptarget_system_targets:
             libomptarget_target))
         config.substitutions.append(("%libomptarget-run-" + \
             libomptarget_target, \
             "%t-" + libomptarget_target))
         config.substitutions.append(("%libomptarget-run-fail-" + \
             libomptarget_target, \
             "%not --crash %t-" + libomptarget_target))
         config.substitutions.append(("%clangxx-" + libomptarget_target, \
-            "%clangxx %openmp_flags %flags -fopenmp-targets=" + libomptarget_target))
+                                     "%clangxx %openmp_flags %flags -fopenmp-targets=" + libomptarget_target \
+                                     + " " + xopenmp_target_flags))
         config.substitutions.append(("%clang-" + libomptarget_target, \
-            "%clang %openmp_flags %flags -fopenmp-targets=" + libomptarget_target))
+                                     "%clang %openmp_flags %flags -fopenmp-targets=" + libomptarget_target \
+                                     + " " + xopenmp_target_flags))
         config.substitutions.append(("%fcheck-" + libomptarget_target, \
             config.libomptarget_filecheck + " %s"))
     else:
         config.substitutions.append(("%libomptarget-compile-run-and-check-" + \
             libomptarget_target, \
             "echo ignored-command"))
         config.substitutions.append(("%libomptarget-compilexx-run-and-check-" + \
             libomptarget_target, \

openmp/libomptarget/test/lit.site.cfg.in

 config.omp_host_rtl_directory = "@LIBOMPTARGET_OPENMP_HOST_RTL_FOLDER@"
 config.operating_system = "@CMAKE_SYSTEM_NAME@"
 config.libomptarget_all_targets = "@LIBOMPTARGET_ALL_TARGETS@".split()
 config.libomptarget_system_targets = "@LIBOMPTARGET_SYSTEM_TARGETS@".split()
 config.libomptarget_filecheck = "@OPENMP_FILECHECK_EXECUTABLE@"
 config.libomptarget_not = "@OPENMP_NOT_EXECUTABLE@"
 config.libomptarget_debug = @LIBOMPTARGET_DEBUG@
 
+# amdgcn target requires -march option to be passed explicitly
+config.libomptarget_amdgcn_test_target = "@LIBOMPTARGET_AMDGCN_TEST_TARGET@"
+
 # Let the main config do the real work.
 lit_config.load_config(config, "@CMAKE_CURRENT_SOURCE_DIR@/lit.cfg")

openmp/libomptarget/test/mapping/data_absent_at_exit.c

 // RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-nvptx64-nvidia-cuda
+// RUN: %libomptarget-compile-run-and-check-amdgcn-amd-amdhsa
 
 #include <stdio.h>
 
 // OpenMP 5.1, sec. 2.21.7.1 "map Clause", p. 351 L14-16:
 // "If the map clause appears on a target, target data, or target exit data
 // construct and a corresponding list item of the original list item is not
 // present in the device data environment on exit from the region then the
 // list item is ignored."

openmp/libomptarget/test/mapping/declare_mapper_api.cpp

 // RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
 // RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
+// RUN: %libomptarget-compilexx-run-and-check-amdgcn-amd-amdhsa
 
 #include <cstdio>
 #include <cstdlib>
 #include <vector>
 #include <cinttypes>
 
 // Data structure definitions copied from OpenMP RTL.
 struct MapComponentInfoTy {

openmp/libomptarget/test/mapping/declare_mapper_target.cpp

 // RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-nvptx64-nvidia-cuda
+// RUN: %libomptarget-compile-run-and-check-amdgcn-amd-amdhsa
 
 #include <cstdio>
 #include <cstdlib>
 
 #define NUM 1024
 
 class C {
 public:

openmp/libomptarget/test/mapping/declare_mapper_target_data.cpp

 // RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-nvptx64-nvidia-cuda
+// RUN: %libomptarget-compile-run-and-check-amdgcn-amd-amdhsa
 
 #include <cstdio>
 #include <cstdlib>
 
 #define NUM 1024
 
 class C {
 public:

openmp/libomptarget/test/mapping/declare_mapper_target_data_enter_exit.cpp

 // RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-nvptx64-nvidia-cuda
+// RUN: %libomptarget-compile-run-and-check-amdgcn-amd-amdhsa
 
 #include <cstdio>
 #include <cstdlib>
 
 #define NUM 1024
 
 class C {
 public:

openmp/libomptarget/test/mapping/declare_mapper_target_update.cpp

 // RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-nvptx64-nvidia-cuda
+// RUN: %libomptarget-compile-run-and-check-amdgcn-amd-amdhsa
 
 #include <cstdio>
 #include <cstdlib>
 
 #define NUM 1024
 
 class C {
 public:

openmp/libomptarget/test/mapping/lambda_mapping.cpp

 // RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
 // RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
+// RUN: %libomptarget-compilexx-run-and-check-amdgcn-amd-amdhsa
 
 #include <iostream>
 
 template <typename LOOP_BODY>
 inline void forall(int Begin, int End, LOOP_BODY LoopBody) {
 #pragma omp target parallel for schedule(static)
   for (int I = Begin; I < End; ++I) {
     LoopBody(I);

openmp/libomptarget/test/mapping/pr38704.c

 // RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-nvptx64-nvidia-cuda
+// RUN: %libomptarget-compile-run-and-check-amdgcn-amd-amdhsa
 
 // Clang 6.0 doesn't use the new map interface, undefined behavior when
 // the compiler emits "old" interface code for structures.
 // UNSUPPORTED: clang-6
 
 #include <stdio.h>
 #include <stdlib.h>
 

openmp/libomptarget/test/mapping/private_mapping.c

 // RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-nvptx64-nvidia-cuda
+// RUN: %libomptarget-compile-run-and-check-amdgcn-amd-amdhsa
 
 #include <assert.h>
 #include <stdio.h>
 
 int main() {
   int data1[3] = {1}, data2[3] = {2}, data3[3] = {3};
   int sum[16] = {0};
 #pragma omp target teams distribute parallel for map(tofrom                    \

openmp/libomptarget/test/offloading/assert.cpp

 // RUN: %libomptarget-compilexx-nvptx64-nvidia-cuda && %libomptarget-run-fail-nvptx64-nvidia-cuda
+// RUN: %libomptarget-compilexx-amdgcn-amd-amdhsa && %libomptarget-run-fail-amdgcn-amd-amdhsa
 
 int main(int argc, char *argv[]) {
 #pragma omp target
   { __builtin_trap(); }
 
   return 0;
 }

openmp/libomptarget/test/offloading/bug47654.cpp

 // RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
 // RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
+// RUN: %libomptarget-compilexx-run-and-check-amdgcn-amd-amdhsa
 
 #include <cassert>
 #include <iostream>
 
 int main(int argc, char *argv[]) {
   int i = 0, j = 0;
 
 #pragma omp target map(tofrom : i, j) nowait

openmp/libomptarget/test/offloading/d2d_memcpy.c

 // RUN: %libomptarget-compile-aarch64-unknown-linux-gnu && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-aarch64-unknown-linux-gnu | %fcheck-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-powerpc64-ibm-linux-gnu | %fcheck-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-powerpc64le-ibm-linux-gnu | %fcheck-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-compile-x86_64-pc-linux-gnu && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-x86_64-pc-linux-gnu | %fcheck-x86_64-pc-linux-gnu -allow-empty
 // RUN: %libomptarget-compile-nvptx64-nvidia-cuda && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-nvptx64-nvidia-cuda | %fcheck-nvptx64-nvidia-cuda -allow-empty
+// RUN: %libomptarget-compile-amdgcn-amd-amdhsa && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-amdgcn-amd-amdhsa | %fcheck-amdgcn-amd-amdhsa -allow-empty
 
 #include <assert.h>
 #include <omp.h>
 #include <stdio.h>
 #include <stdlib.h>
 
 const int magic_num = 7;
 

openmp/libomptarget/test/offloading/dynamic_module.c

 // RUN: %libomptarget-compile-aarch64-unknown-linux-gnu -DSHARED -fPIC -shared -o %t.so && %libomptarget-compile-aarch64-unknown-linux-gnu %t.so && %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 | %fcheck-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu -DSHARED -fPIC -shared -o %t.so && %libomptarget-compile-powerpc64-ibm-linux-gnu %t.so && %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 | %fcheck-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu -DSHARED -fPIC -shared -o %t.so && %libomptarget-compile-powerpc64le-ibm-linux-gnu %t.so && %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 | %fcheck-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-compile-x86_64-pc-linux-gnu -DSHARED -fPIC -shared -o %t.so && %libomptarget-compile-x86_64-pc-linux-gnu %t.so && %libomptarget-run-x86_64-pc-linux-gnu 2>&1 | %fcheck-x86_64-pc-linux-gnu
 // RUN: %libomptarget-compile-nvptx64-nvidia-cuda -DSHARED -fPIC -shared -o %t.so && %libomptarget-compile-nvptx64-nvidia-cuda %t.so && %libomptarget-run-nvptx64-nvidia-cuda 2>&1 | %fcheck-nvptx64-nvidia-cuda
+// RUN: %libomptarget-compile-amdgcn-amd-amdhsa -DSHARED -fPIC -shared -o %t.so && %libomptarget-compile-amdgcn-amd-amdhsa %t.so && %libomptarget-run-amdgcn-amd-amdhsa 2>&1 | %fcheck-amdgcn-amd-amdhsa
 
 #ifdef SHARED
 void foo() {}
 #else
 #include <stdio.h>
 int main() {
 #pragma omp target
   ;
   // CHECK: DONE.
   printf("%s\n", "DONE.");
   return 0;
 }
 #endif

openmp/libomptarget/test/offloading/dynamic_module_load.c

 // RUN: %libomptarget-compile-aarch64-unknown-linux-gnu -DSHARED -fPIC -shared -o %t.so && %clang %flags %s -o %t-aarch64-unknown-linux-gnu -ldl && %libomptarget-run-aarch64-unknown-linux-gnu %t.so 2>&1 | %fcheck-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu -DSHARED -fPIC -shared -o %t.so && %clang %flags %s -o %t-powerpc64-ibm-linux-gnu -ldl && %libomptarget-run-powerpc64-ibm-linux-gnu %t.so 2>&1 | %fcheck-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu -DSHARED -fPIC -shared -o %t.so && %clang %flags %s -o %t-powerpc64le-ibm-linux-gnu -ldl && %libomptarget-run-powerpc64le-ibm-linux-gnu %t.so 2>&1 | %fcheck-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-compile-x86_64-pc-linux-gnu -DSHARED -fPIC -shared -o %t.so && %clang %flags %s -o %t-x86_64-pc-linux-gnu -ldl && %libomptarget-run-x86_64-pc-linux-gnu %t.so 2>&1 | %fcheck-x86_64-pc-linux-gnu
 // RUN: %libomptarget-compile-nvptx64-nvidia-cuda -DSHARED -fPIC -shared -o %t.so && %clang %flags %s -o %t-nvptx64-nvidia-cuda -ldl && %libomptarget-run-nvptx64-nvidia-cuda %t.so 2>&1 | %fcheck-nvptx64-nvidia-cuda
+// RUN: %libomptarget-compile-amdgcn-amd-amdhsa -DSHARED -fPIC -shared -o %t.so && %clang %flags %s -o %t-amdgcn-amd-amdhsa -ldl && %libomptarget-run-amdgcn-amd-amdhsa %t.so 2>&1 | %fcheck-amdgcn-amd-amdhsa
 
 #ifdef SHARED
 #include <stdio.h>
 int foo() {
 #pragma omp target
   ;
   printf("%s\n", "DONE.");
   return 0;

openmp/libomptarget/test/offloading/lone_target_exit_data.c

 // Check that a target exit data directive behaves correctly when the runtime
 // has not yet been initialized.
 
 // RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-nvptx64-nvidia-cuda
+// RUN: %libomptarget-compile-run-and-check-amdgcn-amd-amdhsa
 
 #include <stdio.h>
 
 int main() {
   // CHECK: x = 98
   int x = 98;
   #pragma omp target exit data map(from:x)
   printf("x = %d\n", x);
   return 0;
 }

openmp/libomptarget/test/offloading/looptripcnt.c

 // RUN: %libomptarget-compile-aarch64-unknown-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 | %fcheck-aarch64-unknown-linux-gnu -allow-empty -check-prefix=DEBUG
 // RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 | %fcheck-powerpc64-ibm-linux-gnu -allow-empty -check-prefix=DEBUG
 // RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 | %fcheck-powerpc64le-ibm-linux-gnu -allow-empty -check-prefix=DEBUG
 // RUN: %libomptarget-compile-x86_64-pc-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-x86_64-pc-linux-gnu 2>&1 | %fcheck-x86_64-pc-linux-gnu -allow-empty -check-prefix=DEBUG
 // RUN: %libomptarget-compile-nvptx64-nvidia-cuda && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-nvptx64-nvidia-cuda 2>&1 | %fcheck-nvptx64-nvidia-cuda -allow-empty -check-prefix=DEBUG
+// RUN: %libomptarget-compile-amdgcn-amd-amdhsa && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-amdgcn-amd-amdhsa 2>&1 | %fcheck-amdgcn-amd-amdhsa -allow-empty -check-prefix=DEBUG
 // REQUIRES: libomptarget-debug
 
 /*
   Test for looptripcount being popped from runtime stack.
 */
 #include <stdio.h>
 #include <omp.h>
 int main()

openmp/libomptarget/test/offloading/non_contiguous_update.cpp

 // RUN: %libomptarget-compile-aarch64-unknown-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 | %fcheck-aarch64-unknown-linux-gnu -allow-empty -check-prefix=DEBUG
 // RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 | %fcheck-powerpc64-ibm-linux-gnu -allow-empty -check-prefix=DEBUG
 // RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 | %fcheck-powerpc64le-ibm-linux-gnu -allow-empty -check-prefix=DEBUG
 // RUN: %libomptarget-compile-x86_64-pc-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-x86_64-pc-linux-gnu 2>&1 | %fcheck-x86_64-pc-linux-gnu -allow-empty -check-prefix=DEBUG
 // RUN: %libomptarget-compile-nvptx64-nvidia-cuda && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-nvptx64-nvidia-cuda 2>&1 | %fcheck-nvptx64-nvidia-cuda -allow-empty -check-prefix=DEBUG
+// RUN: %libomptarget-compile-amdgcn-amd-amdhsa && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-amdgcn-amd-amdhsa 2>&1 | %fcheck-amdgcn-amd-amdhsa -allow-empty -check-prefix=DEBUG
 // REQUIRES: libomptarget-debug
 
 #include <cstdio>
 #include <cstdlib>
 #include <cassert>
 
 // Data structure definitions copied from OpenMP RTL.
 struct __tgt_target_non_contig {

openmp/libomptarget/test/offloading/offloading_success.c

 // RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-nvptx64-nvidia-cuda
+// RUN: %libomptarget-compile-run-and-check-amdgcn-amd-amdhsa
 
 #include <stdio.h>
 #include <omp.h>
 
 int main(void) {
   int isHost = -1;
 
 #pragma omp target map(from: isHost)

openmp/libomptarget/test/offloading/offloading_success.cpp

 // RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
 // RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
+// RUN: %libomptarget-compilexx-run-and-check-amdgcn-amd-amdhsa
 
 #include <stdio.h>
 #include <omp.h>
 
 int main(void) {
   int isHost = 0;
 
 #pragma omp target map(from: isHost)

openmp/libomptarget/test/offloading/requires.c

 // RUN: %libomptarget-compile-aarch64-unknown-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 | %fcheck-aarch64-unknown-linux-gnu -allow-empty -check-prefix=DEBUG
 // RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 | %fcheck-powerpc64-ibm-linux-gnu -allow-empty -check-prefix=DEBUG
 // RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 | %fcheck-powerpc64le-ibm-linux-gnu -allow-empty -check-prefix=DEBUG
 // RUN: %libomptarget-compile-x86_64-pc-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-x86_64-pc-linux-gnu 2>&1 | %fcheck-x86_64-pc-linux-gnu -allow-empty -check-prefix=DEBUG
 // RUN: %libomptarget-compile-nvptx64-nvidia-cuda && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-nvptx64-nvidia-cuda 2>&1 | %fcheck-nvptx64-nvidia-cuda -allow-empty -check-prefix=DEBUG
+// RUN: %libomptarget-compile-amdgcn-amd-amdhsa && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-amdgcn-amd-amdhsa 2>&1 | %fcheck-amdgcn-amd-amdhsa -allow-empty -check-prefix=DEBUG
 // REQUIRES: libomptarget-debug
 
 /*
   Test for the 'requires' clause check.
   When a target region is used, the requires flags are set in the
   runtime for the entire compilation unit. If the flags are set again,
   (for whatever reason) the set must be consistent with previously
   set values.

openmp/libomptarget/test/offloading/target_depend_nowait.cpp

 // RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
 // RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
+// RUN: %libomptarget-compilexx-run-and-check-amdgcn-amd-amdhsa
 
 #include <omp.h>
 #include <stdio.h>
 
 #define N 1024
 
 int A[N];
 int B[N];