250 lines · c
1// RUN: %libomptarget-compile-run-and-check-generic2// RUN: %libomptarget-compileopt-run-and-check-generic3 4// REQUIRES: gpu5 6// clang-format off7 8#include <omp.h>9#include <stdio.h>10 11#define N 10012#define BLOCK_SHIFT 813 14void print(int *A, int size) {15 for (int i = 0; i < size; ++i) {16 printf("B%dT%d ", A[i] >> BLOCK_SHIFT, A[i] % (1 << BLOCK_SHIFT));17 }18 printf("\n");19}20 21int main() {22 int A[N];23 24#pragma omp target parallel for map(from:A) num_threads(10) schedule(static, 2)25 for (int i = 0; i < N; ++i) {26 A[i] = (omp_get_team_num() << BLOCK_SHIFT) + omp_get_thread_num();27 }28 printf("omp target parallel for thread chunk size %d\n", 2);29 print(A, N);30 31#pragma omp target teams distribute map(from:A) num_teams(10) \32 dist_schedule(static, 2)33 for (int i = 0; i < N; ++i) {34 A[i] = (omp_get_team_num() << BLOCK_SHIFT) + omp_get_thread_num();35 }36 printf("omp target teams distribute block chunk size %d\n", 2);37 print(A, N);38 39#pragma omp target teams distribute parallel for map(from:A) \40 num_teams(10) dist_schedule(static, 2)41 for (int i = 0; i < N; ++i) {42 A[i] = (omp_get_team_num() << BLOCK_SHIFT) + omp_get_thread_num();43 }44 printf("omp target teams distribute parallel for block chunk size %d ", 2);45 printf("thread chunk size default\n");46 print(A, N);47 48#pragma omp target teams distribute parallel for map(from:A) \49 num_teams(10) dist_schedule(static, 2) schedule(static, 3)50 for (int i = 0; i < N; ++i) {51 A[i] = (omp_get_team_num() << BLOCK_SHIFT) + omp_get_thread_num();52 }53 printf("omp target teams distribute parallel for block chunk size %d ", 2);54 printf("thread chunk size %d\n", 3);55 print(A, N);56 57#pragma omp target teams distribute parallel for map(from:A) \58 num_teams(10) dist_schedule(static, 3) schedule(static, 2)59 for (int i = 0; i < N; ++i) {60 A[i] = (omp_get_team_num() << BLOCK_SHIFT) + omp_get_thread_num();61 }62 printf("omp target teams distribute parallel for block chunk size %d ", 3);63 printf("thread chunk size %d\n", 2);64 print(A, N);65 66#pragma omp target teams distribute parallel for map(from:A) \67 num_teams(10) dist_schedule(static, 5) schedule(static, 2)68 for (int i = 0; i < N; ++i) {69 A[i] = (omp_get_team_num() << BLOCK_SHIFT) + omp_get_thread_num();70 }71 printf("omp target teams distribute parallel for block chunk size %d ", 5);72 printf("thread chunk size %d\n", 2);73 print(A, N);74 75#pragma omp target teams distribute parallel for map(from:A) num_teams(10) \76 dist_schedule(static, 49) schedule(static, 2)77 for (int i = 0; i < N; ++i) {78 A[i] = (omp_get_team_num() << BLOCK_SHIFT) + omp_get_thread_num();79 }80 printf("omp target teams distribute parallel for block chunk size %d ", 49);81 printf("thread chunk size %d\n", 2);82 print(A, N);83 84#pragma omp target teams distribute parallel for map(from:A) \85 num_teams(10) num_threads(10) dist_schedule(static, 29)86 for (int i = 0; i < N; ++i) {87 A[i] = (omp_get_team_num() << BLOCK_SHIFT) + omp_get_thread_num();88 }89 printf("omp target teams distribute parallel for block chunk size %d ", 29);90 printf("thread chunk size default\n");91 print(A, N);92 93#pragma omp target teams distribute parallel for map(from:A) \94 num_teams(10) num_threads(10) dist_schedule(static, 101)95 for (int i = 0; i < N; ++i) {96 A[i] = (omp_get_team_num() << BLOCK_SHIFT) + omp_get_thread_num();97 }98 printf("omp target teams distribute parallel for block chunk size %d ", 101);99 printf("thread chunk size default\n");100 print(A, N);101 102#pragma omp target teams distribute parallel for map(from:A) \103 num_teams(9) num_threads(10) schedule(static, 101)104 for (int i = 0; i < N; ++i) {105 A[i] = (omp_get_team_num() << BLOCK_SHIFT) + omp_get_thread_num();106 }107 printf("omp target teams distribute parallel for default block chunk size ");108 printf("thread chunk size %d\n", 101);109 print(A, N);110 return 0;111}112//CHECK: omp target parallel for thread chunk size 2113 114//CHECK-NEXT: B0T0 B0T0 B0T1 B0T1 B0T2 B0T2 B0T3 B0T3 B0T4 B0T4115//CHECK-SAME: B0T5 B0T5 B0T6 B0T6 B0T7 B0T7 B0T8 B0T8 B0T9 B0T9116//CHECK-SAME: B0T0 B0T0 B0T1 B0T1 B0T2 B0T2 B0T3 B0T3 B0T4 B0T4117//CHECK-SAME: B0T5 B0T5 B0T6 B0T6 B0T7 B0T7 B0T8 B0T8 B0T9 B0T9118//CHECK-SAME: B0T0 B0T0 B0T1 B0T1 B0T2 B0T2 B0T3 B0T3 B0T4 B0T4119//CHECK-SAME: B0T5 B0T5 B0T6 B0T6 B0T7 B0T7 B0T8 B0T8 B0T9 B0T9120//CHECK-SAME: B0T0 B0T0 B0T1 B0T1 B0T2 B0T2 B0T3 B0T3 B0T4 B0T4121//CHECK-SAME: B0T5 B0T5 B0T6 B0T6 B0T7 B0T7 B0T8 B0T8 B0T9 B0T9122//CHECK-SAME: B0T0 B0T0 B0T1 B0T1 B0T2 B0T2 B0T3 B0T3 B0T4 B0T4123//CHECK-SAME: B0T5 B0T5 B0T6 B0T6 B0T7 B0T7 B0T8 B0T8 B0T9 B0T9124 125//CHECK: omp target teams distribute block chunk size 2126 127//CHECK-NEXT: B0T0 B0T0 B1T0 B1T0 B2T0 B2T0 B3T0 B3T0 B4T0 B4T0128//CHECK-SAME: B5T0 B5T0 B6T0 B6T0 B7T0 B7T0 B8T0 B8T0 B9T0 B9T0129//CHECK-SAME: B0T0 B0T0 B1T0 B1T0 B2T0 B2T0 B3T0 B3T0 B4T0 B4T0130//CHECK-SAME: B5T0 B5T0 B6T0 B6T0 B7T0 B7T0 B8T0 B8T0 B9T0 B9T0131//CHECK-SAME: B0T0 B0T0 B1T0 B1T0 B2T0 B2T0 B3T0 B3T0 B4T0 B4T0132//CHECK-SAME: B5T0 B5T0 B6T0 B6T0 B7T0 B7T0 B8T0 B8T0 B9T0 B9T0133//CHECK-SAME: B0T0 B0T0 B1T0 B1T0 B2T0 B2T0 B3T0 B3T0 B4T0 B4T0134//CHECK-SAME: B5T0 B5T0 B6T0 B6T0 B7T0 B7T0 B8T0 B8T0 B9T0 B9T0135//CHECK-SAME: B0T0 B0T0 B1T0 B1T0 B2T0 B2T0 B3T0 B3T0 B4T0 B4T0136//CHECK-SAME: B5T0 B5T0 B6T0 B6T0 B7T0 B7T0 B8T0 B8T0 B9T0 B9T0137 138//CHECK: omp target teams distribute parallel for139//CHECK-SAME: block chunk size 2 thread chunk size default140 141//CHECK-NEXT: B0T0 B0T1 B1T0 B1T1 B2T0 B2T1 B3T0 B3T1 B4T0 B4T1142//CHECK-SAME: B5T0 B5T1 B6T0 B6T1 B7T0 B7T1 B8T0 B8T1 B9T0 B9T1143//CHECK-SAME: B0T0 B0T1 B1T0 B1T1 B2T0 B2T1 B3T0 B3T1 B4T0 B4T1144//CHECK-SAME: B5T0 B5T1 B6T0 B6T1 B7T0 B7T1 B8T0 B8T1 B9T0 B9T1145//CHECK-SAME: B0T0 B0T1 B1T0 B1T1 B2T0 B2T1 B3T0 B3T1 B4T0 B4T1146//CHECK-SAME: B5T0 B5T1 B6T0 B6T1 B7T0 B7T1 B8T0 B8T1 B9T0 B9T1147 148//CHECK: omp target teams distribute parallel for149//CHECK-SAME block chunk size 2 thread chunk size 3150 151//CHECK-NEXT: B0T0 B0T0 B1T0 B1T0 B2T0 B2T0 B3T0 B3T0 B4T0 B4T0152//CHECK-SAME: B5T0 B5T0 B6T0 B6T0 B7T0 B7T0 B8T0 B8T0 B9T0 B9T0153//CHECK-SAME: B0T0 B0T0 B1T0 B1T0 B2T0 B2T0 B3T0 B3T0 B4T0 B4T0154//CHECK-SAME: B5T0 B5T0 B6T0 B6T0 B7T0 B7T0 B8T0 B8T0 B9T0 B9T0155//CHECK-SAME: B0T0 B0T0 B1T0 B1T0 B2T0 B2T0 B3T0 B3T0 B4T0 B4T0156//CHECK-SAME: B5T0 B5T0 B6T0 B6T0 B7T0 B7T0 B8T0 B8T0 B9T0 B9T0157//CHECK-SAME: B0T0 B0T0 B1T0 B1T0 B2T0 B2T0 B3T0 B3T0 B4T0 B4T0158//CHECK-SAME: B5T0 B5T0 B6T0 B6T0 B7T0 B7T0 B8T0 B8T0 B9T0 B9T0159//CHECK-SAME: B0T0 B0T0 B1T0 B1T0 B2T0 B2T0 B3T0 B3T0 B4T0 B4T0160//CHECK-SAME: B5T0 B5T0 B6T0 B6T0 B7T0 B7T0 B8T0 B8T0 B9T0 B9T0161 162//CHECK: omp target teams distribute parallel for163//CHECK-SAME: block chunk size 3 thread chunk size 2164 165//CHECK-NEXT: B0T0 B0T0 B0T1 B1T0 B1T0 B1T1 B2T0 B2T0 B2T1166//CHECK-SAME: B3T0 B3T0 B3T1 B4T0 B4T0 B4T1167//CHECK-SAME: B5T0 B5T0 B5T1 B6T0 B6T0 B6T1 B7T0 B7T0 B7T1168//CHECK-SAME: B8T0 B8T0 B8T1 B9T0 B9T0 B9T1169//CHECK-SAME: B0T0 B0T0 B0T1 B1T0 B1T0 B1T1 B2T0 B2T0 B2T1170//CHECK-SAME: B3T0 B3T0 B3T1 B4T0 B4T0 B4T1171//CHECK-SAME: B5T0 B5T0 B5T1 B6T0 B6T0 B6T1 B7T0 B7T0 B7T1172//CHECK-SAME: B8T0 B8T0 B8T1 B9T0 B9T0 B9T1173//CHECK-SAME: B0T0 B0T0 B0T1 B1T0 B1T0 B1T1 B2T0 B2T0 B2T1174//CHECK-SAME: B3T0 B3T0 B3T1 B4T0 B4T0 B4T1175//CHECK-SAME: B5T0 B5T0 B5T1 B6T0 B6T0 B6T1 B7T0 B7T0 B7T1176//CHECK-SAME: B8T0 B8T0 B8T1 B9T0 B9T0 B9T1177//CHECK-SAME: B0T0 B0T0 B0T1 B1T0 B1T0 B1T1 B2T0 B2T0 B2T1 B3T0178 179//CHECK: omp target teams distribute parallel for180//CHECK-SAME: block chunk size 5 thread chunk size 2181 182//CHECK-NEXT: B0T0 B0T0 B0T1 B0T1 B0T2 B1T0 B1T0 B1T1 B1T1 B1T2183//CHECK-SAME: B2T0 B2T0 B2T1 B2T1 B2T2 B3T0 B3T0 B3T1 B3T1 B3T2184//CHECK-SAME: B4T0 B4T0 B4T1 B4T1 B4T2 B5T0 B5T0 B5T1 B5T1 B5T2185//CHECK-SAME: B6T0 B6T0 B6T1 B6T1 B6T2 B7T0 B7T0 B7T1 B7T1 B7T2186//CHECK-SAME: B8T0 B8T0 B8T1 B8T1 B8T2 B9T0 B9T0 B9T1 B9T1 B9T2187//CHECK-SAME: B0T0 B0T0 B0T1 B0T1 B0T2 B1T0 B1T0 B1T1 B1T1 B1T2188//CHECK-SAME: B2T0 B2T0 B2T1 B2T1 B2T2 B3T0 B3T0 B3T1 B3T1 B3T2189//CHECK-SAME: B4T0 B4T0 B4T1 B4T1 B4T2 B5T0 B5T0 B5T1 B5T1 B5T2190//CHECK-SAME: B6T0 B6T0 B6T1 B6T1 B6T2 B7T0 B7T0 B7T1 B7T1 B7T2191//CHECK-SAME: B8T0 B8T0 B8T1 B8T1 B8T2 B9T0 B9T0 B9T1 B9T1 B9T2192 193//CHECK: omp target teams distribute parallel for194//CHECK-SAME: block chunk size 49 thread chunk size 2195 196//CHECK-NEXT: B0T0 B0T0 B0T1 B0T1 B0T2 B0T2 B0T3 B0T3 B0T4 B0T4 B0T5 B0T5197//CHECK-SAME: B0T6 B0T6 B0T7 B0T7 B0T8 B0T8 B0T9 B0T9 B0T10 B0T10 B0T11 B0T11198//CHECK-SAME: B0T12 B0T12 B0T13 B0T13 B0T14 B0T14 B0T15 B0T15 B0T16 B0T16199//CHECK-SAME: B0T17 B0T17 B0T18 B0T18 B0T19 B0T19 B0T20 B0T20 B0T21 B0T21200//CHECK-SAME: B0T22 B0T22 B0T23 B0T23 B0T24201//CHECK-SAME: B1T0 B1T0 B1T1 B1T1 B1T2 B1T2 B1T3 B1T3 B1T4 B1T4 B1T5 B1T5202//CHECK-SAME: B1T6 B1T6 B1T7 B1T7 B1T8 B1T8 B1T9 B1T9 B1T10 B1T10 B1T11 B1T11203//CHECK-SAME: B1T12 B1T12 B1T13 B1T13 B1T14 B1T14 B1T15 B1T15 B1T16 B1T16204//CHECK-SAME: B1T17 B1T17 B1T18 B1T18 B1T19 B1T19 B1T20 B1T20 B1T21 B1T21205//CHECK-SAME: B1T22 B1T22 B1T23 B1T23 B1T24206//CHECK-SAME: B2T0 B2T0207 208//CHECK: omp target teams distribute parallel for209//CHECK-SAME: block chunk size 29 thread chunk size default210 211//CHECK-NEXT: B0T0 B0T1 B0T2 B0T3 B0T4 B0T5 B0T6 B0T7 B0T8 B0T9212//CHECK-SAME: B0T0 B0T1 B0T2 B0T3 B0T4 B0T5 B0T6 B0T7 B0T8 B0T9213//CHECK-SAME: B0T0 B0T1 B0T2 B0T3 B0T4 B0T5 B0T6 B0T7 B0T8214//CHECK-SAME: B1T0 B1T1 B1T2 B1T3 B1T4 B1T5 B1T6 B1T7 B1T8 B1T9215//CHECK-SAME: B1T0 B1T1 B1T2 B1T3 B1T4 B1T5 B1T6 B1T7 B1T8 B1T9216//CHECK-SAME: B1T0 B1T1 B1T2 B1T3 B1T4 B1T5 B1T6 B1T7 B1T8217//CHECK-SAME: B2T0 B2T1 B2T2 B2T3 B2T4 B2T5 B2T6 B2T7 B2T8 B2T9218//CHECK-SAME: B2T0 B2T1 B2T2 B2T3 B2T4 B2T5 B2T6 B2T7 B2T8 B2T9219//CHECK-SAME: B2T0 B2T1 B2T2 B2T3 B2T4 B2T5 B2T6 B2T7 B2T8220//CHECK-SAME: B3T0 B3T1 B3T2 B3T3 B3T4 B3T5 B3T6 B3T7 B3T8 B3T9221//CHECK-SAME: B3T0 B3T1 B3T2222 223//CHECK: omp target teams distribute parallel for224//CHECK-SAME: block chunk size 101 thread chunk size default225 226//CHECK-NEXT: B0T0 B0T1 B0T2 B0T3 B0T4 B0T5 B0T6 B0T7 B0T8 B0T9227//CHECK-SAME: B0T0 B0T1 B0T2 B0T3 B0T4 B0T5 B0T6 B0T7 B0T8 B0T9228//CHECK-SAME: B0T0 B0T1 B0T2 B0T3 B0T4 B0T5 B0T6 B0T7 B0T8 B0T9229//CHECK-SAME: B0T0 B0T1 B0T2 B0T3 B0T4 B0T5 B0T6 B0T7 B0T8 B0T9230//CHECK-SAME: B0T0 B0T1 B0T2 B0T3 B0T4 B0T5 B0T6 B0T7 B0T8 B0T9231//CHECK-SAME: B0T0 B0T1 B0T2 B0T3 B0T4 B0T5 B0T6 B0T7 B0T8 B0T9232//CHECK-SAME: B0T0 B0T1 B0T2 B0T3 B0T4 B0T5 B0T6 B0T7 B0T8 B0T9233//CHECK-SAME: B0T0 B0T1 B0T2 B0T3 B0T4 B0T5 B0T6 B0T7 B0T8 B0T9234//CHECK-SAME: B0T0 B0T1 B0T2 B0T3 B0T4 B0T5 B0T6 B0T7 B0T8 B0T9235//CHECK-SAME: B0T0 B0T1 B0T2 B0T3 B0T4 B0T5 B0T6 B0T7 B0T8 B0T9236 237//CHECK: omp target teams distribute parallel for238//CHECK-SAME: default block chunk size thread chunk size 101239 240//CHECK-NEXT: B0T0 B0T0 B0T0 B0T0 B0T0 B0T0 B0T0 B0T0 B0T0 B0T0241//CHECK-SAME: B1T0 B1T0 B1T0 B1T0 B1T0 B1T0 B1T0 B1T0 B1T0 B1T0242//CHECK-SAME: B2T0 B2T0 B2T0 B2T0 B2T0 B2T0 B2T0 B2T0 B2T0 B2T0243//CHECK-SAME: B3T0 B3T0 B3T0 B3T0 B3T0 B3T0 B3T0 B3T0 B3T0 B3T0244//CHECK-SAME: B4T0 B4T0 B4T0 B4T0 B4T0 B4T0 B4T0 B4T0 B4T0 B4T0245//CHECK-SAME: B5T0 B5T0 B5T0 B5T0 B5T0 B5T0 B5T0 B5T0 B5T0 B5T0246//CHECK-SAME: B6T0 B6T0 B6T0 B6T0 B6T0 B6T0 B6T0 B6T0 B6T0 B6T0247//CHECK-SAME: B7T0 B7T0 B7T0 B7T0 B7T0 B7T0 B7T0 B7T0 B7T0 B7T0248//CHECK-SAME: B8T0 B8T0 B8T0 B8T0 B8T0 B8T0 B8T0 B8T0 B8T0 B8T0249//CHECK-SAME: B0T0 B0T0 B0T0 B0T0 B0T0 B0T0 B0T0 B0T0 B0T0 B0T0250