diff --git a/test/smoke-limbo/usm-locals-xnack-disabled-xnack-any-apu-maps/usm_locals.cpp b/test/smoke-limbo/usm-locals-xnack-disabled-xnack-any-apu-maps/usm_locals.cpp index f29aa6527..afb629436 100644 --- a/test/smoke-limbo/usm-locals-xnack-disabled-xnack-any-apu-maps/usm_locals.cpp +++ b/test/smoke-limbo/usm-locals-xnack-disabled-xnack-any-apu-maps/usm_locals.cpp @@ -16,10 +16,10 @@ int main() { for(size_t t = 0; t < 5; t++) k[t] = -t; - /// CHECK: data_submit_async: {{.*}} 0 ({{.*}} 4, {{.*}}) - /// CHECK: data_submit_async: {{.*}} 0 ({{.*}} 4, {{.*}}) - /// CHECK: data_submit_async: {{.*}} 0 ({{.*}} 40, {{.*}}) - /// CHECK: data_submit_async: {{.*}} 0 ({{.*}} 20, {{.*}}) + /// CHECK-DAG: data_submit_async: {{.*}} 0 ({{.*}} 4, {{.*}}) + /// CHECK-DAG: data_submit_async: {{.*}} 0 ({{.*}} 4, {{.*}}) + /// CHECK-DAG: data_submit_async: {{.*}} 0 ({{.*}} 40, {{.*}}) + /// CHECK-DAG: data_submit_async: {{.*}} 0 ({{.*}} 20, {{.*}}) #pragma omp target update to(z[:10]) #pragma omp target map(tofrom:k[:5]) map(always, tofrom:x) map(tofrom:y) @@ -30,10 +30,13 @@ int main() { z[t]++; printf("Device pointer for k = %p, k[3] = %d\n", k, k[3]); } - /// CHECK: data_retrieve_async: {{.*}} 0 ({{.*}} 20, {{.*}}) - /// CHECK: data_retrieve_async: {{.*}} 0 ({{.*}} 40, {{.*}}) - /// CHECK: data_retrieve_async: {{.*}} 0 ({{.*}} 4, {{.*}}) - /// CHECK: data_retrieve_async: {{.*}} 0 ({{.*}} 4, {{.*}}) + // Note separate the two DAGs to ensure that retrieving is happening after + // submitting. + /// CHECK: launch_kernel: + /// CHECK-DAG: data_retrieve_async: {{.*}} 0 ({{.*}} 20, {{.*}}) + /// CHECK-DAG: data_retrieve_async: {{.*}} 0 ({{.*}} 40, {{.*}}) + /// CHECK-DAG: data_retrieve_async: {{.*}} 0 ({{.*}} 4, {{.*}}) + /// CHECK-DAG: data_retrieve_async: {{.*}} 0 ({{.*}} 4, {{.*}}) #pragma omp target update from(z[:10]) // Note: when the output is redirected rather than printed at the console, diff --git a/test/smoke-limbo/usm-locals-xnack-disabled-xnack-any/usm_locals.cpp b/test/smoke-limbo/usm-locals-xnack-disabled-xnack-any/usm_locals.cpp index 1901a9af5..daf36c7ef 100644 --- a/test/smoke-limbo/usm-locals-xnack-disabled-xnack-any/usm_locals.cpp +++ b/test/smoke-limbo/usm-locals-xnack-disabled-xnack-any/usm_locals.cpp @@ -16,10 +16,10 @@ int main() { for(size_t t = 0; t < 5; t++) k[t] = -t; - /// CHECK: data_submit_async: {{.*}} 0 ({{.*}} 4, {{.*}}) - /// CHECK: data_submit_async: {{.*}} 0 ({{.*}} 4, {{.*}}) - /// CHECK: data_submit_async: {{.*}} 0 ({{.*}} 40, {{.*}}) - /// CHECK: data_submit_async: {{.*}} 0 ({{.*}} 20, {{.*}}) + /// CHECK-DAG: data_submit_async: {{.*}} 0 ({{.*}} 4, {{.*}}) + /// CHECK-DAG: data_submit_async: {{.*}} 0 ({{.*}} 4, {{.*}}) + /// CHECK-DAG: data_submit_async: {{.*}} 0 ({{.*}} 40, {{.*}}) + /// CHECK-DAG: data_submit_async: {{.*}} 0 ({{.*}} 20, {{.*}}) #pragma omp target update to(z[:10]) #pragma omp target map(tofrom:k[:5]) map(always, tofrom:x) map(tofrom:y) @@ -30,10 +30,13 @@ int main() { z[t]++; printf("Device pointer for k = %p, k[3] = %d\n", k, k[3]); } - /// CHECK: data_retrieve_async: {{.*}} 0 ({{.*}} 20, {{.*}}) - /// CHECK: data_retrieve_async: {{.*}} 0 ({{.*}} 40, {{.*}}) - /// CHECK: data_retrieve_async: {{.*}} 0 ({{.*}} 4, {{.*}}) - /// CHECK: data_retrieve_async: {{.*}} 0 ({{.*}} 4, {{.*}}) + // Note separate the two DAGs to ensure that retrieving is happening after + // submitting. + /// CHECK: launch_kernel: + /// CHECK-DAG: data_retrieve_async: {{.*}} 0 ({{.*}} 20, {{.*}}) + /// CHECK-DAG: data_retrieve_async: {{.*}} 0 ({{.*}} 40, {{.*}}) + /// CHECK-DAG: data_retrieve_async: {{.*}} 0 ({{.*}} 4, {{.*}}) + /// CHECK-DAG: data_retrieve_async: {{.*}} 0 ({{.*}} 4, {{.*}}) #pragma omp target update from(z[:10]) // Note: when the output is redirected rather than printed at the console, diff --git a/test/smoke-limbo/usm-locals-xnack-disabled-xnack-minus/usm_locals.cpp b/test/smoke-limbo/usm-locals-xnack-disabled-xnack-minus/usm_locals.cpp index 72b1f788c..858e5b0db 100644 --- a/test/smoke-limbo/usm-locals-xnack-disabled-xnack-minus/usm_locals.cpp +++ b/test/smoke-limbo/usm-locals-xnack-disabled-xnack-minus/usm_locals.cpp @@ -16,10 +16,10 @@ int main() { for(size_t t = 0; t < 5; t++) k[t] = -t; - /// CHECK: data_submit_async: {{.*}} 0 ({{.*}} 4, {{.*}}) - /// CHECK: data_submit_async: {{.*}} 0 ({{.*}} 4, {{.*}}) - /// CHECK: data_submit_async: {{.*}} 0 ({{.*}} 40, {{.*}}) - /// CHECK: data_submit_async: {{.*}} 0 ({{.*}} 20, {{.*}}) + /// CHECK-DAG: data_submit_async: {{.*}} 0 ({{.*}} 4, {{.*}}) + /// CHECK-DAG: data_submit_async: {{.*}} 0 ({{.*}} 4, {{.*}}) + /// CHECK-DAG: data_submit_async: {{.*}} 0 ({{.*}} 40, {{.*}}) + /// CHECK-DAG: data_submit_async: {{.*}} 0 ({{.*}} 20, {{.*}}) #pragma omp target update to(z[:10]) #pragma omp target map(tofrom:k[:5]) map(always, tofrom:x) map(tofrom:y) @@ -30,10 +30,13 @@ int main() { z[t]++; printf("Device pointer for k = %p, k[3] = %d\n", k, k[3]); } - /// CHECK: data_retrieve_async: {{.*}} 0 ({{.*}} 20, {{.*}}) - /// CHECK: data_retrieve_async: {{.*}} 0 ({{.*}} 40, {{.*}}) - /// CHECK: data_retrieve_async: {{.*}} 0 ({{.*}} 4, {{.*}}) - /// CHECK: data_retrieve_async: {{.*}} 0 ({{.*}} 4, {{.*}}) + // Note separate the two DAGs to ensure that retrieving is happening after + // submitting. + /// CHECK: launch_kernel: + /// CHECK-DAG: data_retrieve_async: {{.*}} 0 ({{.*}} 20, {{.*}}) + /// CHECK-DAG: data_retrieve_async: {{.*}} 0 ({{.*}} 40, {{.*}}) + /// CHECK-DAG: data_retrieve_async: {{.*}} 0 ({{.*}} 4, {{.*}}) + /// CHECK-DAG: data_retrieve_async: {{.*}} 0 ({{.*}} 4, {{.*}}) #pragma omp target update from(z[:10]) // Note: when the output is redirected rather than printed at the console, diff --git a/test/smoke-limbo/usm-locals-xnack-enabled-xnack-any-non-apu/usm_locals.cpp b/test/smoke-limbo/usm-locals-xnack-enabled-xnack-any-non-apu/usm_locals.cpp index 3a1899a5b..bff625d92 100644 --- a/test/smoke-limbo/usm-locals-xnack-enabled-xnack-any-non-apu/usm_locals.cpp +++ b/test/smoke-limbo/usm-locals-xnack-enabled-xnack-any-non-apu/usm_locals.cpp @@ -16,10 +16,10 @@ int main() { for(size_t t = 0; t < 5; t++) k[t] = -t; - /// CHECK: data_submit_async: {{.*}} 0 ({{.*}} 4, {{.*}}) - /// CHECK: data_submit_async: {{.*}} 0 ({{.*}} 4, {{.*}}) - /// CHECK: data_submit_async: {{.*}} 0 ({{.*}} 40, {{.*}}) - /// CHECK: data_submit_async: {{.*}} 0 ({{.*}} 20, {{.*}}) + /// CHECK-DAG: data_submit_async: {{.*}} 0 ({{.*}} 4, {{.*}}) + /// CHECK-DAG: data_submit_async: {{.*}} 0 ({{.*}} 4, {{.*}}) + /// CHECK-DAG: data_submit_async: {{.*}} 0 ({{.*}} 40, {{.*}}) + /// CHECK-DAG: data_submit_async: {{.*}} 0 ({{.*}} 20, {{.*}}) #pragma omp target update to(z[:10]) #pragma omp target map(to:k[:5]) map(always, tofrom:x) map(tofrom:y) @@ -32,9 +32,12 @@ int main() { printf("Device pointer for k = %p, k[3] = %d\n", k, k[3]); } #pragma omp target update from(z[:10]) - /// CHECK: data_retrieve_async: {{.*}} 0 ({{.*}} 40, {{.*}}) - /// CHECK: data_retrieve_async: {{.*}} 0 ({{.*}} 4, {{.*}}) - /// CHECK: data_retrieve_async: {{.*}} 0 ({{.*}} 4, {{.*}}) + // Note separate the two DAGs to ensure that retrieving is happening after + // submitting. + /// CHECK: launch_kernel: + /// CHECK-DAG: data_retrieve_async: {{.*}} 0 ({{.*}} 40, {{.*}}) + /// CHECK-DAG: data_retrieve_async: {{.*}} 0 ({{.*}} 4, {{.*}}) + /// CHECK-DAG: data_retrieve_async: {{.*}} 0 ({{.*}} 4, {{.*}}) // Note: when the output is redirected rather than printed at the console, // the printf'd strings are printed AFTER all the OpenMP runtime library diff --git a/test/smoke-limbo/usm-locals-xnack-enabled-xnack-plus-non-apu/usm_locals.cpp b/test/smoke-limbo/usm-locals-xnack-enabled-xnack-plus-non-apu/usm_locals.cpp index 18b61889b..1ae59f42c 100644 --- a/test/smoke-limbo/usm-locals-xnack-enabled-xnack-plus-non-apu/usm_locals.cpp +++ b/test/smoke-limbo/usm-locals-xnack-enabled-xnack-plus-non-apu/usm_locals.cpp @@ -16,10 +16,10 @@ int main() { for(size_t t = 0; t < 5; t++) k[t] = -t; - /// CHECK: data_submit_async: {{.*}} 0 ({{.*}} 4, {{.*}}) - /// CHECK: data_submit_async: {{.*}} 0 ({{.*}} 4, {{.*}}) - /// CHECK: data_submit_async: {{.*}} 0 ({{.*}} 40, {{.*}}) - /// CHECK: data_submit_async: {{.*}} 0 ({{.*}} 20, {{.*}}) + /// CHECK-DAG: data_submit_async: {{.*}} 0 ({{.*}} 4, {{.*}}) + /// CHECK-DAG: data_submit_async: {{.*}} 0 ({{.*}} 4, {{.*}}) + /// CHECK-DAG: data_submit_async: {{.*}} 0 ({{.*}} 40, {{.*}}) + /// CHECK-DAG: data_submit_async: {{.*}} 0 ({{.*}} 20, {{.*}}) #pragma omp target update to(z[:10]) #pragma omp target map(to:k[:5]) map(always, tofrom:x) map(tofrom:y) @@ -31,9 +31,12 @@ int main() { printf("Device pointer for k = %p, k[3] = %d\n", k, k[3]); } #pragma omp target update from(z[:10]) - /// CHECK: data_retrieve_async: {{.*}} 0 ({{.*}} 40, {{.*}}) - /// CHECK: data_retrieve_async: {{.*}} 0 ({{.*}} 4, {{.*}}) - /// CHECK: data_retrieve_async: {{.*}} 0 ({{.*}} 4, {{.*}}) + // Note separate the two DAGs to ensure that retrieving is happening after + // submitting. + /// CHECK: launch_kernel: + /// CHECK-DAG: data_retrieve_async: {{.*}} 0 ({{.*}} 40, {{.*}}) + /// CHECK-DAG: data_retrieve_async: {{.*}} 0 ({{.*}} 4, {{.*}}) + /// CHECK-DAG: data_retrieve_async: {{.*}} 0 ({{.*}} 4, {{.*}}) // Note: when the output is redirected rather than printed at the console, // the printf'd strings are printed AFTER all the OpenMP runtime library