[OpenMP] Enable optimization level for memory_manager.cpp OpenMP test.

[doru1004]

Summary

Enable -O3 opt level for this test.

[JonChesterfield]

Please add a reason why to the commit message. Is O3 different to O2 in terms of behaviour?

Are we defaulting to O0 for all the runtime tests on amdgpu? That seems unlikely as O0 usually doesn't work, but equally might explain why some tests are disabled.

[tianshilei1992]

I'm not fully sold by it as well. Why do we need O3 for this?

[doru1004]

In D159119#4625678, @JonChesterfield wrote:

Please add a reason why to the commit message. Is O3 different to O2 in terms of behaviour?

Are we defaulting to O0 for all the runtime tests on amdgpu? That seems unlikely as O0 usually doesn't work, but equally might explain why some tests are disabled.

The initial test doesn't have an optimization level specified. Does it? If it had O2 it would be fine, it would be passing. Please check this issue here: https://github.com/llvm/llvm-project/issues/65077

The rationale is this test is currently disabled but passes with O2 or O3 explicitly passed in. The test fails to compile with O0 and fails to run consistently with O1 or no opt level specified.

[jdoerfert]

openmp/libomptarget/test/offloading/memory_manager.cpp

1

we have something like this.

[jplehr]

I think we can accept this and monitor whether it is still flaky with the opt level given. If required, we can re-disable the test.

As discussed in the linked issue, there seem to be different problems here.

[doru1004]

In D159119#4627860, @jplehr wrote:

I think we can accept this and monitor whether it is still flaky with the opt level given. If required, we can re-disable the test.

As discussed in the linked issue, there seem to be different problems here.

Thanks JP. Everyone ok with this plan?

[protze.joachim]

The test actually has an OpenMP issue:

The memory is allocated on device 0, but no device is specified for the target region. I think, all target constructs should explicitly have a device(0) clause. At the moment, the test relies on implementation-defined behavior.

[doru1004]

In D159119#4630718, @protze.joachim wrote:

The test actually has an OpenMP issue:

The memory is allocated on device 0, but no device is specified for the target region. I think, all target constructs should explicitly have a device(0) clause. At the moment, the test relies on implementation-defined behavior.

Added the additional device(0) clauses. This fixed the -O1 fails, I could no longer see intermittent fails with -O1 and with no optimization level the intermittent fails were significantly reduced to 1-2 in 500 tries whereas before there were 4-5 in every 100 tries. I will update the issue.

[tianshilei1992]

In D159119#4630718, @protze.joachim wrote:

The test actually has an OpenMP issue:

The memory is allocated on device 0, but no device is specified for the target region. I think, all target constructs should explicitly have a device(0) clause. At the moment, the test relies on implementation-defined behavior.

I don’t think that is the cause of the issue. The default device is 0, at least for OpenMP.

[protze.joachim]

In D159119#4631583, @tianshilei1992 wrote:

I don’t think that is the cause of the issue. The default device is 0, at least for OpenMP.

The default device is not 0. At least not according to the OpenMP spec. Check OpenMP-5.2, page 41, lines 10-13:

If target-offload-var is mandatory and the number of non-host devices is zero then the default-device-var is initialized to omp_invalid_device. Otherwise, the initial value is an implementation-defined non-negative integer that is less than or, if target-offload-var is not mandatory, equal to omp_get_initial_device().

I didn't claim, that this caused the issue. I just noted that according to the OpenMP spec, the test was sailing in implementation-defined waters.

Thanks Doru for fixing the test :D

[tianshilei1992]

In D159119#4631612, @protze.joachim wrote:

In D159119#4631583, @tianshilei1992 wrote:

I don’t think that is the cause of the issue. The default device is 0, at least for OpenMP.

The default device is not 0. At least not according to the OpenMP spec. Check OpenMP-5.2, page 41, lines 10-13:

If target-offload-var is mandatory and the number of non-host devices is zero then the default-device-var is initialized to omp_invalid_device. Otherwise, the initial value is an implementation-defined non-negative integer that is less than or, if target-offload-var is not mandatory, equal to omp_get_initial_device().

I didn't claim, that this caused the issue. I just noted that according to the OpenMP spec, the test was sailing in implementation-defined waters.

Thanks Doru for fixing the test :D

I meant to say "for LLVM OpenMP" but somehow missed the "LLVM". ;-)

On paper you are right. In practice, for LLVM OpenMP, omp_get_default_device returns 0 *by default*. This still holds for the case we are discussing here because:

• Our test config doesn't change the value via env.
• For each test we require the corresponding "test feature", and also only compile for the corresponding target.

All of this can guarantee that device 0 used for API corresponds to the default device for target region. I think that is also the reason that almost all libomptarget tests don't explicitly specify device number on the target construct.

[doru1004]

On paper you are right.

I added them anyway just to be sure :)

openmp/libomptarget/test/offloading/memory_manager.cpp

-// RUN: %libomptarget-compilexx-run-and-check-generic
+// RUN: %libomptarget-compileoptxx-run-and-check-generic
 
-// UNSUPPORTED: amdgcn-amd-amdhsa
 // UNSUPPORTED: x86_64-pc-linux-gnu
 // UNSUPPORTED: x86_64-pc-linux-gnu-LTO
 
 #include <omp.h>
 
 #include <cassert>
 #include <iostream>
 
 int main(int argc, char *argv[]) {
 #pragma omp parallel for
   for (int i = 0; i < 16; ++i) {
     for (int n = 1; n < (1 << 13); n <<= 1) {
       void *p = omp_target_alloc(n * sizeof(int), 0);
       omp_target_free(p, 0);
     }
   }
 
 #pragma omp parallel for
   for (int i = 0; i < 16; ++i) {
     for (int n = 1; n < (1 << 13); n <<= 1) {
       int *p = (int *)omp_target_alloc(n * sizeof(int), 0);
-#pragma omp target teams distribute parallel for is_device_ptr(p)
+#pragma omp target teams distribute parallel for is_device_ptr(p) device(0)
       for (int j = 0; j < n; ++j) {
         p[j] = i;
       }
       int buffer[n];
 #pragma omp target teams distribute parallel for is_device_ptr(p)              \
-    map(from : buffer)
+    map(from : buffer) device(0)
       for (int j = 0; j < n; ++j) {
         buffer[j] = p[j];
       }
       for (int j = 0; j < n; ++j) {
         assert(buffer[j] == i);
       }
       omp_target_free(p, 0);
     }
   }
 
   std::cout << "PASS\n";
   return 0;
 }
 
 // CHECK: PASS