From 890ffa8bb2f26e679a2bac3601556084061e0e0b Mon Sep 17 00:00:00 2001 From: Kewen Meng Date: Thu, 2 Jan 2025 19:33:14 -0600 Subject: [PATCH 1/2] [Test] Add smoke test to cover new runtime envar This smoke test is to cover the new runtime envar OMPX_AMD_MEMORY_MANAGER_THRESHOLD_EXP_2 which is designed for users to specify the threshold for the memory size handled by AMDGPUMemoryManager. --- .../AMDGPUMemoryManager_threshold_env.c | 38 +++++++++++++++++++ .../rt-memory-threshold-env/Makefile | 18 +++++++++ 2 files changed, 56 insertions(+) create mode 100644 test/smoke-dev/rt-memory-threshold-env/AMDGPUMemoryManager_threshold_env.c create mode 100644 test/smoke-dev/rt-memory-threshold-env/Makefile diff --git a/test/smoke-dev/rt-memory-threshold-env/AMDGPUMemoryManager_threshold_env.c b/test/smoke-dev/rt-memory-threshold-env/AMDGPUMemoryManager_threshold_env.c new file mode 100644 index 000000000..dc37754f6 --- /dev/null +++ b/test/smoke-dev/rt-memory-threshold-env/AMDGPUMemoryManager_threshold_env.c @@ -0,0 +1,38 @@ +#include +#include + +int main() { + int N = 10; + + int a[N]; + int b[N]; + + int i; + + for (i = 0; i < N; i++) + a[i] = 0; + + for (i = 0; i < N; i++) + b[i] = i; + +#pragma omp target parallel for + { + for (int j = 0; j < N; j++) + a[j] = b[j]; + } + + int rc = 0; + for (i = 0; i < N; i++) + if (a[i] != b[i]) { + rc++; + printf("Wrong varlue: a[%d]=%d\n", i, a[i]); + } + + if (!rc) + printf("Success\n"); + + return rc; +} + +/// CHECK: TARGET AMDGPU RTL --> AMDGPUMemoryManager threshhold was set to: +/// 1048576 B diff --git a/test/smoke-dev/rt-memory-threshold-env/Makefile b/test/smoke-dev/rt-memory-threshold-env/Makefile new file mode 100644 index 000000000..88ed2c632 --- /dev/null +++ b/test/smoke-dev/rt-memory-threshold-env/Makefile @@ -0,0 +1,18 @@ +include ../../Makefile.defs + +TESTNAME = AMDGPUMemoryManager_threshold_env +TESTSRC_MAIN = AMDGPUMemoryManager_threshold_env.c +TESTSRC_AUX = +TESTSRC_ALL = $(TESTSRC_MAIN) $(TESTSRC_AUX) +RUNENV += LIBOMPTARGET_DEBUG=1 +RUNENV += OMPX_AMD_MEMORY_MANAGER_THRESHOLD_EXP_2=20 + +RUNCMD = ./$(TESTNAME) 2>&1 | $(FILECHECK) $(TESTSRC_MAIN) + +CLANG ?= clang +OMP_BIN = $(AOMP)/bin/$(CLANG) +CC = $(OMP_BIN) $(VERBOSE) +#-ccc-print-phases +#"-\#\#\#" + +include ../Makefile.rules From 73395a2f923969dec19315b7590d90e9e753654c Mon Sep 17 00:00:00 2001 From: Kewen12 Date: Mon, 6 Jan 2025 13:15:18 -0800 Subject: [PATCH 2/2] Fix the typo. Co-authored-by: Jan Patrick Lehr --- .../rt-memory-threshold-env/AMDGPUMemoryManager_threshold_env.c | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/test/smoke-dev/rt-memory-threshold-env/AMDGPUMemoryManager_threshold_env.c b/test/smoke-dev/rt-memory-threshold-env/AMDGPUMemoryManager_threshold_env.c index dc37754f6..044d8e72a 100644 --- a/test/smoke-dev/rt-memory-threshold-env/AMDGPUMemoryManager_threshold_env.c +++ b/test/smoke-dev/rt-memory-threshold-env/AMDGPUMemoryManager_threshold_env.c @@ -25,7 +25,7 @@ int main() { for (i = 0; i < N; i++) if (a[i] != b[i]) { rc++; - printf("Wrong varlue: a[%d]=%d\n", i, a[i]); + printf("Wrong value: a[%d]=%d\n", i, a[i]); } if (!rc)