341 lines · c
1// expected-no-diagnostics2 3#ifndef HEADER4#define HEADER5 6// RUN: %clang_cc1 -DCK1 -verify -fopenmp -triple x86_64-unknown-linux -emit-llvm %s -o - | FileCheck %s --check-prefix=CK17// RUN: %clang_cc1 -DCK1 -fopenmp -x c -triple x86_64-unknown-linux -emit-pch -o %t -fopenmp-version=45 %s8// RUN: %clang_cc1 -DCK1 -fopenmp -x c -triple x86_64-unknown-linux -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=45 | FileCheck %s --check-prefix=CK19// RUN: %clang_cc1 -DCK1 -fopenmp -x c -triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s --check-prefix=CK110// RUN: %clang_cc1 -DCK1 -fopenmp -x c -triple x86_64-unknown-linux -fopenmp-targets=spirv64 -emit-llvm %s -o - | FileCheck %s --check-prefix=CK111// RUN: %clang_cc1 -DCK1 -fopenmp -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --check-prefix=CK112// RUN: %clang_cc1 -DCK1 -fopenmp -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-pch -o %t %s13// RUN: %clang_cc1 -DCK1 -fopenmp -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CK114 15// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -triple x86_64-unknown-linux -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"16// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c -triple x86_64-unknown-linux -emit-pch -o %t -fopenmp-version=45 %s17// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c -triple x86_64-unknown-linux -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=45 | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"18// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c -triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"19// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c -triple x86_64-unknown-linux -fopenmp-targets=spirv64 -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"20// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"21// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-pch -o %t %s22// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"23 24#ifdef CK125 26#define N 10027 28void p_vxv(int *v1, int *v2, int *v3, int n);29void t_vxv(int *v1, int *v2, int *v3, int n);30 31#pragma omp declare variant(t_vxv) match(construct={target})32#pragma omp declare variant(p_vxv) match(construct={parallel})33void vxv(int *v1, int *v2, int *v3, int n) {34 for (int i = 0; i < n; i++) v3[i] = v1[i] * v2[i];35}36// CK1: define dso_local void @vxv37 38void p_vxv(int *v1, int *v2, int *v3, int n) {39#pragma omp for40 for (int i = 0; i < n; i++) v3[i] = v1[i] * v2[i] * 3;41}42// CK1: define dso_local void @p_vxv43 44#pragma omp declare target45void t_vxv(int *v1, int *v2, int *v3, int n) {46#pragma distribute simd47 for (int i = 0; i < n; i++) v3[i] = v1[i] * v2[i] * 2;48}49#pragma omp end declare target50// CK1: define dso_local void @t_vxv51 52 53// CK1-LABEL: define {{[^@]+}}@test54int test(void) {55 int v1[N], v2[N], v3[N];56 57 // init58 for (int i = 0; i < N; i++) {59 v1[i] = (i + 1);60 v2[i] = -(i + 1);61 v3[i] = 0;62 }63 64#pragma omp target teams map(to: v1[:N],v2[:N]) map(from: v3[:N])65 {66 vxv(v1, v2, v3, N);67 }68// CK1: call void @__omp_offloading_[[OFFLOAD:.+]]({{.+}})69 70 vxv(v1, v2, v3, N);71// CK1: call void @vxv72 73#pragma omp parallel74 {75 vxv(v1, v2, v3, N);76 }77// CK1: call void ({{.+}}) @__kmpc_fork_call(ptr {{.+}}, i32 3, ptr [[PARALLEL_REGION:@[^,]+]]78 79 return 0;80}81 82// CK1: define internal void @__omp_offloading_[[OFFLOAD]]({{.+}})83// CK1: call void ({{.+}}) @__kmpc_fork_teams(ptr {{.+}}, i32 3, ptr [[TARGET_REGION:@[^,]+]]84// CK1: define internal void [[TARGET_REGION]](85// CK1: call void @t_vxv86 87// CK1: define internal void [[PARALLEL_REGION]](88// CK1: call void @p_vxv89#endif // CK190 91// RUN: %clang_cc1 -DCK2 -verify -fopenmp -triple x86_64-unknown-linux -emit-llvm %s -o - | FileCheck %s --check-prefix=CK292// RUN: %clang_cc1 -DCK2 -fopenmp -x c -triple x86_64-unknown-linux -emit-pch -o %t -fopenmp-version=45 %s93// RUN: %clang_cc1 -DCK2 -fopenmp -x c -triple x86_64-unknown-linux -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=45 | FileCheck %s --check-prefix=CK294// RUN: %clang_cc1 -DCK2 -fopenmp -x c -triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s --check-prefix=CK295// RUN: %clang_cc1 -DCK2 -fopenmp -x c -triple x86_64-unknown-linux -fopenmp-targets=spirv64 -emit-llvm %s -o - | FileCheck %s --check-prefix=CK296// RUN: %clang_cc1 -DCK2 -fopenmp -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --check-prefix=CK297// RUN: %clang_cc1 -DCK2 -fopenmp -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-pch -o %t %s98// RUN: %clang_cc1 -DCK2 -fopenmp -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CK299 100// RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -triple x86_64-unknown-linux -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"101// RUN: %clang_cc1 -DCK2 -fopenmp-simd -x c -triple x86_64-unknown-linux -emit-pch -o %t -fopenmp-version=45 %s102// RUN: %clang_cc1 -DCK2 -fopenmp-simd -x c -triple x86_64-unknown-linux -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=45 | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"103// RUN: %clang_cc1 -DCK2 -fopenmp-simd -x c -triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"104// RUN: %clang_cc1 -DCK2 -fopenmp-simd -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"105// RUN: %clang_cc1 -DCK2 -fopenmp-simd -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-pch -o %t %s106// RUN: %clang_cc1 -DCK2 -fopenmp-simd -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"107 108#ifdef CK2109 110void test_teams(int ***v1, int ***v2, int ***v3, int n);111void test_target(int ***v1, int ***v2, int ***v3, int n);112void test_parallel(int ***v1, int ***v2, int ***v3, int n);113 114#pragma omp declare variant(test_teams) match(construct = {teams})115#pragma omp declare variant(test_target) match(construct = {target})116#pragma omp declare variant(test_parallel) match(construct = {parallel})117void test_base(int ***v1, int ***v2, int ***v3, int n) {118 for (int i = 0; i < n; i++)119 for (int j = 0; j < n; ++j)120 for (int k = 0; k < n; ++k)121 v3[i][j][k] = v1[i][j][k] * v2[i][j][k];122}123 124#pragma omp declare target125void test_teams(int ***v1, int ***v2, int ***v3, int n) {126#pragma omp distribute parallel for simd collapse(2)127 for (int i = 0; i < n; ++i)128 for (int j = 0; j < n; ++j)129 for (int k = 0; k < n; ++k)130 v3[i][j][k] = v1[i][j][k] * v2[i][j][k];131}132#pragma omp end declare target133 134#pragma omp declare target135void test_target(int ***v1, int ***v2, int ***v3, int n) {136#pragma omp parallel for simd collapse(3)137 for (int i = 0; i < n; ++i)138 for (int j = 0; j < n; ++j)139 for (int k = 0; k < n; ++k)140 v3[i][j][k] = v1[i][j][k] * v2[i][j][k];141}142#pragma omp end declare target143 144void test_parallel(int ***v1, int ***v2, int ***v3, int n) {145#pragma omp for collapse(3)146 for (int i = 0; i < n; ++i)147 for (int j = 0; j < n; ++j)148 for (int k = 0; k < n; ++k)149 v3[i][j][k] = v1[i][j][k] * v2[i][j][k];150}151 152// CK2-LABEL: define {{[^@]+}}@test153void test(int ***v1, int ***v2, int ***v3, int n) {154 int i;155 156#pragma omp target157#pragma omp teams158 {159 test_base(v1, v2, v3, 0);160 }161// CK2: call void @__omp_offloading_[[OFFLOAD_1:.+]]({{.+}})162 163#pragma omp target164 {165 test_base(v1, v2, v3, 0);166 }167// CK2: call void @__omp_offloading_[[OFFLOAD_2:.+]]({{.+}})168 169#pragma omp parallel170 {171 test_base(v1, v2, v3, 0);172 }173// CK2: call void ({{.+}}) @__kmpc_fork_call(ptr {{.+}}, i32 3, ptr [[PARALLEL_REGION:@[^,]+]]174}175 176// CK2: define internal void @__omp_offloading_[[OFFLOAD_1]]({{.+}})177// CK2: call void ({{.+}}) @__kmpc_fork_teams(ptr {{.+}}, i32 3, ptr [[TARGET_REGION_1:@[^,]+]]178// CK2: define internal void [[TARGET_REGION_1]](179// CK2: call void @test_teams180 181// CK2: define internal void @__omp_offloading_[[OFFLOAD_2]]({{.+}})182// CK2: call void @test_target183 184// CK2: define internal void [[PARALLEL_REGION]](185// CK2: call void @test_parallel186 187#endif // CK2188 189// RUN: %clang_cc1 -DCK3 -verify -fopenmp -triple x86_64-unknown-linux -emit-llvm %s -o - | FileCheck %s --check-prefix=CK3190// RUN: %clang_cc1 -DCK3 -fopenmp -x c -triple x86_64-unknown-linux -emit-pch -o %t -fopenmp-version=45 %s191// RUN: %clang_cc1 -DCK3 -fopenmp -x c -triple x86_64-unknown-linux -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=45 | FileCheck %s --check-prefix=CK3192// RUN: %clang_cc1 -DCK3 -fopenmp -x c -triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s --check-prefix=CK3193// RUN: %clang_cc1 -DCK3 -fopenmp -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --check-prefix=CK3194// RUN: %clang_cc1 -DCK3 -fopenmp -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-pch -o %t %s195// RUN: %clang_cc1 -DCK3 -fopenmp -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CK3196 197// RUN: %clang_cc1 -DCK3 -verify -fopenmp-simd -triple x86_64-unknown-linux -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"198// RUN: %clang_cc1 -DCK3 -fopenmp-simd -x c -triple x86_64-unknown-linux -emit-pch -o %t -fopenmp-version=45 %s199// RUN: %clang_cc1 -DCK3 -fopenmp-simd -x c -triple x86_64-unknown-linux -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=45 | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"200// RUN: %clang_cc1 -DCK3 -fopenmp-simd -x c -triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"201// RUN: %clang_cc1 -DCK3 -fopenmp-simd -x c -triple x86_64-unknown-linux -fopenmp-targets=spirv64 -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"202// RUN: %clang_cc1 -DCK3 -fopenmp-simd -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"203// RUN: %clang_cc1 -DCK3 -fopenmp-simd -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-pch -o %t %s204// RUN: %clang_cc1 -DCK3 -fopenmp-simd -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"205 206#ifdef CK3207 208#define N 100209 210int t_for(int *v1, int *v2, int *v3, int n);211int t_simd(int *v1, int *v2, int *v3, int n);212 213#pragma omp declare variant(t_simd) match(construct = {simd})214#pragma omp declare variant(t_for) match(construct = {for})215int t(int *v1, int *v2, int *v3, int idx) {216 return v1[idx] * v2[idx];217}218 219int t_for(int *v1, int *v2, int *v3, int idx) {220 return v1[idx] * v2[idx];221}222 223#pragma omp declare simd224int t_simd(int *v1, int *v2, int *v3, int idx) {225 return v1[idx] * v2[idx];226}227 228// CK3-LABEL: define {{[^@]+}}@test229void test(void) {230 int v1[N], v2[N], v3[N];231 232 // init233 for (int i = 0; i < N; i++) {234 v1[i] = (i + 1);235 v2[i] = -(i + 1);236 v3[i] = 0;237 }238 239#pragma omp simd240 for (int i = 0; i < N; i++) {241 v3[i] = t(v1, v2, v3, i);242 }243// CK3: call = call i32 @t_simd244 245 246#pragma omp for247 for (int i = 0; i < N; i++) {248 v3[i] = t(v1, v2, v3, i);249 }250// CK3: call{{.+}} = call i32 @t_for251}252 253#endif // CK3254 255// RUN: %clang_cc1 -DCK4 -verify -fopenmp -triple x86_64-unknown-linux -emit-llvm %s -o - | FileCheck %s --check-prefix=CK4256// RUN: %clang_cc1 -DCK4 -fopenmp -x c -triple x86_64-unknown-linux -emit-pch -o %t -fopenmp-version=45 %s257// RUN: %clang_cc1 -DCK4 -fopenmp -x c -triple x86_64-unknown-linux -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=45 | FileCheck %s --check-prefix=CK4258// RUN: %clang_cc1 -DCK4 -fopenmp -x c -triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s --check-prefix=CK4259// RUN: %clang_cc1 -DCK4 -fopenmp -x c -triple x86_64-unknown-linux -fopenmp-targets=spirv64 -emit-llvm %s -o - | FileCheck %s --check-prefix=CK4260// RUN: %clang_cc1 -DCK4 -fopenmp -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --check-prefix=CK4261// RUN: %clang_cc1 -DCK4 -fopenmp -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-pch -o %t %s262// RUN: %clang_cc1 -DCK4 -fopenmp -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CK4263 264// RUN: %clang_cc1 -DCK4 -verify -fopenmp-simd -triple x86_64-unknown-linux -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"265// RUN: %clang_cc1 -DCK4 -fopenmp-simd -x c -triple x86_64-unknown-linux -emit-pch -o %t -fopenmp-version=45 %s266// RUN: %clang_cc1 -DCK4 -fopenmp-simd -x c -triple x86_64-unknown-linux -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=45 | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"267// RUN: %clang_cc1 -DCK4 -fopenmp-simd -x c -triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"268// RUN: %clang_cc1 -DCK4 -fopenmp-simd -x c -triple x86_64-unknown-linux -fopenmp-targets=spirv64 -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"269// RUN: %clang_cc1 -DCK4 -fopenmp-simd -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"270// RUN: %clang_cc1 -DCK4 -fopenmp-simd -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-pch -o %t %s271// RUN: %clang_cc1 -DCK4 -fopenmp-simd -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}"272 273#ifdef CK4274 275#define N 100276 277void not_selected_vxv(int *v1, int *v2, int *v3, int n);278void combined_vxv(int *v1, int *v2, int *v3, int n);279void all_vxv(int *v1, int *v2, int *v3, int n);280 281#pragma omp declare variant(all_vxv) match(construct={target,teams,parallel,for,simd})282#pragma omp declare variant(combined_vxv) match(construct={target,teams,parallel,for})283#pragma omp declare variant(not_selected_vxv) match(construct={parallel,for})284void vxv(int *v1, int *v2, int *v3, int n) {285 for (int i = 0; i < n; i++) v3[i] = v1[i] * v2[i];286}287 288void not_selected_vxv(int *v1, int *v2, int *v3, int n) {289 for (int i = 0; i < n; i++) v3[i] = v1[i] * v2[i] * 3;290}291 292#pragma omp declare target293void combined_vxv(int *v1, int *v2, int *v3, int n) {294 for (int i = 0; i < n; i++) v3[i] = v1[i] * v2[i] * 2;295}296#pragma omp end declare target297 298#pragma omp declare target299void all_vxv(int *v1, int *v2, int *v3, int n) {300 for (int i = 0; i < n; i++) v3[i] = v1[i] * v2[i] * 4;301}302#pragma omp end declare target303 304// CK4-LABEL: define {{[^@]+}}@test305void test(void) {306 int v1[N], v2[N], v3[N];307 308 //init309 for (int i = 0; i < N; i++) {310 v1[i] = (i + 1);311 v2[i] = -(i + 1);312 v3[i] = 0;313 }314 315#pragma omp target teams map(to: v1[:N],v2[:N]) map(from: v3[:N])316 {317#pragma omp parallel for318 for (int i = 0; i < N; i++)319 vxv(v1, v2, v3, N);320 }321// CK4: call void @__omp_offloading_[[OFFLOAD_1:.+]]({{.+}})322 323#pragma omp simd324 for (int i = 0; i < N; i++)325 vxv(v1, v2, v3, N);326// CK4: call void @vxv327 328#pragma omp target teams distribute parallel for simd map(from: v3[:N])329 for (int i = 0; i < N; i++)330 for (int i = 0; i < N; i++)331 for (int i = 0; i < N; i++)332 vxv(v1, v2, v3, N);333// CK4: call void @__omp_offloading_[[OFFLOAD_2:.+]]({{.+}})334}335// CK4-DAG: call void @all_vxv336// CK4-DAG: call void @combined_vxv337 338#endif // CK4339 340#endif // HEADER341