diff options
author | Gheorghe-Teodor Bercea <doru.bercea@amd.com> | 2024-02-19 11:23:53 -0500 |
---|---|---|
committer | GitHub <noreply@github.com> | 2024-02-19 11:23:53 -0500 |
commit | 12ade6fc606c269101b9b13e52093c868e1a4924 (patch) | |
tree | a356b93aa999904b5261333c970f5971095e20ce /openmp | |
parent | bb56f052a37d9fa0a87c9636c1a94b85b72a52e0 (diff) | |
download | llvm-12ade6fc606c269101b9b13e52093c868e1a4924.zip llvm-12ade6fc606c269101b9b13e52093c868e1a4924.tar.gz llvm-12ade6fc606c269101b9b13e52093c868e1a4924.tar.bz2 |
[libomptarget][test] Add support for APU testing feature. (#82054)
Add test and support for `// REQUIRES: apu` for the category of tests
which exercise APU specific behavior.
Note: when running on an actual APU you may have to use the following if
the architecture ID is not enough to determine if the underlying device
is an APU:
```
IS_APU=1 ninja check-openmp
```
Diffstat (limited to 'openmp')
-rw-r--r-- | openmp/libomptarget/test/lit.cfg | 13 | ||||
-rw-r--r-- | openmp/libomptarget/test/mapping/auto_zero_copy_apu.cpp | 57 |
2 files changed, 70 insertions, 0 deletions
diff --git a/openmp/libomptarget/test/lit.cfg b/openmp/libomptarget/test/lit.cfg index fc1d436..565556e 100644 --- a/openmp/libomptarget/test/lit.cfg +++ b/openmp/libomptarget/test/lit.cfg @@ -34,6 +34,11 @@ if 'OMP_TARGET_OFFLOAD' in os.environ: if 'HSA_ENABLE_SDMA' in os.environ: config.environment['HSA_ENABLE_SDMA'] = os.environ['HSA_ENABLE_SDMA'] +# Architectures like gfx942 may or may not be APUs so an additional environment +# variable is required as some tests can be APU specific. +if 'IS_APU' in os.environ: + config.environment['IS_APU'] = os.environ['IS_APU'] + # set default environment variables for test if 'CHECK_OPENMP_ENV' in os.environ: test_env = os.environ['CHECK_OPENMP_ENV'].split() @@ -111,6 +116,7 @@ if config.libomptarget_has_libc: # For CUDA, this is the case with compute capability 70 (Volta) or higher. # For all other targets, we currently assume it is. supports_unified_shared_memory = True +supports_apu = False if config.libomptarget_current_target.startswith('nvptx'): try: cuda_arch = int(config.cuda_test_arch[:3]) @@ -126,8 +132,15 @@ elif config.libomptarget_current_target.startswith('amdgcn'): config.amdgpu_test_arch.startswith("gfx940") or config.amdgpu_test_arch.startswith("gfx942")): supports_unified_shared_memory = False + # check if AMD architecture is an APU: + if (config.amdgpu_test_arch.startswith("gfx940") or + (config.amdgpu_test_arch.startswith("gfx942") and + evaluate_bool_env(config.environment['IS_APU']))): + supports_apu = True if supports_unified_shared_memory: config.available_features.add('unified_shared_memory') +if supports_apu: + config.available_features.add('apu') # Setup environment to find dynamic library at runtime if config.operating_system == 'Windows': diff --git a/openmp/libomptarget/test/mapping/auto_zero_copy_apu.cpp b/openmp/libomptarget/test/mapping/auto_zero_copy_apu.cpp new file mode 100644 index 0000000..48360e4 --- /dev/null +++ b/openmp/libomptarget/test/mapping/auto_zero_copy_apu.cpp @@ -0,0 +1,57 @@ +// clang-format off +// RUN: %libomptarget-compilexx-generic +// RUN: env HSA_XNACK=1 LIBOMPTARGET_INFO=30 %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic -check-prefix=INFO_ZERO -check-prefix=CHECK + +// RUN: %libomptarget-compilexx-generic +// RUN: env HSA_XNACK=0 LIBOMPTARGET_INFO=30 %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic -check-prefix=INFO_COPY -check-prefix=CHECK + +// UNSUPPORTED: aarch64-unknown-linux-gnu +// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO +// UNSUPPORTED: nvptx64-nvidia-cuda +// UNSUPPORTED: nvptx64-nvidia-cuda-LTO +// UNSUPPORTED: x86_64-pc-linux-gnu +// UNSUPPORTED: x86_64-pc-linux-gnu-LTO + +// REQUIRES: apu + +// clang-format on + +#include <cstdio> + +int main() { + int n = 1024; + + // test various mapping types + int *a = new int[n]; + int k = 3; + int b[n]; + + for (int i = 0; i < n; i++) + b[i] = i; + + // clang-format off + // INFO_ZERO: Return HstPtrBegin 0x{{.*}} Size=4096 for unified shared memory + // INFO_ZERO: Return HstPtrBegin 0x{{.*}} Size=4096 for unified shared memory + + // INFO_COPY: Creating new map entry with HstPtrBase=0x{{.*}}, HstPtrBegin=0x{{.*}}, TgtAllocBegin=0x{{.*}}, TgtPtrBegin=0x{{.*}}, Size=4096, + // INFO_COPY: Creating new map entry with HstPtrBase=0x{{.*}}, HstPtrBegin=0x{{.*}}, TgtAllocBegin=0x{{.*}}, TgtPtrBegin=0x{{.*}}, Size=4096, + // INFO_COPY: Mapping exists with HstPtrBegin=0x{{.*}}, TgtPtrBegin=0x{{.*}}, Size=4096, DynRefCount=1 (update suppressed) + // INFO_COPY: Mapping exists with HstPtrBegin=0x{{.*}}, TgtPtrBegin=0x{{.*}}, Size=4096, DynRefCount=1 (update suppressed) +// clang-format on +#pragma omp target teams distribute parallel for map(tofrom : a[ : n]) \ + map(to : b[ : n]) + for (int i = 0; i < n; i++) + a[i] = i + b[i] + k; + + int err = 0; + for (int i = 0; i < n; i++) + if (a[i] != i + b[i] + k) + err++; + + // CHECK: PASS + if (err == 0) + printf("PASS\n"); + return err; +} |