83 lines · c
1// RUN: %libomptarget-compile-generic -fopenmp-offload-mandatory -O32// RUN: %libomptarget-run-generic3// RUN: %libomptarget-compileopt-generic -fopenmp-offload-mandatory -O34// RUN: %libomptarget-run-generic5 6// REQUIRES: gpu7//8// FIXME: https://github.com/llvm/llvm-project/issues/1612659// UNSUPPORTED: gpu10 11#include <omp.h>12#include <stdio.h>13 14struct IdentTy;15void __kmpc_barrier_simple_spmd(struct IdentTy *Loc, int32_t TId);16void __kmpc_barrier_simple_generic(struct IdentTy *Loc, int32_t TId);17 18#pragma omp begin declare target device_type(nohost)19static int A[512] __attribute__((address_space(3), loader_uninitialized));20static int B[512 * 32] __attribute__((loader_uninitialized));21#pragma omp end declare target22 23int main() {24 printf("Testing simple spmd barrier\n");25 for (int r = 0; r < 50; r++) {26#pragma omp target teams distribute thread_limit(512) num_teams(440)27 for (int j = 0; j < 512 * 32; ++j) {28#pragma omp parallel firstprivate(j)29 {30 int TId = omp_get_thread_num();31 int TeamId = omp_get_team_num();32 int NT = omp_get_num_threads();33 // Sequential34 for (int i = 0; i < NT; ++i) {35 // Test shared memory globals36 if (TId == i)37 A[i] = i + j;38 __kmpc_barrier_simple_spmd(0, TId);39 if (A[i] != i + j)40 __builtin_trap();41 __kmpc_barrier_simple_spmd(0, TId);42 // Test generic globals43 if (TId == i)44 B[TeamId] = i;45 __kmpc_barrier_simple_spmd(0, TId);46 if (B[TeamId] != i)47 __builtin_trap();48 __kmpc_barrier_simple_spmd(0, TId);49 }50 }51 }52 }53 54 printf("Testing simple generic barrier\n");55 for (int r = 0; r < 50; r++) {56#pragma omp target teams distribute thread_limit(512) num_teams(440)57 for (int j = 0; j < 512 * 32; ++j) {58#pragma omp parallel firstprivate(j)59 {60 int TId = omp_get_thread_num();61 int TeamId = omp_get_team_num();62 int NT = omp_get_num_threads();63 // Sequential64 for (int i = 0; i < NT; ++i) {65 if (TId == i)66 A[i] = i + j;67 __kmpc_barrier_simple_generic(0, TId);68 if (A[i] != i + j)69 __builtin_trap();70 __kmpc_barrier_simple_generic(0, TId);71 if (TId == i)72 B[TeamId] = i;73 __kmpc_barrier_simple_generic(0, TId);74 if (B[TeamId] != i)75 __builtin_trap();76 __kmpc_barrier_simple_generic(0, TId);77 }78 }79 }80 }81 return 0;82}83