diff --git a/openmp/libomptarget/CMakeLists.txt b/openmp/libomptarget/CMakeLists.txt
--- a/openmp/libomptarget/CMakeLists.txt
+++ b/openmp/libomptarget/CMakeLists.txt
@@ -35,8 +35,13 @@
   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")
@@ -73,7 +78,6 @@
 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)
diff --git a/openmp/libomptarget/deviceRTLs/CMakeLists.txt b/openmp/libomptarget/deviceRTLs/CMakeLists.txt
--- a/openmp/libomptarget/deviceRTLs/CMakeLists.txt
+++ b/openmp/libomptarget/deviceRTLs/CMakeLists.txt
@@ -9,5 +9,7 @@
 # Build a device RTL for each available machine.
 #
 ##===----------------------------------------------------------------------===##
-
+if(LIBOMPTARGET_ENABLE_EXPERIMENTAL_AMDGCN_TARGET)
+    add_subdirectory(amdgcn)
+endif()
 add_subdirectory(nvptx)
diff --git a/openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt b/openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt
--- a/openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt
+++ b/openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt
@@ -91,7 +91,7 @@
 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()
diff --git a/openmp/libomptarget/plugins/CMakeLists.txt b/openmp/libomptarget/plugins/CMakeLists.txt
--- a/openmp/libomptarget/plugins/CMakeLists.txt
+++ b/openmp/libomptarget/plugins/CMakeLists.txt
@@ -70,7 +70,6 @@
 endmacro()
 
 add_subdirectory(aarch64)
-add_subdirectory(amdgpu)
 add_subdirectory(cuda)
 add_subdirectory(ppc64)
 add_subdirectory(ppc64le)
@@ -78,6 +77,11 @@
 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)
diff --git a/openmp/libomptarget/test/env/base_ptr_ref_count.c b/openmp/libomptarget/test/env/base_ptr_ref_count.c
--- a/openmp/libomptarget/test/env/base_ptr_ref_count.c
+++ b/openmp/libomptarget/test/env/base_ptr_ref_count.c
@@ -3,6 +3,8 @@
 // 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>
diff --git a/openmp/libomptarget/test/env/omp_target_debug.c b/openmp/libomptarget/test/env/omp_target_debug.c
--- a/openmp/libomptarget/test/env/omp_target_debug.c
+++ b/openmp/libomptarget/test/env/omp_target_debug.c
@@ -8,6 +8,8 @@
 // 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) {
diff --git a/openmp/libomptarget/test/lit.cfg b/openmp/libomptarget/test/lit.cfg
--- a/openmp/libomptarget/test/lit.cfg
+++ b/openmp/libomptarget/test/lit.cfg
@@ -84,6 +84,11 @@
 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, \
@@ -116,9 +121,11 @@
             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:
diff --git a/openmp/libomptarget/test/lit.site.cfg.in b/openmp/libomptarget/test/lit.site.cfg.in
--- a/openmp/libomptarget/test/lit.site.cfg.in
+++ b/openmp/libomptarget/test/lit.site.cfg.in
@@ -16,5 +16,8 @@
 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")
diff --git a/openmp/libomptarget/test/mapping/data_absent_at_exit.c b/openmp/libomptarget/test/mapping/data_absent_at_exit.c
--- a/openmp/libomptarget/test/mapping/data_absent_at_exit.c
+++ b/openmp/libomptarget/test/mapping/data_absent_at_exit.c
@@ -3,6 +3,7 @@
 // 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>
 
diff --git a/openmp/libomptarget/test/mapping/declare_mapper_api.cpp b/openmp/libomptarget/test/mapping/declare_mapper_api.cpp
--- a/openmp/libomptarget/test/mapping/declare_mapper_api.cpp
+++ b/openmp/libomptarget/test/mapping/declare_mapper_api.cpp
@@ -3,6 +3,7 @@
 // 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>
diff --git a/openmp/libomptarget/test/mapping/declare_mapper_target.cpp b/openmp/libomptarget/test/mapping/declare_mapper_target.cpp
--- a/openmp/libomptarget/test/mapping/declare_mapper_target.cpp
+++ b/openmp/libomptarget/test/mapping/declare_mapper_target.cpp
@@ -3,6 +3,7 @@
 // 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>
diff --git a/openmp/libomptarget/test/mapping/declare_mapper_target_data.cpp b/openmp/libomptarget/test/mapping/declare_mapper_target_data.cpp
--- a/openmp/libomptarget/test/mapping/declare_mapper_target_data.cpp
+++ b/openmp/libomptarget/test/mapping/declare_mapper_target_data.cpp
@@ -3,6 +3,7 @@
 // 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>
diff --git a/openmp/libomptarget/test/mapping/declare_mapper_target_data_enter_exit.cpp b/openmp/libomptarget/test/mapping/declare_mapper_target_data_enter_exit.cpp
--- a/openmp/libomptarget/test/mapping/declare_mapper_target_data_enter_exit.cpp
+++ b/openmp/libomptarget/test/mapping/declare_mapper_target_data_enter_exit.cpp
@@ -3,6 +3,7 @@
 // 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>
diff --git a/openmp/libomptarget/test/mapping/declare_mapper_target_update.cpp b/openmp/libomptarget/test/mapping/declare_mapper_target_update.cpp
--- a/openmp/libomptarget/test/mapping/declare_mapper_target_update.cpp
+++ b/openmp/libomptarget/test/mapping/declare_mapper_target_update.cpp
@@ -3,6 +3,7 @@
 // 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>
diff --git a/openmp/libomptarget/test/mapping/lambda_mapping.cpp b/openmp/libomptarget/test/mapping/lambda_mapping.cpp
--- a/openmp/libomptarget/test/mapping/lambda_mapping.cpp
+++ b/openmp/libomptarget/test/mapping/lambda_mapping.cpp
@@ -3,6 +3,7 @@
 // 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>
 
diff --git a/openmp/libomptarget/test/mapping/pr38704.c b/openmp/libomptarget/test/mapping/pr38704.c
--- a/openmp/libomptarget/test/mapping/pr38704.c
+++ b/openmp/libomptarget/test/mapping/pr38704.c
@@ -3,6 +3,7 @@
 // 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.
diff --git a/openmp/libomptarget/test/mapping/private_mapping.c b/openmp/libomptarget/test/mapping/private_mapping.c
--- a/openmp/libomptarget/test/mapping/private_mapping.c
+++ b/openmp/libomptarget/test/mapping/private_mapping.c
@@ -3,6 +3,7 @@
 // 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>
diff --git a/openmp/libomptarget/test/offloading/assert.cpp b/openmp/libomptarget/test/offloading/assert.cpp
--- a/openmp/libomptarget/test/offloading/assert.cpp
+++ b/openmp/libomptarget/test/offloading/assert.cpp
@@ -1,4 +1,5 @@
 // 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
diff --git a/openmp/libomptarget/test/offloading/bug47654.cpp b/openmp/libomptarget/test/offloading/bug47654.cpp
--- a/openmp/libomptarget/test/offloading/bug47654.cpp
+++ b/openmp/libomptarget/test/offloading/bug47654.cpp
@@ -3,6 +3,7 @@
 // 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>
diff --git a/openmp/libomptarget/test/offloading/d2d_memcpy.c b/openmp/libomptarget/test/offloading/d2d_memcpy.c
--- a/openmp/libomptarget/test/offloading/d2d_memcpy.c
+++ b/openmp/libomptarget/test/offloading/d2d_memcpy.c
@@ -3,6 +3,7 @@
 // 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>
diff --git a/openmp/libomptarget/test/offloading/dynamic_module.c b/openmp/libomptarget/test/offloading/dynamic_module.c
--- a/openmp/libomptarget/test/offloading/dynamic_module.c
+++ b/openmp/libomptarget/test/offloading/dynamic_module.c
@@ -3,6 +3,7 @@
 // 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() {}
diff --git a/openmp/libomptarget/test/offloading/dynamic_module_load.c b/openmp/libomptarget/test/offloading/dynamic_module_load.c
--- a/openmp/libomptarget/test/offloading/dynamic_module_load.c
+++ b/openmp/libomptarget/test/offloading/dynamic_module_load.c
@@ -3,6 +3,7 @@
 // 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>
diff --git a/openmp/libomptarget/test/offloading/lone_target_exit_data.c b/openmp/libomptarget/test/offloading/lone_target_exit_data.c
--- a/openmp/libomptarget/test/offloading/lone_target_exit_data.c
+++ b/openmp/libomptarget/test/offloading/lone_target_exit_data.c
@@ -6,6 +6,7 @@
 // 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>
 
diff --git a/openmp/libomptarget/test/offloading/looptripcnt.c b/openmp/libomptarget/test/offloading/looptripcnt.c
--- a/openmp/libomptarget/test/offloading/looptripcnt.c
+++ b/openmp/libomptarget/test/offloading/looptripcnt.c
@@ -3,6 +3,7 @@
 // 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
 
 /*
diff --git a/openmp/libomptarget/test/offloading/non_contiguous_update.cpp b/openmp/libomptarget/test/offloading/non_contiguous_update.cpp
--- a/openmp/libomptarget/test/offloading/non_contiguous_update.cpp
+++ b/openmp/libomptarget/test/offloading/non_contiguous_update.cpp
@@ -3,6 +3,7 @@
 // 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>
diff --git a/openmp/libomptarget/test/offloading/offloading_success.c b/openmp/libomptarget/test/offloading/offloading_success.c
--- a/openmp/libomptarget/test/offloading/offloading_success.c
+++ b/openmp/libomptarget/test/offloading/offloading_success.c
@@ -3,6 +3,7 @@
 // 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>
diff --git a/openmp/libomptarget/test/offloading/offloading_success.cpp b/openmp/libomptarget/test/offloading/offloading_success.cpp
--- a/openmp/libomptarget/test/offloading/offloading_success.cpp
+++ b/openmp/libomptarget/test/offloading/offloading_success.cpp
@@ -3,6 +3,7 @@
 // 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>
diff --git a/openmp/libomptarget/test/offloading/requires.c b/openmp/libomptarget/test/offloading/requires.c
--- a/openmp/libomptarget/test/offloading/requires.c
+++ b/openmp/libomptarget/test/offloading/requires.c
@@ -3,6 +3,7 @@
 // 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
 
 /*
diff --git a/openmp/libomptarget/test/offloading/target_depend_nowait.cpp b/openmp/libomptarget/test/offloading/target_depend_nowait.cpp
--- a/openmp/libomptarget/test/offloading/target_depend_nowait.cpp
+++ b/openmp/libomptarget/test/offloading/target_depend_nowait.cpp
@@ -3,6 +3,7 @@
 // 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>