Gericom/teak-llvm
1
Fork 0
mirror of https://github.com/Gericom/teak-llvm.git synced 2025-06-19 11:35:51 -04:00
Code Issues Actions Packages Projects Releases Wiki Activity
master
teak-llvm/clang/test/OpenMP/nvptx_SPMD_codegen.cpp
Alexey Bataev 5e2879320d [OPENMP][NVPTX]Run combined constructs with if clause in SPMD mode. 
Combined constructs with parallel and if clauses without modifiers may
be executed in SPMD mode since if the condition is true for the target
region, it is also true for parallel region and the threads must be run
in parallel.

llvm-svn: 358503
2019-04-16 15:39:12 +00:00

486 lines
17 KiB
C++
Raw Permalink Blame History

// Test target codegen - host bc file has to be created first.
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s
// expected-no-diagnostics
#ifndef HEADER
#define HEADER
int a;
// CHECK-NOT: @__omp_offloading_{{.+}}_exec_mode = weak constant i8 1
// CHECK-DAG: [[DISTR_LIGHT:@.+]] = private unnamed_addr constant %struct.ident_t { i32 0, i32 2050, i32 3, i32 0, i8* getelementptr inbounds
// CHECK-DAG: [[FOR_LIGHT:@.+]] = private unnamed_addr constant %struct.ident_t { i32 0, i32 514, i32 3, i32 0, i8* getelementptr inbounds
// CHECK-DAG: [[LIGHT:@.+]] = private unnamed_addr constant %struct.ident_t { i32 0, i32 2, i32 3, i32 0, i8* getelementptr inbounds
// CHECK-DAG: [[DISTR_FULL:@.+]] = private unnamed_addr constant %struct.ident_t { i32 0, i32 2050, i32 1, i32 0, i8* getelementptr inbounds
// CHECK-DAG: [[FULL:@.+]] = private unnamed_addr constant %struct.ident_t { i32 0, i32 2, i32 1, i32 0, i8* getelementptr inbounds
// CHECK-DAG: [[BAR_LIGHT:@.+]] = private unnamed_addr constant %struct.ident_t { i32 0, i32 66, i32 3, i32 0, i8* getelementptr inbounds
// CHECK-DAG: [[BAR_FULL:@.+]] = private unnamed_addr constant %struct.ident_t { i32 0, i32 66, i32 1, i32 0, i8* getelementptr inbounds
// CHECK-NOT: @__omp_offloading_{{.+}}_exec_mode = weak constant i8 1
void foo() {
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
#pragma omp target teams distribute parallel for simd if(a)
for (int i = 0; i < 10; ++i)
;
#pragma omp target teams distribute parallel for simd schedule(static)
for (int i = 0; i < 10; ++i)
;
#pragma omp target teams distribute parallel for simd schedule(static, 1)
for (int i = 0; i < 10; ++i)
;
#pragma omp target teams distribute parallel for simd schedule(auto)
for (int i = 0; i < 10; ++i)
;
#pragma omp target teams distribute parallel for simd schedule(runtime)
for (int i = 0; i < 10; ++i)
;
#pragma omp target teams distribute parallel for simd schedule(dynamic)
for (int i = 0; i < 10; ++i)
;
#pragma omp target teams distribute parallel for simd schedule(guided)
for (int i = 0; i < 10; ++i)
;
int a;
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
#pragma omp target teams distribute parallel for lastprivate(a)
for (int i = 0; i < 10; ++i)
a = i;
#pragma omp target teams distribute parallel for schedule(static)
for (int i = 0; i < 10; ++i)
;
#pragma omp target teams distribute parallel for schedule(static, 1)
for (int i = 0; i < 10; ++i)
;
#pragma omp target teams distribute parallel for schedule(auto)
for (int i = 0; i < 10; ++i)
;
#pragma omp target teams distribute parallel for schedule(runtime)
for (int i = 0; i < 10; ++i)
;
#pragma omp target teams distribute parallel for schedule(dynamic)
for (int i = 0; i < 10; ++i)
;
#pragma omp target teams distribute parallel for schedule(guided)
for (int i = 0; i < 10; ++i)
;
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
#pragma omp target teams
{
int b;
#pragma omp distribute parallel for simd
for (int i = 0; i < 10; ++i)
;
;
}
#pragma omp target teams
{
int b[] = {2, 3, sizeof(int)};
#pragma omp distribute parallel for simd schedule(static)
for (int i = 0; i < 10; ++i)
;
}
#pragma omp target teams
{
int b;
#pragma omp distribute parallel for simd schedule(static, 1)
for (int i = 0; i < 10; ++i)
;
int &c = b;
}
#pragma omp target teams
#pragma omp distribute parallel for simd schedule(auto)
for (int i = 0; i < 10; ++i)
;
#pragma omp target teams
#pragma omp distribute parallel for simd schedule(runtime)
for (int i = 0; i < 10; ++i)
;
#pragma omp target teams
#pragma omp distribute parallel for simd schedule(dynamic)
for (int i = 0; i < 10; ++i)
;
#pragma omp target teams
#pragma omp distribute parallel for simd schedule(guided)
for (int i = 0; i < 10; ++i)
;
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
#pragma omp target teams
#pragma omp distribute parallel for
for (int i = 0; i < 10; ++i)
;
#pragma omp target teams
#pragma omp distribute parallel for schedule(static)
for (int i = 0; i < 10; ++i)
;
#pragma omp target teams
#pragma omp distribute parallel for schedule(static, 1)
for (int i = 0; i < 10; ++i)
;
#pragma omp target teams
#pragma omp distribute parallel for schedule(auto)
for (int i = 0; i < 10; ++i)
;
#pragma omp target teams
#pragma omp distribute parallel for schedule(runtime)
for (int i = 0; i < 10; ++i)
;
#pragma omp target teams
#pragma omp distribute parallel for schedule(dynamic)
for (int i = 0; i < 10; ++i)
;
#pragma omp target teams
#pragma omp distribute parallel for schedule(guided)
for (int i = 0; i < 10; ++i)
;
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
#pragma omp target
#pragma omp teams
#pragma omp distribute parallel for
for (int i = 0; i < 10; ++i)
;
#pragma omp target
#pragma omp teams
#pragma omp distribute parallel for schedule(static)
for (int i = 0; i < 10; ++i)
;
#pragma omp target
#pragma omp teams
#pragma omp distribute parallel for schedule(static, 1)
for (int i = 0; i < 10; ++i)
;
#pragma omp target
#pragma omp teams
#pragma omp distribute parallel for schedule(auto)
for (int i = 0; i < 10; ++i)
;
#pragma omp target
#pragma omp teams
#pragma omp distribute parallel for schedule(runtime)
for (int i = 0; i < 10; ++i)
;
#pragma omp target
#pragma omp teams
#pragma omp distribute parallel for schedule(dynamic)
for (int i = 0; i < 10; ++i)
;
#pragma omp target
#pragma omp teams
#pragma omp distribute parallel for schedule(guided)
for (int i = 0; i < 10; ++i)
;
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[FULL]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[FULL]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[FULL]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[FULL]]
#pragma omp target parallel for if(a)
for (int i = 0; i < 10; ++i)
;
#pragma omp target parallel for schedule(static)
for (int i = 0; i < 10; ++i)
;
#pragma omp target parallel for schedule(static, 1)
for (int i = 0; i < 10; ++i)
;
#pragma omp target parallel for schedule(auto)
for (int i = 0; i < 10; ++i)
;
#pragma omp target parallel for schedule(runtime)
for (int i = 0; i < 10; ++i)
;
#pragma omp target parallel for schedule(dynamic)
for (int i = 0; i < 10; ++i)
;
#pragma omp target parallel for schedule(guided)
for (int i = 0; i < 10; ++i)
;
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK-DAG: [[BAR_LIGHT]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK-DAG: [[BAR_LIGHT]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK-DAG: [[BAR_LIGHT]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[FULL]]
// CHECK-DAG: [[BAR_FULL]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[FULL]]
// CHECK-DAG: [[BAR_FULL]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[FULL]]
// CHECK-DAG: [[BAR_FULL]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[FULL]]
// CHECK-DAG: [[BAR_FULL]]
#pragma omp target parallel if(a)
#pragma omp for simd
for (int i = 0; i < 10; ++i)
;
#pragma omp target parallel
#pragma omp for simd schedule(static)
for (int i = 0; i < 10; ++i)
;
#pragma omp target parallel
#pragma omp for simd schedule(static, 1)
for (int i = 0; i < 10; ++i)
;
#pragma omp target parallel
#pragma omp for simd schedule(auto)
for (int i = 0; i < 10; ++i)
;
#pragma omp target parallel
#pragma omp for simd schedule(runtime)
for (int i = 0; i < 10; ++i)
;
#pragma omp target parallel
#pragma omp for simd schedule(dynamic)
for (int i = 0; i < 10; ++i)
;
#pragma omp target parallel
#pragma omp for simd schedule(guided)
for (int i = 0; i < 10; ++i)
;
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[FULL]]
// CHECK-DAG: [[BAR_FULL]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK-DAG: [[BAR_LIGHT]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK-DAG: [[BAR_LIGHT]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[FULL]]
// CHECK-DAG: [[BAR_FULL]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[FULL]]
// CHECK-DAG: [[BAR_FULL]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[FULL]]
// CHECK-DAG: [[BAR_FULL]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[FULL]]
// CHECK-DAG: [[BAR_FULL]]
#pragma omp target
#pragma omp parallel
#pragma omp for simd ordered
for (int i = 0; i < 10; ++i)
;
#pragma omp target
#pragma omp parallel
#pragma omp for simd schedule(static)
for (int i = 0; i < 10; ++i)
;
#pragma omp target
#pragma omp parallel
#pragma omp for simd schedule(static, 1)
for (int i = 0; i < 10; ++i)
;
#pragma omp target
#pragma omp parallel
#pragma omp for simd schedule(auto)
for (int i = 0; i < 10; ++i)
;
#pragma omp target
#pragma omp parallel
#pragma omp for simd schedule(runtime)
for (int i = 0; i < 10; ++i)
;
#pragma omp target
#pragma omp parallel
#pragma omp for simd schedule(dynamic)
for (int i = 0; i < 10; ++i)
;
#pragma omp target
#pragma omp parallel
#pragma omp for simd schedule(guided)
for (int i = 0; i < 10; ++i)
;
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0)
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[FULL]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[FULL]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[FULL]]
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
// CHECK-DAG: [[FULL]]
#pragma omp target
#pragma omp parallel for
for (int i = 0; i < 10; ++i)
;
#pragma omp target
#pragma omp parallel for schedule(static)
for (int i = 0; i < 10; ++i)
;
#pragma omp target
#pragma omp parallel for schedule(static, 1)
for (int i = 0; i < 10; ++i)
;
#pragma omp target
#pragma omp parallel for schedule(auto)
for (int i = 0; i < 10; ++i)
;
#pragma omp target
#pragma omp parallel for schedule(runtime)
for (int i = 0; i < 10; ++i)
;
#pragma omp target
#pragma omp parallel for schedule(dynamic)
for (int i = 0; i < 10; ++i)
;
#pragma omp target
#pragma omp parallel for schedule(guided)
for (int i = 0; i < 10; ++i)
;
}
#endif
Reference in New Issue View Git Blame Copy Permalink