From 9956102ff29892a80786ba46ff83f24658b2c918 Mon Sep 17 00:00:00 2001 From: Nicole Aschenbrenner Date: Thu, 23 Apr 2026 03:05:36 -0500 Subject: [PATCH 1/5] [smoke-dev] Remove multi-device test Remove xteam-red-min-max-float-multi-device test which depends on -fopenmp-target-multi-device flag. This flag is being removed from ROCm/llvm-project (PR #2277) as part of driver convergence effort. Test purpose was to verify multi-device compilation compatibility with Xteam min/max reduction. With multi-device feature removed, this test becomes redundant. Depends on: ROCm/llvm-project#2277 --- .../Makefile | 19 ------- .../xteam-red-min-max-float-multi-device.c | 49 ------------------- 2 files changed, 68 deletions(-) delete mode 100644 test/smoke-dev/xteam-red-min-max-float-multi-device/Makefile delete mode 100644 test/smoke-dev/xteam-red-min-max-float-multi-device/xteam-red-min-max-float-multi-device.c diff --git a/test/smoke-dev/xteam-red-min-max-float-multi-device/Makefile b/test/smoke-dev/xteam-red-min-max-float-multi-device/Makefile deleted file mode 100644 index d90453c20c..0000000000 --- a/test/smoke-dev/xteam-red-min-max-float-multi-device/Makefile +++ /dev/null @@ -1,19 +0,0 @@ -include ../../Makefile.defs - -TESTNAME = xteam-red-min-max-float-multi-device -TESTSRC_MAIN = xteam-red-min-max-float-multi-device.c -TESTSRC_AUX = -TESTSRC_ALL = $(TESTSRC_MAIN) $(TESTSRC_AUX) -RUNENV += HSA_XNACK=1 OMPX_APU_MAPS=1 LIBOMPTARGET_NUM_MULTI_DEVICES=1 LIBOMPTARGET_KERNEL_TRACE=1 - -RUNCMD = ./$(TESTNAME) 2>&1 | $(FILECHECK) $(TESTSRC_MAIN) - -CFLAGS += -fopenmp-target-multi-device -EXTRA_LDFLAGS += -lm -CLANG = clang -OMP_BIN = $(AOMP)/bin/$(CLANG) -CC = $(OMP_BIN) $(VERBOSE) -#-ccc-print-phases -#"-\#\#\#" - -include ../Makefile.rules diff --git a/test/smoke-dev/xteam-red-min-max-float-multi-device/xteam-red-min-max-float-multi-device.c b/test/smoke-dev/xteam-red-min-max-float-multi-device/xteam-red-min-max-float-multi-device.c deleted file mode 100644 index eea219a04b..0000000000 --- a/test/smoke-dev/xteam-red-min-max-float-multi-device/xteam-red-min-max-float-multi-device.c +++ /dev/null @@ -1,49 +0,0 @@ -/* - * Test multi-device min/max reduction on floats using minf/maxf. - * There are 2 target regions in this program, the first has min/max reduction - * and the second has a sum reduction. The program is compiled with multi-device - * ON. Since multi-device compilation may be incompatible with Xteam min/max, the - * first target region does not use Xteam reduction. The second one, however, does. - */ -#include -#include -#include -#include - -int main() -{ - int N = 1000; - - float a[N]; - - for (int i = 0; i < N; i++) - a[i] = i + 11; - - float max1 = 0; - float min1 = 1000000; - float sum1 = 0; - -#pragma omp target teams distribute parallel for reduction(max : max1) reduction(min : min1) - for (int i = 0; i < N; i = i + 1) - { - max1 = fmaxf(max1, a[i]); - min1 = fminf(min1, a[i]); - } - - #pragma omp target teams distribute parallel for reduction(+ : sum1) - for (int i = 0; i < N; i = i + 1) - sum1 += a[i]; - - printf("max1 = %f min1 = %f sum1 = %f\n", max1, min1, sum1); - int rc = (max1 != 1010) || (min1 != 11) || (sum1 != 510500); - - if (!rc) - printf("Success\n"); - else - printf("Failed\n"); - - return rc; -} - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 \ No newline at end of file From 025163fc74ed60c5b14682979bbd1442e6ca4b48 Mon Sep 17 00:00:00 2001 From: Nicole Aschenbrenner Date: Thu, 28 May 2026 04:47:09 -0500 Subject: [PATCH 2/5] [smoke] Remove xteam-scan tests These tests exercise -fopenmp-target-xteam-scan and -fopenmp-target-xteam-no-loop-scan, which are deprecated in ROCm/llvm-project#2704. The flags and their implementation will be removed once the upstream cross-team reduction rewrite lands in amd-staging. --- test/smoke-dev/xteam-scan-1/Makefile | 18 -- test/smoke-dev/xteam-scan-1/xteam_scan_1.cpp | 199 ----------------- test/smoke-dev/xteam-scan-no-loop/Makefile | 18 -- .../xteam-scan-no-loop/xteam_scan_no_loop.cpp | 201 ------------------ .../xteam-red-min-max-scan/Makefile | 15 -- .../xteam-red-min-max-scan.cpp | 79 ------- 6 files changed, 530 deletions(-) delete mode 100644 test/smoke-dev/xteam-scan-1/Makefile delete mode 100644 test/smoke-dev/xteam-scan-1/xteam_scan_1.cpp delete mode 100644 test/smoke-dev/xteam-scan-no-loop/Makefile delete mode 100644 test/smoke-dev/xteam-scan-no-loop/xteam_scan_no_loop.cpp delete mode 100644 test/smoke-fails/xteam-red-min-max-scan/Makefile delete mode 100644 test/smoke-fails/xteam-red-min-max-scan/xteam-red-min-max-scan.cpp diff --git a/test/smoke-dev/xteam-scan-1/Makefile b/test/smoke-dev/xteam-scan-1/Makefile deleted file mode 100644 index f8cb1727dd..0000000000 --- a/test/smoke-dev/xteam-scan-1/Makefile +++ /dev/null @@ -1,18 +0,0 @@ -include ../../Makefile.defs - -TESTNAME = xteam_scan_1 -TESTSRC_MAIN = xteam_scan_1.cpp -TESTSRC_AUX = -TESTSRC_ALL = $(TESTSRC_MAIN) $(TESTSRC_AUX) -RUNENV += LIBOMPTARGET_KERNEL_TRACE=1 - -RUNCMD = ./$(TESTNAME) 2>&1 | $(FILECHECK) $(TESTSRC_MAIN) - -CFLAGS += -fopenmp-target-xteam-scan -lm -O0 -CLANG = clang++ -OMP_BIN = $(AOMP)/bin/$(CLANG) -CC = $(OMP_BIN) $(VERBOSE) -#-ccc-print-phases -#"-\#\#\#" - -include ../Makefile.rules diff --git a/test/smoke-dev/xteam-scan-1/xteam_scan_1.cpp b/test/smoke-dev/xteam-scan-1/xteam_scan_1.cpp deleted file mode 100644 index e7a9c5ec4e..0000000000 --- a/test/smoke-dev/xteam-scan-1/xteam_scan_1.cpp +++ /dev/null @@ -1,199 +0,0 @@ -#include -#include -#include - -#define N 2000000 - -template -void run_test() { - T *in = (T*)malloc(sizeof(T) * N); - T *out1 = (T*)malloc(sizeof(T) * N); // For inclusive scan - T *out2 = (T *)malloc(sizeof(T) * N); // For exclusive scan - - for (int i = 0; i < N; i++) { - in[i] = 10; - out1[i] = 0; - } - - T sum1 = T(0); - -#pragma omp target teams distribute parallel for reduction(inscan, +:sum1) map(tofrom: in[0:N], out1[0:N]) num_threads(1024) - for (int i = 0; i < N; i++) { - sum1 += in[i]; // input phase -#pragma omp scan inclusive(sum1) - out1[i] = sum1; // scan phase - } - - T checksum = T(0); - for (int i = 0; i < N; i++) { - checksum += in[i]; - if (checksum != out1[i]) { - printf("Inclusive Scan: Failure. Wrong Result at %d. Exiting...\n", i); - return; - } - } - free(out1); - printf("Inclusive Scan: Success!\n"); - - T sum2 = T(0); - -#pragma omp target teams distribute parallel for reduction(inscan, +:sum2) map(tofrom: in[0:N], out2[0:N]) num_threads(1024) - for (int i = 0; i < N; i++) { - out2[i] = sum2; // scan phase -#pragma omp scan exclusive(sum2) - sum2 += in[i]; // input phase - } - - checksum = T(0); - for (int i = 0; i < N; i++) { - if (checksum != out2[i]) { - printf("Exclusive Scan: Failure. Wrong Result at %d. Exiting...\n", i); - return; - } - checksum += in[i]; - } - free(in); - free(out2); - printf("Exclusive Scan: Success!\n"); -} - -int main() { - printf("Testing for datatype int\n"); - run_test(); - - printf("Testing for datatype uint32_t\n"); - run_test(); - - printf("Testing for datatype uint64_t\n"); - run_test(); - - printf("Testing for datatype long\n"); - run_test(); - - printf("Testing for datatype double\n"); - run_test(); - - printf("Testing for datatype float\n"); - run_test(); - return 0; -} - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args:10 teamsXthrds:( 104X1024) -/// CHECK: n:__omp_offloading_[[MANGLED:.*i.*]]_l20 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args:10 teamsXthrds:( 104X1024) -/// CHECK: n:__omp_offloading_[[MANGLED]]_l20_1 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args:10 teamsXthrds:( 104X1024) -/// CHECK: n:__omp_offloading_[[MANGLED:.*i.*]]_l40 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args:10 teamsXthrds:( 104X1024) -/// CHECK: n:__omp_offloading_[[MANGLED]]_l40_1 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args:10 teamsXthrds:( 104X1024) -/// CHECK: n:__omp_offloading_[[MANGLED:.*j.*]]_l20 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args:10 teamsXthrds:( 104X1024) -/// CHECK: n:__omp_offloading_[[MANGLED]]_l20_1 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args:10 teamsXthrds:( 104X1024) -/// CHECK: n:__omp_offloading_[[MANGLED:.*j.*]]_l40 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args:10 teamsXthrds:( 104X1024) -/// CHECK: n:__omp_offloading_[[MANGLED]]_l40_1 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args:10 teamsXthrds:( 104X1024) -/// CHECK: n:__omp_offloading_[[MANGLED:.*m.*]]_l20 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args:10 teamsXthrds:( 104X1024) -/// CHECK: n:__omp_offloading_[[MANGLED]]_l20_1 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args:10 teamsXthrds:( 104X1024) -/// CHECK: n:__omp_offloading_[[MANGLED:.*m.*]]_l40 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args:10 teamsXthrds:( 104X1024) -/// CHECK: n:__omp_offloading_[[MANGLED]]_l40_1 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args:10 teamsXthrds:( 104X1024) -/// CHECK: n:__omp_offloading_[[MANGLED:.*l.*]]_l20 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args:10 teamsXthrds:( 104X1024) -/// CHECK: n:__omp_offloading_[[MANGLED]]_l20_1 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args:10 teamsXthrds:( 104X1024) -/// CHECK: n:__omp_offloading_[[MANGLED:.*l.*]]_l40 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args:10 teamsXthrds:( 104X1024) -/// CHECK: n:__omp_offloading_[[MANGLED]]_l40_1 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args:10 teamsXthrds:( 104X1024) -/// CHECK: n:__omp_offloading_[[MANGLED:.*d.*]]_l20 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args:10 teamsXthrds:( 104X1024) -/// CHECK: n:__omp_offloading_[[MANGLED]]_l20_1 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args:10 teamsXthrds:( 104X1024) -/// CHECK: n:__omp_offloading_[[MANGLED:.*d.*]]_l40 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args:10 teamsXthrds:( 104X1024) -/// CHECK: n:__omp_offloading_[[MANGLED]]_l40_1 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args:10 teamsXthrds:( 104X1024) -/// CHECK: n:__omp_offloading_[[MANGLED:.*f.*]]_l20 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args:10 teamsXthrds:( 104X1024) -/// CHECK: n:__omp_offloading_[[MANGLED]]_l20_1 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args:10 teamsXthrds:( 104X1024) -/// CHECK: n:__omp_offloading_[[MANGLED:.*f.*]]_l40 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args:10 teamsXthrds:( 104X1024) -/// CHECK: n:__omp_offloading_[[MANGLED]]_l40_1 - -/// CHECK: Testing for datatype int -/// CHECK: Inclusive Scan: Success! -/// CHECK: Exclusive Scan: Success! - -/// CHECK: Testing for datatype uint32_t -/// CHECK: Inclusive Scan: Success! -/// CHECK: Exclusive Scan: Success! - -/// CHECK: Testing for datatype uint64_t -/// CHECK: Inclusive Scan: Success! -/// CHECK: Exclusive Scan: Success! - -/// CHECK: Testing for datatype long -/// CHECK: Inclusive Scan: Success! -/// CHECK: Exclusive Scan: Success! - -/// CHECK: Testing for datatype double -/// CHECK: Inclusive Scan: Success! -/// CHECK: Exclusive Scan: Success! - -/// CHECK: Testing for datatype float -/// CHECK: Inclusive Scan: Success! -/// CHECK: Exclusive Scan: Success! diff --git a/test/smoke-dev/xteam-scan-no-loop/Makefile b/test/smoke-dev/xteam-scan-no-loop/Makefile deleted file mode 100644 index c864751a64..0000000000 --- a/test/smoke-dev/xteam-scan-no-loop/Makefile +++ /dev/null @@ -1,18 +0,0 @@ -include ../../Makefile.defs - -TESTNAME = xteam_scan_no_loop -TESTSRC_MAIN = xteam_scan_no_loop.cpp -TESTSRC_AUX = -TESTSRC_ALL = $(TESTSRC_MAIN) $(TESTSRC_AUX) -RUNENV += LIBOMPTARGET_KERNEL_TRACE=1 - -RUNCMD = ./$(TESTNAME) 2>&1 | $(FILECHECK) $(TESTSRC_MAIN) - -CFLAGS += -fopenmp-target-xteam-no-loop-scan -lm -O0 -CLANG = clang++ -OMP_BIN = $(AOMP)/bin/$(CLANG) -CC = $(OMP_BIN) $(VERBOSE) -#-ccc-print-phases -#"-\#\#\#" - -include ../Makefile.rules diff --git a/test/smoke-dev/xteam-scan-no-loop/xteam_scan_no_loop.cpp b/test/smoke-dev/xteam-scan-no-loop/xteam_scan_no_loop.cpp deleted file mode 100644 index d5a2b8346b..0000000000 --- a/test/smoke-dev/xteam-scan-no-loop/xteam_scan_no_loop.cpp +++ /dev/null @@ -1,201 +0,0 @@ -#include -#include -#include - -#define NUM_TEAMS 250 -#define NUM_THREADS 256 -#define N NUM_THREADS *NUM_TEAMS - -template -void run_test() { - T *in = (T*)malloc(sizeof(T) * N); - T *out1 = (T*)malloc(sizeof(T) * N); // For inclusive scan - T *out2 = (T *)malloc(sizeof(T) * N); // For exclusive scan - - for (int i = 0; i < N; i++) { - in[i] = 10; - out1[i] = 0; - } - - T sum1 = T(0); - -#pragma omp target teams distribute parallel for reduction(inscan, +:sum1) map(tofrom: in[0:N], out1[0:N]) num_teams(NUM_TEAMS) num_threads(NUM_THREADS) - for (int i = 0; i < N; i++) { - sum1 += in[i]; // input phase -#pragma omp scan inclusive(sum1) - out1[i] = sum1; // scan phase - } - - T checksum = T(0); - for (int i = 0; i < N; i++) { - checksum += in[i]; - if (checksum != out1[i]) { - printf("Inclusive Scan: Failure. Wrong Result at %d. Exiting...\n", i); - return; - } - } - free(out1); - printf("Inclusive Scan: Success!\n"); - - T sum2 = T(0); - -#pragma omp target teams distribute parallel for reduction(inscan, +:sum2) map(tofrom: in[0:N], out2[0:N]) num_teams(NUM_TEAMS) num_threads(NUM_THREADS) - for (int i = 0; i < N; i++) { - out2[i] = sum2; // scan phase -#pragma omp scan exclusive(sum2) - sum2 += in[i]; // input phase - } - - checksum = T(0); - for (int i = 0; i < N; i++) { - if (checksum != out2[i]) { - printf("Exclusive Scan: Failure. Wrong Result at %d. Exiting...\n", i); - return; - } - checksum += in[i]; - } - free(in); - free(out2); - printf("Exclusive Scan: Success!\n"); -} - -int main() { - printf("Testing for datatype int\n"); - run_test(); - - printf("Testing for datatype uint32_t\n"); - run_test(); - - printf("Testing for datatype uint64_t\n"); - run_test(); - - printf("Testing for datatype long\n"); - run_test(); - - printf("Testing for datatype double\n"); - run_test(); - - printf("Testing for datatype float\n"); - run_test(); - return 0; -} - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args: 9 teamsXthrds:( 250X 256) -/// CHECK: n:__omp_offloading_[[MANGLED:.*i.*]]_l22 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args: 9 teamsXthrds:( 250X 256) -/// CHECK: n:__omp_offloading_[[MANGLED]]_l22_1 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args: 9 teamsXthrds:( 250X 256) -/// CHECK: n:__omp_offloading_[[MANGLED:.*i.*]]_l42 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args: 9 teamsXthrds:( 250X 256) -/// CHECK: n:__omp_offloading_[[MANGLED]]_l42_1 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args: 9 teamsXthrds:( 250X 256) -/// CHECK: n:__omp_offloading_[[MANGLED:.*j.*]]_l22 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args: 9 teamsXthrds:( 250X 256) -/// CHECK: n:__omp_offloading_[[MANGLED]]_l22_1 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args: 9 teamsXthrds:( 250X 256) -/// CHECK: n:__omp_offloading_[[MANGLED:.*j.*]]_l42 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args: 9 teamsXthrds:( 250X 256) -/// CHECK: n:__omp_offloading_[[MANGLED]]_l42_1 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args: 9 teamsXthrds:( 250X 256) -/// CHECK: n:__omp_offloading_[[MANGLED:.*m.*]]_l22 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args: 9 teamsXthrds:( 250X 256) -/// CHECK: n:__omp_offloading_[[MANGLED]]_l22_1 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args: 9 teamsXthrds:( 250X 256) -/// CHECK: n:__omp_offloading_[[MANGLED:.*m.*]]_l42 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args: 9 teamsXthrds:( 250X 256) -/// CHECK: n:__omp_offloading_[[MANGLED]]_l42_1 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args: 9 teamsXthrds:( 250X 256) -/// CHECK: n:__omp_offloading_[[MANGLED:.*l.*]]_l22 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args: 9 teamsXthrds:( 250X 256) -/// CHECK: n:__omp_offloading_[[MANGLED]]_l22_1 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args: 9 teamsXthrds:( 250X 256) -/// CHECK: n:__omp_offloading_[[MANGLED:.*l.*]]_l42 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args: 9 teamsXthrds:( 250X 256) -/// CHECK: n:__omp_offloading_[[MANGLED]]_l42_1 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args: 9 teamsXthrds:( 250X 256) -/// CHECK: n:__omp_offloading_[[MANGLED:.*d.*]]_l22 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args: 9 teamsXthrds:( 250X 256) -/// CHECK: n:__omp_offloading_[[MANGLED]]_l22_1 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args: 9 teamsXthrds:( 250X 256) -/// CHECK: n:__omp_offloading_[[MANGLED:.*d.*]]_l42 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args: 9 teamsXthrds:( 250X 256) -/// CHECK: n:__omp_offloading_[[MANGLED]]_l42_1 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args: 9 teamsXthrds:( 250X 256) -/// CHECK: n:__omp_offloading_[[MANGLED:.*f.*]]_l22 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args: 9 teamsXthrds:( 250X 256) -/// CHECK: n:__omp_offloading_[[MANGLED]]_l22_1 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args: 9 teamsXthrds:( 250X 256) -/// CHECK: n:__omp_offloading_[[MANGLED:.*f.*]]_l42 - -/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:8 -/// CHECK: args: 9 teamsXthrds:( 250X 256) -/// CHECK: n:__omp_offloading_[[MANGLED]]_l42_1 - -/// CHECK: Testing for datatype int -/// CHECK: Inclusive Scan: Success! -/// CHECK: Exclusive Scan: Success! - -/// CHECK: Testing for datatype uint32_t -/// CHECK: Inclusive Scan: Success! -/// CHECK: Exclusive Scan: Success! - -/// CHECK: Testing for datatype uint64_t -/// CHECK: Inclusive Scan: Success! -/// CHECK: Exclusive Scan: Success! - -/// CHECK: Testing for datatype long -/// CHECK: Inclusive Scan: Success! -/// CHECK: Exclusive Scan: Success! - -/// CHECK: Testing for datatype double -/// CHECK: Inclusive Scan: Success! -/// CHECK: Exclusive Scan: Success! - -/// CHECK: Testing for datatype float -/// CHECK: Inclusive Scan: Success! -/// CHECK: Exclusive Scan: Success! \ No newline at end of file diff --git a/test/smoke-fails/xteam-red-min-max-scan/Makefile b/test/smoke-fails/xteam-red-min-max-scan/Makefile deleted file mode 100644 index 8aa3008ac2..0000000000 --- a/test/smoke-fails/xteam-red-min-max-scan/Makefile +++ /dev/null @@ -1,15 +0,0 @@ -include ../../Makefile.defs - -TESTNAME = xteam-red-min-max-scan -TESTSRC_MAIN = xteam-red-min-max-scan.cpp -TESTSRC_AUX = -TESTSRC_ALL = $(TESTSRC_MAIN) $(TESTSRC_AUX) - -CFLAGS += -fopenmp-target-xteam-scan -lm -O0 -CLANG = clang++ -OMP_BIN = $(AOMP)/bin/$(CLANG) -CC = $(OMP_BIN) $(VERBOSE) -#-ccc-print-phases -#"-\#\#\#" - -include ../Makefile.rules diff --git a/test/smoke-fails/xteam-red-min-max-scan/xteam-red-min-max-scan.cpp b/test/smoke-fails/xteam-red-min-max-scan/xteam-red-min-max-scan.cpp deleted file mode 100644 index 6161d77314..0000000000 --- a/test/smoke-fails/xteam-red-min-max-scan/xteam-red-min-max-scan.cpp +++ /dev/null @@ -1,79 +0,0 @@ -#include -#include -#include -#include - -/* Scan and min/max not yet supported */ - -#define N 2000000 - -int run_test() { - int rc = 0; - float *in = new float[N]; - float *out1 = new float[N]; // For inclusive scan - float *out2 = new float[N]; // For exclusive scan - - for (int i = 0; i < N; i++) { - in[i] = 10; - out1[i] = 0; - } - - float max1 = 0; - float min1 = 3000000; - -#pragma omp target teams distribute parallel for reduction(inscan, min : min1) map(tofrom: in[0:N], out1[0:N]) - for (int i = 0; i < N; i++) { - min1 = fminf(min1, in[i]); // input phase -#pragma omp scan inclusive(min1) - out1[i] = min1; // scan phase - } - - float checksum = 3000000; - for (int i = 0; i < N; i++) { - checksum = fminf(checksum, in[i]); - if (checksum != out1[i]) { - printf("Inclusive Scan: Failure. Wrong Result at %d. Exiting...\n", i); - rc = 1; - break; - } - } - if (!rc) - printf("Inclusive Scan: Success!\n"); - if (min1 != 10) { - printf("Min failed: found %f, expected 10\n", min1); - rc = 1; - } - delete [] out1; - -#pragma omp target teams distribute parallel for reduction(inscan, max: max1) map(tofrom: in[0:N], out2[0:N]) - for (int i = 0; i < N; i++) { - out2[i] = max1; // scan phase -#pragma omp scan exclusive(max1) - max1 = fmaxf(max1, in[i]); // input phase - } - - checksum = 0; - for (int i = 0; i < N; i++) { - if (checksum != out2[i]) { - printf("Exclusive Scan: Failure. Wrong Result at %d. Exiting...\n", i); - rc = 1; - break; - } - checksum = fmaxf(checksum, in[i]); - } - if (!rc) - printf("Exclusive Scan: Success!\n"); - if (max1 != 10) { - printf("Max failed: found %f, expected 10\n", max1); - rc = 1; - } - delete [] in; - delete [] out2; - - return rc; -} - -int main() { - int rc = run_test(); - return rc; -} From 72ff4ed0874768718285a6ea08666f3fcdd11e1e Mon Sep 17 00:00:00 2001 From: agozillon Date: Fri, 29 May 2026 22:31:56 -0500 Subject: [PATCH 3/5] [Flang][OpenMP][Fortran-map-tests] Add size zero presence check test and do some tidying --- test/fortran-map-tests/UMT-reproducer-6.f90 | 84 ------------------- .../size-zero-presence-check.f90 | 44 ++++++++++ test/fortran-map-tests/test.sh | 26 +++--- 3 files changed, 55 insertions(+), 99 deletions(-) delete mode 100644 test/fortran-map-tests/UMT-reproducer-6.f90 create mode 100644 test/fortran-map-tests/size-zero-presence-check.f90 diff --git a/test/fortran-map-tests/UMT-reproducer-6.f90 b/test/fortran-map-tests/UMT-reproducer-6.f90 deleted file mode 100644 index 4f6a0f8393..0000000000 --- a/test/fortran-map-tests/UMT-reproducer-6.f90 +++ /dev/null @@ -1,84 +0,0 @@ -! Copyright © Advanced Micro Devices, Inc., or its affiliates. -! -! SPDX-License-Identifier: MIT - -program prog_a - implicit none - type :: dtype_a - integer :: var_a - end type dtype_a - - type :: dtype_b - integer, pointer, contiguous :: var_b(:) => null() - real(8), pointer, contiguous :: var_c(:,:) => null() - integer :: var_d - real(8), pointer, contiguous :: var_e(:,:,:) => null() - end type dtype_b - - type :: dtype_c - type(dtype_b), pointer :: var_f(:) => null() - type(dtype_a), pointer :: var_g(:) => null() - integer :: var_h - end type dtype_c - - integer :: var_i, var_j - type(dtype_c), pointer :: var_k => null() - - var_j = 4 - allocate(var_k) - allocate(var_k%var_f(var_j)) - allocate(var_k%var_g(var_j)) - loop_a: do var_i = 1, var_j - allocate(var_k%var_f(var_i)%var_e(2,2,2)) - allocate(var_k%var_f(var_i)%var_b(2)) - allocate(var_k%var_f(var_i)%var_c(10,10)) - enddo loop_a - - - !$omp target enter data map(to:var_k) - !$omp target enter data map(always,to:var_k%var_f) - !$omp target enter data map(always,to:var_k%var_g) - do var_i = 1, var_j - !$omp target enter data map(always,to:var_k% var_f(var_i)% var_c) - !$omp target enter data map(always,to:var_k% var_f(var_i)% var_b) - !$omp target enter data map(always,to:var_k% var_f(var_i)% var_e) - end do - - do var_i = 1, var_j - if (associated(var_k%var_f(var_i)%var_c)) then - print *, "var_c IS ASSOCIATED 1" - endif - enddo - - loop_b: do var_i = 1, var_j -!$omp target update from(var_k% var_f(var_i)% var_d) -!!$omp target update from(var_k% var_f(var_i)% var_c) -!$omp target update from(var_k% var_f(var_i)% var_e) - deallocate(var_k%var_f(var_i)%var_c) - deallocate(var_k%var_f(var_i)%var_e) - deallocate(var_k%var_f(var_i)%var_b) - enddo loop_b - - do var_i = 1, var_j - if (associated(var_k%var_f(var_i)%var_c)) then - print *, "======= FORTRAN Test Failed! =======" - stop 1 - endif - enddo - - do var_i = 1, var_j - if (associated(var_k%var_f(var_i)%var_e)) then - print *, "======= FORTRAN Test Failed! =======" - stop 1 - endif - enddo - - do var_i = 1, var_j - if (associated(var_k%var_f(var_i)%var_b)) then - print *, "======= FORTRAN Test Failed! =======" - stop 1 - endif - enddo - - print *, "======= FORTRAN Test Passed! =======" -end program prog_a diff --git a/test/fortran-map-tests/size-zero-presence-check.f90 b/test/fortran-map-tests/size-zero-presence-check.f90 new file mode 100644 index 0000000000..a2c53ad129 --- /dev/null +++ b/test/fortran-map-tests/size-zero-presence-check.f90 @@ -0,0 +1,44 @@ +! Copyright © Advanced Micro Devices, Inc., or its affiliates. +! +! SPDX-License-Identifier: MIT + +program prog_a + implicit none + integer :: var_a + integer :: var_b + real, allocatable :: var_c(:) + integer :: var_d + + var_a = 10 + var_b = 9 + + ! Make it size 0, the main purpose of this test is to make sure + ! the presence check on a size 0 array returns true and lets us + ! pass to the end of the program. + allocate(var_c(var_a:var_b)) + + do var_d = var_a, var_b + var_c(var_d) = real(var_d) + end do + + !$omp target enter data map(to: var_c) + + !$omp target teams distribute parallel do map(present,alloc: var_c) + do var_d = var_a, var_b + var_c(var_d) = var_c(var_d) * 2.0 + end do + !$omp end target teams distribute parallel do + + !$omp target exit data map(from: var_c) + + do var_d = var_a, var_b + if (var_c(var_d) /= real(var_d) * 2.0) then + print *, "======= FORTRAN Test Failed! =======" + stop 1 + end if + end do + + deallocate(var_c) + + print*, "======= FORTRAN Test Passed! =======" +end program diff --git a/test/fortran-map-tests/test.sh b/test/fortran-map-tests/test.sh index f63c90823c..a44e2129e4 100755 --- a/test/fortran-map-tests/test.sh +++ b/test/fortran-map-tests/test.sh @@ -653,10 +653,6 @@ echo "compiling UMT-reproducer-5.f90" $AOMP/bin/flang --offload-arch=$AOMP_GPU -fopenmp UMT-reproducer-5.f90 -o UMT-reproducer-5.out -echo "compiling UMT-reproducer-6.f90" - -$AOMP/bin/flang --offload-arch=$AOMP_GPU -fopenmp UMT-reproducer-6.f90 -o UMT-reproducer-6.out - echo "compiling reproducer-3-47779-v2.f90" $AOMP/bin/flang --offload-arch=$AOMP_GPU -fopenmp reproducer-3-47779-v2.f90 -o reproducer-3-47779-v2.out @@ -829,6 +825,10 @@ echo "compiling lcompiler-1645.f90" $AOMP/bin/flang -fopenmp --offload-arch=$AOMP_GPU lcompiler-1645.f90 -o lcompiler-1645.out +echo "compiling size-zero-presence-check.f90" + +$AOMP/bin/flang -fopenmp --offload-arch=$AOMP_GPU size-zero-presence-check.f90 -o size-zero-presence-check.out + echo "basic exp map" echo "RUNNING TEST: basic-exp-map" @@ -1638,11 +1638,6 @@ echo "UMT Reproducer 5, make sure data is appropriately passed across over multi echo "RUNNING TEST: UMT-reproducer-5" ./UMT-reproducer-5.out -echo "UMT Reproducer 6 map back re-associate bug" - -echo "RUNNING TEST: UMT-reproducer-6" -./UMT-reproducer-6.out - echo "Test mapping of allocatable char array in derived type" echo "RUNNING TEST: char-array-map-test-v1" @@ -1731,7 +1726,7 @@ unset LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS echo "attach never map test" -LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS=0 +export LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS=0 echo "RUNNING TEST: attach_never" ./attach_never.out @@ -1748,11 +1743,6 @@ echo "Test deallocation of a pointer array inside of a derived type using ref_pt echo "RUNNING TEST: SWDEV-564425" ./SWDEV-564425.out -echo "Test deallocation of multiple pointer members inside of a derived type using ref_ptr/ptee" - -echo "RUNNING TEST: SWDEV-564425-v2" -./SWDEV-564425-v2.out - echo "ref_ptr_ptee_alloca_struct_map-v1" echo "RUNNING TEST: ref_ptr_ptee_alloca_struct_map-v1" @@ -1842,6 +1832,12 @@ echo "SWDEV-579431 write-back issue caused by inter-mixing of implicit declare m echo "RUNNING TEST: SWDEV-579431" ./SWDEV-579431.out +echo "size-zero-presence-check, this test verifies present tests via map on size zero arrays return true" + +echo "RUNNING TEST: size-zero-presence-check" + +./size-zero-presence-check.out + # Tests that require XNACK/USM to pass echo "test declare target enter/to usm works reasonably in simple cases" From e07d762247d42e07a34c9b1718e2091012816797 Mon Sep 17 00:00:00 2001 From: Jan Patrick Lehr Date: Thu, 4 Jun 2026 21:40:19 +0200 Subject: [PATCH 4/5] [MiniQMC] Update remote, improve script (#2225) Reverting my last change of the miniQMC remote, as suggested by QMCPack maintainer. The original goal to enable both the QMCPack/miniqmc as a stable/legacy version of the mini-app and the current development version were unsuccessfull. It seems I had done something wrong when I originally posted the PR with the change. This changes a few situation in how we handle missing files or not-set environment variables. It also updates the CMake command with now required cmake options. --- bin/run_miniQMC.sh | 53 +++++++++++++++++++++++++++++++++------------- 1 file changed, 38 insertions(+), 15 deletions(-) diff --git a/bin/run_miniQMC.sh b/bin/run_miniQMC.sh index 2843568746..7e864432e4 100755 --- a/bin/run_miniQMC.sh +++ b/bin/run_miniQMC.sh @@ -23,8 +23,12 @@ export AOMP_USE_CCACHE=0 # Control how many OpenMP threads are used by MiniQMCPack : "${MQMC_OMP_NUM_THREADS:=32}" +if [ -z "$AOMP_GPU" ]; then + echo "Error: Set AOMP_GPU to the target GPU architecture (e.g., gfx90a)" + exit 1 +fi + export PATH=$AOMP/bin:$PATH -#export PATH=/home/janplehr/rocm/trunk/bin:$PATH # We export all these paths, so they will be picked-up in the CMake command. # For libraries that AOMP provides (hsa-runtime, hip, comgr, device-libs), use AOMP. @@ -43,17 +47,23 @@ export rocsolver_DIR=${ROCM}/lib/cmake/rocsolver/ : "${MQMC_BUILD_PREFIX:=$AOMP_REPOS_TEST/miniqmc_build}" # Set the default build directory name : "${MQMC_BUILD_DIR:=${MQMC_BUILD_PREFIX}/build_aomp_clang}" -# Path to the miniqmc source directory -: "${MQMC_SOURCE_DIR:=$AOMP_REPOS_TEST/miniqmc_src}" # how many threads should be used for building miniqmc : "${MQMC_NUM_BUILD_PROCS:=32}" # We pin the version by default, so we have only AOMP as moving target : "${MQMC_GIT_TAG:=9d9d7d3}" +# Path to the miniqmc source directory +: "${MQMC_SOURCE_DIR:=$AOMP_REPOS_TEST/miniqmc_src}" if [ ! -d "$MQMC_SOURCE_DIR" ]; then - git clone https://github.com/QMCPACK/miniqmc.git "$MQMC_SOURCE_DIR" - git checkout "${MQMC_GIT_TAG}" + git clone https://github.com/ye-luo/miniqmc.git "$MQMC_SOURCE_DIR" + pushd "$MQMC_SOURCE_DIR" || exit + git checkout OMP_offload + popd || exit +else + pushd "$MQMC_SOURCE_DIR" || exit + git pull + popd || exit fi rm -rf "${MQMC_BUILD_DIR}" @@ -63,8 +73,8 @@ rm -rf "${MQMC_BUILD_DIR}" cmake -B "${MQMC_BUILD_DIR}" -S "${MQMC_SOURCE_DIR}" \ -DCMAKE_PREFIX_PATH="${AOMP}/lib/cmake;${ROCM}/lib/cmake" \ -DCMAKE_CXX_COMPILER=clang++ \ - -DENABLE_OFFLOAD=ON \ - -DQMC_ENABLE_ROCM=ON \ + -DQMC_GPU="openmp" \ + -DQMC_GPU_ARCHS="${AOMP_GPU}" \ -DCMAKE_CXX_FLAGS='-fopenmp-assume-no-nested-parallelism -DCUDART_VERSION=10000 -DcudaMemoryTypeManaged=hipMemoryTypeManaged ' \ -DAMDGPU_DISABLE_HOST_DEVMEM=ON \ -DCMAKE_VERBOSE_MAKEFILE=ON @@ -81,13 +91,26 @@ echo "Running Tests" # Ensure AOMP runtime libraries are found before /opt/rocm libraries to avoid ABI mismatches export LD_LIBRARY_PATH="${AOMP}/lib:${AOMP}/lib/llvm/lib:${LD_LIBRARY_PATH}" -echo "OMP_NUM_THREADS=${MQMC_OMP_NUM_THREADS} ${MQMC_BUILD_DIR}/bin/check_spo_batched_reduction -n 10" -OMP_NUM_THREADS=${MQMC_OMP_NUM_THREADS} "${MQMC_BUILD_DIR}"/bin/check_spo_batched_reduction -n 10 +# We intentionally continue running even if some binaries are missing. +if [ ! -f "${MQMC_BUILD_DIR}/bin/check_spo_batched_reduction" ]; then + echo "Error: check_spo_batched_reduction binary not found in ${MQMC_BUILD_DIR}/bin" +else + echo "OMP_NUM_THREADS=${MQMC_OMP_NUM_THREADS} ${MQMC_BUILD_DIR}/bin/check_spo_batched_reduction -n 10" + OMP_NUM_THREADS=${MQMC_OMP_NUM_THREADS} "${MQMC_BUILD_DIR}"/bin/check_spo_batched_reduction -n 10 +fi -echo "" -echo "OMP_NUM_THREADS=${MQMC_OMP_NUM_THREADS} ${MQMC_BUILD_DIR}/bin/miniqmc" -OMP_NUM_THREADS=${MQMC_OMP_NUM_THREADS} "${MQMC_BUILD_DIR}"/bin/miniqmc -v +if [ ! -f "${MQMC_BUILD_DIR}/bin/miniqmc" ]; then + echo "Error: miniqmc binary not found in ${MQMC_BUILD_DIR}/bin" +else + echo "" + echo "OMP_NUM_THREADS=${MQMC_OMP_NUM_THREADS} ${MQMC_BUILD_DIR}/bin/miniqmc" + OMP_NUM_THREADS=${MQMC_OMP_NUM_THREADS} "${MQMC_BUILD_DIR}"/bin/miniqmc -v +fi -echo "" -echo "OMP_NUM_THREADS=${MQMC_OMP_NUM_THREADS} ${MQMC_BUILD_DIR}/bin/check_spo -n 10" -OMP_NUM_THREADS=${MQMC_OMP_NUM_THREADS} "${MQMC_BUILD_DIR}"/bin/check_spo -n 10 -v +if [ ! -f "${MQMC_BUILD_DIR}/bin/check_spo" ]; then + echo "Error: check_spo binary not found in ${MQMC_BUILD_DIR}/bin" +else + echo "" + echo "OMP_NUM_THREADS=${MQMC_OMP_NUM_THREADS} ${MQMC_BUILD_DIR}/bin/check_spo -n 10" + OMP_NUM_THREADS=${MQMC_OMP_NUM_THREADS} "${MQMC_BUILD_DIR}"/bin/check_spo -n 10 -v +fi From ae613e4ac47a06b1155ba74e5260f3388365f1fb Mon Sep 17 00:00:00 2001 From: agozillon Date: Thu, 4 Jun 2026 20:18:56 -0500 Subject: [PATCH 5/5] [Flang][OpenMP][fortran-map-tests] Add initialized common block test It turns out initialized common blocks are generated to slightly different LLVM-IR than a zero initialized common block, former is a tuple, latter is a byte array. So good to have a test verifying this also functions and continues to do so. --- .../initialized-common-block-mapping.f90 | 110 ++++++++++++++++++ test/fortran-map-tests/test.sh | 9 ++ 2 files changed, 119 insertions(+) create mode 100644 test/fortran-map-tests/initialized-common-block-mapping.f90 diff --git a/test/fortran-map-tests/initialized-common-block-mapping.f90 b/test/fortran-map-tests/initialized-common-block-mapping.f90 new file mode 100644 index 0000000000..a06e5a0d37 --- /dev/null +++ b/test/fortran-map-tests/initialized-common-block-mapping.f90 @@ -0,0 +1,110 @@ +! Copyright © Advanced Micro Devices, Inc., or its affiliates. +! +! SPDX-License-Identifier: MIT + +! Common blocks can have a special "edge-case" where they are initialized +! common blocks, which basically just means we give them some values on +! initialize. This, however, means the LLVM-IR for a common block is actually +! generated differently, the common block goes from a byte array into a typed +! tuple, so we need to be sure this still lowers and handles fine. + +block data init_blocks + implicit none + + integer :: var_a + real :: var_b(5) + common /cblock_a/ var_a, var_b + + integer :: var_c + real :: var_d(5) + common /cblock_b/ var_c, var_d + + data var_a / 100 / + data var_b / 1.0, 2.0, 3.0, 4.0, 5.0 / + + data var_c / 200 / + data var_d / 10.0, 20.0, 30.0, 40.0, 50.0 / + +end block data + +program prog_a + implicit none + + integer :: var_a + real :: var_b(5) + common /cblock_a/ var_a, var_b + !$omp declare target link(/cblock_a/) + + integer :: var_c + real :: var_d(5) + common /cblock_b/ var_c, var_d + + integer :: i + + print *, "=== Initial values (before target region) ===" + print *, "" + print *, "cblock_a (with declare target link):" + print *, " var_a =", var_a + print *, " var_b =", var_b + print *, "" + print *, "cblock_b (without declare target):" + print *, " var_c =", var_c + print *, " var_d =", var_d + print *, "" + + !$omp target map(tofrom: /cblock_a/, /cblock_b/) + var_a = var_a + 1000 + do i = 1, 5 + var_b(i) = var_b(i) * 2.0 + end do + + var_c = var_c + 2000 + do i = 1, 5 + var_d(i) = var_d(i) * 3.0 + end do + !$omp end target + + print *, "=== Final values (after target region) ===" + print *, "" + print *, "cblock_a (with declare target link):" + print *, " var_a =", var_a + print *, " var_b =", var_b + print *, "" + print *, "cblock_b (without declare target):" + print *, " var_c =", var_c + print *, " var_d =", var_d + print *, "" + + print *, "=== Validating results ===" + + if (var_a /= 1100) then + print *, "FAILED: var_a expected 1100, got", var_a + print *, "======= FORTRAN Test Failed! =======" + stop 1 + end if + + if (var_b(1) /= 2.0 .or. var_b(2) /= 4.0 .or. & + var_b(3) /= 6.0 .or. var_b(4) /= 8.0 .or. & + var_b(5) /= 10.0) then + print *, "FAILED: var_b expected [2.0, 4.0, 6.0, 8.0, 10.0], got", var_b + print *, "======= FORTRAN Test Failed! =======" + stop 1 + end if + + if (var_c /= 2200) then + print *, "FAILED: var_c expected 2200, got", var_c + print *, "======= FORTRAN Test Failed! =======" + stop 1 + end if + + if (var_d(1) /= 30.0 .or. var_d(2) /= 60.0 .or. & + var_d(3) /= 90.0 .or. var_d(4) /= 120.0 .or. & + var_d(5) /= 150.0) then + print *, "FAILED: var_d expected [30.0, 60.0, 90.0, 120.0, 150.0], got", var_d + print *, "======= FORTRAN Test Failed! =======" + stop 1 + end if + + print *, "======= FORTRAN Test Passed! =======" + +end program prog_a diff --git a/test/fortran-map-tests/test.sh b/test/fortran-map-tests/test.sh index a44e2129e4..4c999ed48f 100755 --- a/test/fortran-map-tests/test.sh +++ b/test/fortran-map-tests/test.sh @@ -829,6 +829,10 @@ echo "compiling size-zero-presence-check.f90" $AOMP/bin/flang -fopenmp --offload-arch=$AOMP_GPU size-zero-presence-check.f90 -o size-zero-presence-check.out +echo "compiling initialized-common-block-mapping.f90" + +$AOMP/bin/flang -fopenmp --offload-arch=$AOMP_GPU initialized-common-block-mapping.f90 -o initialized-common-block-mapping.out + echo "basic exp map" echo "RUNNING TEST: basic-exp-map" @@ -1838,6 +1842,11 @@ echo "RUNNING TEST: size-zero-presence-check" ./size-zero-presence-check.out +echo "initialized-common-block-mapping, tests initialized COMMON blocks with declare target link and explicit mapping" + +echo "RUNNING TEST: initialized-common-block-mapping" +./initialized-common-block-mapping.out + # Tests that require XNACK/USM to pass echo "test declare target enter/to usm works reasonably in simple cases"