[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 ↗

(On Diff #334417)

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 ↗

(On Diff #334417)

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 ↗

(On Diff #334417)

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

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

19 ↗

(On Diff #334417)

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 ↗

(On Diff #334417)

I will create a separate patch for it.

openmp/libomptarget/plugins/CMakeLists.txt

80 ↗

(On Diff #334417)

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 ↗

(On Diff #334417)

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

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

 
 # 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()
 
 # 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 "")

openmp/libomptarget/test/mapping/data_member_ref.cpp

 // RUN: %libomptarget-compilexx-run-and-check-generic
 
+// amdgcn does not have printf definition
+// XFAIL: amdgcn-amd-amdhsa
+
 #include <stdio.h>
 
 struct View {
   int Data;
 };
 
 struct ViewPtr {
   int *Data;

openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers.cpp

 // RUN: %libomptarget-compilexx-run-and-check-generic
 
+// amdgcn does not have printf definition
+// XFAIL: amdgcn-amd-amdhsa
+
 #include <cstdio>
 #include <cstdlib>
 
 typedef struct {
   int a;
   double *b;
 } C1;
 #pragma omp declare mapper(C1 s) map(to : s.a) map(from : s.b [0:2])

openmp/libomptarget/test/mapping/declare_mapper_nested_mappers.cpp

 // RUN: %libomptarget-compilexx-run-and-check-generic
 
+// amdgcn does not have printf definition
+// XFAIL: amdgcn-amd-amdhsa
+
 #include <cstdio>
 #include <cstdlib>
 
 typedef struct {
   int a;
   double *b;
 } C;
 #pragma omp declare mapper(id1 : C s) map(to : s.a) map(from : s.b [0:2])

openmp/libomptarget/test/mapping/delete_inf_refcount.c

 // RUN: %libomptarget-compile-run-and-check-generic
 
+// fails with error message 'Unable to generate target entries' on amdgcn
+// XFAIL: amdgcn-amd-amdhsa
+
 #include <stdio.h>
 #include <omp.h>
 
 #pragma omp declare target
 int isHost;
 #pragma omp end declare target
 
 int main(void) {

openmp/libomptarget/test/mapping/ptr_and_obj_motion.c

 // RUN: %libomptarget-compile-run-and-check-generic
 
+// amdgcn does not have printf definition
+// XFAIL: amdgcn-amd-amdhsa
+
 #include <stdio.h>
 
 typedef struct {
   double *dataptr;
   int dummy1;
   int dummy2;
 } DV;
 

openmp/libomptarget/test/offloading/host_as_target.c

 // Check that specifying device as omp_get_initial_device():
 // - Doesn't cause the runtime to fail.
 // - Offloads code to the host.
 // - Doesn't transfer data.  In this case, just check that neither host data nor
 //   default device data are affected by the specified transfers.
 // - Works whether it's specified directly or as the default device.
 
 // RUN: %libomptarget-compile-run-and-check-generic
 
+// amdgcn does not have printf definition
+// XFAIL: amdgcn-amd-amdhsa
+
 #include <stdio.h>
 #include <omp.h>
 
 static void check(char *X, int Dev) {
   printf("  host X = %c\n", *X);
   #pragma omp target device(Dev)
   printf("device X = %c\n", *X);
 }

openmp/libomptarget/test/unified_shared_memory/api.c

 // RUN: %libomptarget-compile-run-and-check-generic
 // XFAIL: nvptx64-nvidia-cuda
+// XFAIL: amdgcn-amd-amdhsa
 
 #include <stdio.h>
 #include <omp.h>
 
 // ---------------------------------------------------------------------------
 // Various definitions copied from OpenMP RTL
 
 extern void __tgt_register_requires(int64_t);

openmp/libomptarget/test/unified_shared_memory/close_enter_exit.c

 // RUN: %libomptarget-compile-run-and-check-generic
 
 // UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
-// XFAIL: nvptx64-nvidia-cuda
+// XFAIL: nvptx64-nvidia-cuda, amdgcn-amd-amdhsa
 
 #include <omp.h>
 #include <stdio.h>
 
 #pragma omp requires unified_shared_memory
 
 #define N 1024
 

openmp/libomptarget/test/unified_shared_memory/close_modifier.c

 // RUN: %libomptarget-compile-run-and-check-generic
-// XFAIL: nvptx64-nvidia-cuda
+// XFAIL: nvptx64-nvidia-cuda, amdgcn-amd-amdhsa
 
 // UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
 
 #include <omp.h>
 #include <stdio.h>
 
 #pragma omp requires unified_shared_memory
 

openmp/libomptarget/test/unified_shared_memory/shared_update.c

 // RUN: %libomptarget-compile-run-and-check-generic
-// XFAIL: nvptx64-nvidia-cuda
+// XFAIL: nvptx64-nvidia-cuda, amdgcn-amd-amdhsa
 
 #include <stdio.h>
 #include <omp.h>
 
 // ---------------------------------------------------------------------------
 // Various definitions copied from OpenMP RTL
 
 extern void __tgt_register_requires(int64_t);