| 1 | // REQUIRES: powerpc-registered-target |
|---|
| 2 | |
|---|
| 3 | // RUN: %clang_cc1 -verify -fopenmp -x c -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-unknown-unknown \ |
|---|
| 4 | // RUN: -emit-llvm-bc %s -o %t-ppc-host.bc |
|---|
| 5 | // RUN: %clang_cc1 -verify -fopenmp -x ir -triple powerpc64le-unknown-unknown -emit-llvm \ |
|---|
| 6 | // RUN: %t-ppc-host.bc -o - | FileCheck %s -check-prefixes HOST,OUTLINED |
|---|
| 7 | // RUN: %clang_cc1 -verify -fopenmp -x c -triple powerpc64le-unknown-unknown -emit-llvm -fopenmp-is-device \ |
|---|
| 8 | // RUN: %s -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s -check-prefixes DEVICE,OUTLINED |
|---|
| 9 | |
|---|
| 10 | // RUN: %clang_cc1 -verify -fopenmp-simd -x c -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-unknown-unknown -emit-llvm-bc %s -o %t-ppc-host.bc |
|---|
| 11 | // RUN: %clang_cc1 -verify -fopenmp-simd -x ir -triple powerpc64le-unknown-unknown -emit-llvm %t-ppc-host.bc -o - | FileCheck --check-prefix SIMD-ONLY0 %s |
|---|
| 12 | // RUN: %clang_cc1 -verify -fopenmp-simd -x c -triple powerpc64le-unknown-unknown -emit-llvm -fopenmp-is-device %s -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck --check-prefix SIMD-ONLY0 %s |
|---|
| 13 | // SIMD-ONLY0-NOT: {{__kmpc|__tgt}} |
|---|
| 14 | |
|---|
| 15 | // expected-no-diagnostics |
|---|
| 16 | int check() { |
|---|
| 17 | int host = omp_is_initial_device(); |
|---|
| 18 | int device; |
|---|
| 19 | #pragma omp target map(tofrom: device) |
|---|
| 20 | { |
|---|
| 21 | device = omp_is_initial_device(); |
|---|
| 22 | } |
|---|
| 23 | |
|---|
| 24 | return host + device; |
|---|
| 25 | } |
|---|
| 26 | |
|---|
| 27 | // The host should get a value of 1: |
|---|
| 28 | // HOST: define{{.*}} @check() |
|---|
| 29 | // HOST: [[HOST:%.*]] = alloca i32 |
|---|
| 30 | // HOST: store i32 1, i32* [[HOST]] |
|---|
| 31 | |
|---|
| 32 | // OUTLINED: define{{.*}} @{{.*}}omp_offloading{{.*}}(i32*{{.*}} [[DEVICE_ARGUMENT:%.*]]) |
|---|
| 33 | // OUTLINED: [[DEVICE_ADDR_STORAGE:%.*]] = alloca i32* |
|---|
| 34 | // OUTLINED: store i32* [[DEVICE_ARGUMENT]], i32** [[DEVICE_ADDR_STORAGE]] |
|---|
| 35 | // OUTLINED: [[DEVICE_ADDR:%.*]] = load i32*, i32** [[DEVICE_ADDR_STORAGE]] |
|---|
| 36 | |
|---|
| 37 | // The outlined function that is called as fallback also runs on the host: |
|---|
| 38 | // HOST: store i32 1, i32* [[DEVICE_ADDR]] |
|---|
| 39 | |
|---|
| 40 | // The device should get a value of 0: |
|---|
| 41 | // DEVICE: store i32 0, i32* [[DEVICE_ADDR]] |
|---|
| 42 | |
|---|