diff options
Diffstat (limited to 'final/libomptarget/deviceRTLs/nvptx/test/api')
3 files changed, 156 insertions, 0 deletions
diff --git a/final/libomptarget/deviceRTLs/nvptx/test/api/ignored.c b/final/libomptarget/deviceRTLs/nvptx/test/api/ignored.c new file mode 100644 index 0000000..1fa9ae0 --- /dev/null +++ b/final/libomptarget/deviceRTLs/nvptx/test/api/ignored.c @@ -0,0 +1,38 @@ +// RUN: %compile-run-and-check + +#include <omp.h> +#include <stdio.h> + +const int MaxThreads = 1024; + +int main(int argc, char *argv[]) { + int cancellation = -1, dynamic = -1, nested = -1, maxActiveLevels = -1; + + #pragma omp target map(cancellation, dynamic, nested, maxActiveLevels) + { + // libomptarget-nvptx doesn't support cancellation. + cancellation = omp_get_cancellation(); + + // No support for dynamic adjustment of the number of threads. + omp_set_dynamic(1); + dynamic = omp_get_dynamic(); + + // libomptarget-nvptx doesn't support nested parallelism. + omp_set_nested(1); + nested = omp_get_nested(); + + omp_set_max_active_levels(42); + maxActiveLevels = omp_get_max_active_levels(); + } + + // CHECK: cancellation = 0 + printf("cancellation = %d\n", cancellation); + // CHECK: dynamic = 0 + printf("dynamic = %d\n", dynamic); + // CHECK: nested = 0 + printf("nested = %d\n", nested); + // CHECK: maxActiveLevels = 1 + printf("maxActiveLevels = %d\n", maxActiveLevels); + + return 0; +} diff --git a/final/libomptarget/deviceRTLs/nvptx/test/api/max_threads.c b/final/libomptarget/deviceRTLs/nvptx/test/api/max_threads.c new file mode 100644 index 0000000..d0d9f31 --- /dev/null +++ b/final/libomptarget/deviceRTLs/nvptx/test/api/max_threads.c @@ -0,0 +1,46 @@ +// RUN: %compile-run-and-check + +#include <omp.h> +#include <stdio.h> + +int main(int argc, char *argv[]) { + int MaxThreadsL1 = -1, MaxThreadsL2 = -1; + +#pragma omp declare reduction(unique:int \ + : omp_out = (omp_in == 1 ? omp_in : omp_out)) \ + initializer(omp_priv = -1) + + // Non-SPMD mode. +#pragma omp target teams map(MaxThreadsL1, MaxThreadsL2) thread_limit(32) \ + num_teams(1) + { + MaxThreadsL1 = omp_get_max_threads(); +#pragma omp parallel reduction(unique : MaxThreadsL2) + { MaxThreadsL2 = omp_get_max_threads(); } + } + + // CHECK: Non-SPMD MaxThreadsL1 = 32 + printf("Non-SPMD MaxThreadsL1 = %d\n", MaxThreadsL1); + // CHECK: Non-SPMD MaxThreadsL2 = 1 + printf("Non-SPMD MaxThreadsL2 = %d\n", MaxThreadsL2); + + // SPMD mode with full runtime + MaxThreadsL2 = -1; +#pragma omp target parallel reduction(unique : MaxThreadsL2) + { MaxThreadsL2 = omp_get_max_threads(); } + + // CHECK: SPMD with full runtime MaxThreadsL2 = 1 + printf("SPMD with full runtime MaxThreadsL2 = %d\n", MaxThreadsL2); + + // SPMD mode without runtime + MaxThreadsL2 = -1; +#pragma omp target parallel for reduction(unique : MaxThreadsL2) + for (int I = 0; I < 2; ++I) { + MaxThreadsL2 = omp_get_max_threads(); + } + + // CHECK: SPMD without runtime MaxThreadsL2 = 1 + printf("SPMD without runtime MaxThreadsL2 = %d\n", MaxThreadsL2); + + return 0; +} diff --git a/final/libomptarget/deviceRTLs/nvptx/test/api/thread_limit.c b/final/libomptarget/deviceRTLs/nvptx/test/api/thread_limit.c new file mode 100644 index 0000000..626d620 --- /dev/null +++ b/final/libomptarget/deviceRTLs/nvptx/test/api/thread_limit.c @@ -0,0 +1,72 @@ +// RUN: %compile-run-and-check + +#include <omp.h> +#include <stdio.h> + +int main(int argc, char *argv[]) { + int ThreadLimitL0 = -1, ThreadLimitL1 = -1, ThreadLimitL2 = -1; + +#pragma omp declare reduction(unique64:int \ + : omp_out = (omp_in == 64 ? omp_in : omp_out)) \ + initializer(omp_priv = -1) +#pragma omp declare reduction(unique32:int \ + : omp_out = (omp_in == 32 ? omp_in : omp_out)) \ + initializer(omp_priv = -1) + + // Non-SPMD mode. +#pragma omp target teams map(ThreadLimitL0, ThreadLimitL1, ThreadLimitL2) \ + thread_limit(64) num_teams(1) + { + ThreadLimitL0 = omp_get_thread_limit(); +#pragma omp parallel reduction(unique64 \ + : ThreadLimitL1, ThreadLimitL2) num_threads(32) + { + ThreadLimitL1 = omp_get_thread_limit(); +#pragma omp parallel reduction(unique64 : ThreadLimitL2) + { ThreadLimitL2 = omp_get_thread_limit(); } + } + } + + // CHECK: Non-SPMD ThreadLimitL0 = 64 + printf("Non-SPMD ThreadLimitL0 = %d\n", ThreadLimitL0); + // CHECK: Non-SPMD ThreadLimitL1 = 64 + printf("Non-SPMD ThreadLimitL1 = %d\n", ThreadLimitL1); + // CHECK: Non-SPMD ThreadLimitL2 = 64 + printf("Non-SPMD ThreadLimitL2 = %d\n", ThreadLimitL2); + + // SPMD mode with full runtime + ThreadLimitL1 = -1; + ThreadLimitL2 = -1; +#pragma omp target parallel reduction(unique32 \ + : ThreadLimitL1, ThreadLimitL2) \ + num_threads(32) + { + ThreadLimitL1 = omp_get_thread_limit(); +#pragma omp parallel reduction(unique32 : ThreadLimitL2) + { ThreadLimitL2 = omp_get_thread_limit(); } + } + + // CHECK: SPMD with full runtime ThreadLimitL1 = 32 + printf("SPMD with full runtime ThreadLimitL1 = %d\n", ThreadLimitL1); + // CHECK: SPMD with full runtime ThreadLimitL2 = 32 + printf("SPMD with full runtime ThreadLimitL2 = %d\n", ThreadLimitL2); + + // SPMD mode without runtime + ThreadLimitL1 = -1; + ThreadLimitL2 = -1; +#pragma omp target parallel for reduction(unique32 \ + : ThreadLimitL1, ThreadLimitL2) \ + num_threads(32) + for (int I = 0; I < 2; ++I) { + ThreadLimitL1 = omp_get_thread_limit(); +#pragma omp parallel reduction(unique32 : ThreadLimitL2) + { ThreadLimitL2 = omp_get_thread_limit(); } + } + + // CHECK: SPMD without runtime ThreadLimitL1 = 32 + printf("SPMD without runtime ThreadLimitL1 = %d\n", ThreadLimitL1); + // CHECK: SPMD without runtime ThreadLimitL2 = 32 + printf("SPMD without runtime ThreadLimitL2 = %d\n", ThreadLimitL2); + + return 0; +} |