298 lines · c
1// RUN: %clang_cc1 -verify -fopenmp -triple x86_64-unknown-linux-gnu -x c -std=c99 -ast-print %s -o - | FileCheck %s --check-prefix=DEFAULT2 3// RUN: %clang_cc1 -verify -fopenmp-simd -triple x86_64-unknown-linux-gnu -x c -std=c99 -ast-print %s -o - | FileCheck %s --check-prefix=DEFAULT4 5// RUN: %clang_cc1 -verify -fopenmp -triple amdgcn-amd-amdhsa -x c -std=c99 -ast-print %s -o - | FileCheck %s --check-prefix=DEFAULT-GPU6 7// RUN: %clang_cc1 -verify -fopenmp-simd -triple amdgcn-amd-amdhsa -x c -std=c99 -ast-print %s -o - | FileCheck %s --check-prefix=DEFAULT-GPU8 9// RUN: %clang_cc1 -verify -fopenmp -triple spirv64-intel -x c -std=c99 -ast-print %s -o - | FileCheck %s --check-prefix=DEFAULT-GPU10 11// RUN: %clang_cc1 -verify -fopenmp-simd -triple spirv64-intel -x c -std=c99 -ast-print %s -o - | FileCheck %s --check-prefix=DEFAULT-GPU12 13// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=52 -DOMP52 -triple x86_64-unknown-linux-gnu -x c -std=c99 -ast-print %s -o - | FileCheck %s --check-prefix=OMP5214 15// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=52 -DOMP52 -triple x86_64-unknown-linux-gnu -x c -std=c99 -ast-print %s -o - | FileCheck %s --check-prefix=OMP5216 17// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=52 -DOMP52 -triple amdgcn-amd-amdhsa -x c -std=c99 -ast-print %s -o - | FileCheck %s --check-prefix=OMP52-GPU18 19// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=52 -DOMP52 -triple amdgcn-amd-amdhsa -x c -std=c99 -ast-print %s -o - | FileCheck %s --check-prefix=OMP52-GPU20 21// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=52 -DOMP52 -triple spirv64-intel -x c -std=c99 -ast-print %s -o - | FileCheck %s --check-prefix=OMP52-GPU22 23// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=52 -DOMP52 -triple spirv64-intel -x c -std=c99 -ast-print %s -o - | FileCheck %s --check-prefix=OMP52-GPU24// expected-no-diagnostics25 26#ifndef HEADER27#define HEADER28 29#ifdef OMP5230void bar(void);31 32#define N 1033void foo1(void) {34#pragma omp metadirective when(device = {kind(cpu)} \35 : parallel) otherwise()36 bar();37#pragma omp metadirective when(implementation = {vendor(score(0) \38 : llvm)}, \39 device = {kind(cpu)} \40 : parallel) otherwise(target teams)41 bar();42#pragma omp metadirective when(device = {kind(gpu)} \43 : target teams) when(implementation = {vendor(llvm)} \44 : parallel) otherwise()45 bar();46#pragma omp metadirective otherwise(target) when(implementation = {vendor(score(5) \47 : llvm)}, \48 device = {kind(cpu, host)} \49 : parallel)50 bar();51#pragma omp metadirective when(user = {condition(N > 10)} \52 : target) when(user = {condition(N == 10)} \53 : parallel)54 bar();55#pragma omp metadirective when(device = {kind(host)} \56 : parallel for)57 for (int i = 0; i < 100; i++)58 ;59#pragma omp metadirective when(implementation = {extension(match_all)} \60 : parallel) otherwise(parallel for)61 for (int i = 0; i < 100; i++)62 ;63#pragma omp metadirective when(implementation = {extension(match_any)} \64 : parallel) otherwise(parallel for)65 for (int i = 0; i < 100; i++)66 ;67#pragma omp metadirective when(implementation = {extension(match_none)} \68 : parallel) otherwise(parallel for)69 for (int i = 0; i < 100; i++)70 ;71 72// Test metadirective with nested OpenMP directive.73 int array[16];74 #pragma omp metadirective when(user = {condition(1)} \75 : parallel for)76 for (int i = 0; i < 16; i++) {77 #pragma omp simd78 for (int j = 0; j < 16; j++)79 array[i] = i;80 }81 82#pragma omp metadirective when(device={arch("amdgcn")}: \83 teams distribute parallel for)\84 otherwise(parallel for)85 for (int i = 0; i < 100; i++)86 ;87 88#pragma omp metadirective when(device={arch("spirv64")}: \89 teams distribute parallel for)\90 otherwise(parallel for)91 for (int i = 0; i < 100; i++)92 ;93 94#pragma omp metadirective when(implementation = {extension(match_all)} \95 : nothing) otherwise(parallel for)96 for (int i = 0; i < 16; i++)97 ;98 99#pragma omp metadirective when(implementation = {extension(match_any)} \100 : parallel) otherwise(nothing)101 for (int i = 0; i < 16; i++)102 ;103 104 105#pragma omp metadirective when(user = {condition(0)} \106 : parallel for)107 for (int i=0; i<10; i++)108 ;109#pragma omp metadirective when(user = {condition(0)} \110 : parallel for) when(implementation = {extension(match_none)} \111 : parallel) otherwise(parallel for)112 for (int i=0; i<10; i++)113 ;114 115 116#pragma omp metadirective when(user = {condition(1)} \117 : parallel for)118 for (int i=0; i<10; i++)119 ;120#pragma omp metadirective when(user = {condition(1)} \121 : parallel for) when(implementation = {extension(match_none)} \122 : parallel) otherwise(parallel for)123 for (int i=0; i<10; i++)124 ;125}126 127// OMP52: void bar(void);128// OMP52: void foo1(void)129// OMP52-NEXT: #pragma omp parallel130// OMP52-NEXT: bar()131// OMP52-NEXT: #pragma omp parallel132// OMP52-NEXT: bar()133// OMP52-NEXT: #pragma omp parallel134// OMP52-NEXT: bar()135// OMP52-NEXT: #pragma omp parallel136// OMP52-NEXT: bar()137// OMP52-NEXT: #pragma omp parallel138// OMP52-NEXT: bar()139// OMP52-NEXT: #pragma omp parallel for140// OMP52-NEXT: for (int i = 0; i < 100; i++)141// OMP52: #pragma omp parallel142// OMP52-NEXT: for (int i = 0; i < 100; i++)143// OMP52: #pragma omp parallel for144// OMP52-NEXT: for (int i = 0; i < 100; i++)145// OMP52: #pragma omp parallel146// OMP52-NEXT: for (int i = 0; i < 100; i++)147// OMP52: #pragma omp parallel for148// OMP52-NEXT: for (int i = 0; i < 16; i++) {149// OMP52-NEXT: #pragma omp simd150// OMP52-NEXT: for (int j = 0; j < 16; j++)151// OMP52-GPU: #pragma omp teams distribute parallel for152// OMP52-GPU-NEXT: for (int i = 0; i < 100; i++)153// OMP52: for (int i = 0; i < 16; i++)154// OMP52: for (int i = 0; i < 16; i++)155 156#else157void bar(void);158 159#define N 10160void foo2(void) {161#pragma omp metadirective when(device = {kind(cpu)} \162 : parallel) default()163 bar();164#pragma omp metadirective when(implementation = {vendor(score(0) \165 : llvm)}, \166 device = {kind(cpu)} \167 : parallel) default(target teams)168 bar();169#pragma omp metadirective when(device = {kind(gpu)} \170 : target teams) when(implementation = {vendor(llvm)} \171 : parallel) default()172 bar();173#pragma omp metadirective default(target) when(implementation = {vendor(score(5) \174 : llvm)}, \175 device = {kind(cpu, host)} \176 : parallel)177 bar();178#pragma omp metadirective when(user = {condition(N > 10)} \179 : target) when(user = {condition(N == 10)} \180 : parallel)181 bar();182#pragma omp metadirective when(device = {kind(host)} \183 : parallel for)184 for (int i = 0; i < 100; i++)185 ;186#pragma omp metadirective when(implementation = {extension(match_all)} \187 : parallel) default(parallel for)188 for (int i = 0; i < 100; i++)189 ;190#pragma omp metadirective when(implementation = {extension(match_any)} \191 : parallel) default(parallel for)192 for (int i = 0; i < 100; i++)193 ;194#pragma omp metadirective when(implementation = {extension(match_none)} \195 : parallel) default(parallel for)196 for (int i = 0; i < 100; i++)197 ;198 199// Test metadirective with nested OpenMP directive.200 int array[16];201 #pragma omp metadirective when(user = {condition(1)} \202 : parallel for)203 for (int i = 0; i < 16; i++) {204 #pragma omp simd205 for (int j = 0; j < 16; j++)206 array[i] = i;207 }208 209#pragma omp metadirective when(device={arch("amdgcn")}: \210 teams distribute parallel for)\211 default(parallel for)212 for (int i = 0; i < 100; i++)213 ;214 215#pragma omp metadirective when(device={arch("spirv64")}: \216 teams distribute parallel for)\217 default(parallel for)218 for (int i = 0; i < 100; i++)219 ; 220 221#pragma omp metadirective when(implementation = {extension(match_all)} \222 : nothing) default(parallel for)223 for (int i = 0; i < 16; i++)224 ;225 226#pragma omp metadirective when(implementation = {extension(match_any)} \227 : parallel) default(nothing)228 for (int i = 0; i < 16; i++)229 ;230 231 232#pragma omp metadirective when(user = {condition(0)} \233 : parallel for)234 for (int i=0; i<10; i++)235 ;236#pragma omp metadirective when(user = {condition(0)} \237 : parallel for) when(implementation = {extension(match_none)} \238 : parallel) default(parallel for)239 for (int i=0; i<10; i++)240 ;241 242 243#pragma omp metadirective when(user = {condition(1)} \244 : parallel for)245 for (int i=0; i<10; i++)246 ;247#pragma omp metadirective when(user = {condition(1)} \248 : parallel for) when(implementation = {extension(match_none)} \249 : parallel) default(parallel for)250 for (int i=0; i<10; i++)251 ;252#if _OPENMP >= 202111253 #pragma omp metadirective when(user = {condition(0)} \254 : parallel for) otherwise()255 for (int i=0; i<10; i++)256 ;257 258 #pragma omp metadirective when(user = {condition(1)} \259 : parallel for) otherwise()260 for (int i=0; i<10; i++)261 ;262#endif263}264 265// DEFAULT: void bar(void);266// DEFAULT: void foo2(void)267// DEFAULT-NEXT: #pragma omp parallel268// DEFAULT-NEXT: bar()269// DEFAULT-NEXT: #pragma omp parallel270// DEFAULT-NEXT: bar()271// DEFAULT-NEXT: #pragma omp parallel272// DEFAULT-NEXT: bar()273// DEFAULT-NEXT: #pragma omp parallel274// DEFAULT-NEXT: bar()275// DEFAULT-NEXT: #pragma omp parallel276// DEFAULT-NEXT: bar()277// DEFAULT-NEXT: #pragma omp parallel for278// DEFAULT-NEXT: for (int i = 0; i < 100; i++)279// DEFAULT: #pragma omp parallel280// DEFAULT-NEXT: for (int i = 0; i < 100; i++)281// DEFAULT: #pragma omp parallel for282// DEFAULT-NEXT: for (int i = 0; i < 100; i++)283// DEFAULT: #pragma omp parallel284// DEFAULT-NEXT: for (int i = 0; i < 100; i++)285// DEFAULT: #pragma omp parallel for286// DEFAULT-NEXT: for (int i = 0; i < 16; i++) {287// DEFAULT-NEXT: #pragma omp simd288// DEFAULT-NEXT: for (int j = 0; j < 16; j++)289// DEFAULT-GPU: #pragma omp teams distribute parallel for290// DEFAULT-GPU-NEXT: for (int i = 0; i < 100; i++)291// DEFAULT: for (int i = 0; i < 16; i++)292// DEFAULT: for (int i = 0; i < 16; i++)293 294 295#endif296#endif297 298