Skip to content

Commit 4f428d3

Browse files
sarnexsarnex
authored
[clang][lit] Add SPIR-V to some OpenMP offload tests (#165775)
Just to get a little more test coverage. Signed-off-by: Nick Sarnie <nick.sarnie@intel.com>
1 parent 240fe7e commit 4f428d3

File tree

4 files changed

+86
-53
lines changed

4 files changed

+86
-53
lines changed

‎clang/test/OpenMP/metadirective_ast_print.c‎

Lines changed: 28 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -2,17 +2,25 @@
22

33
// RUN: %clang_cc1 -verify -fopenmp-simd -triple x86_64-unknown-linux-gnu -x c -std=c99 -ast-print %s -o - | FileCheck %s --check-prefix=DEFAULT
44

5-
// RUN: %clang_cc1 -verify -fopenmp -triple amdgcn-amd-amdhsa -x c -std=c99 -ast-print %s -o - | FileCheck %s --check-prefix=DEFAULT-AMDGCN
5+
// RUN: %clang_cc1 -verify -fopenmp -triple amdgcn-amd-amdhsa -x c -std=c99 -ast-print %s -o - | FileCheck %s --check-prefix=DEFAULT-GPU
66

7-
// RUN: %clang_cc1 -verify -fopenmp-simd -triple amdgcn-amd-amdhsa -x c -std=c99 -ast-print %s -o - | FileCheck %s --check-prefix=DEFAULT-AMDGCN
7+
// RUN: %clang_cc1 -verify -fopenmp-simd -triple amdgcn-amd-amdhsa -x c -std=c99 -ast-print %s -o - | FileCheck %s --check-prefix=DEFAULT-GPU
8+
9+
// RUN: %clang_cc1 -verify -fopenmp -triple spirv64-intel -x c -std=c99 -ast-print %s -o - | FileCheck %s --check-prefix=DEFAULT-GPU
10+
11+
// RUN: %clang_cc1 -verify -fopenmp-simd -triple spirv64-intel -x c -std=c99 -ast-print %s -o - | FileCheck %s --check-prefix=DEFAULT-GPU
812

913
// 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=OMP52
1014

1115
// 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=OMP52
1216

13-
// 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-AMDGCN
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-GPU
18+
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-GPU
20+
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-GPU
1422

15-
// 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-AMDGCN
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-GPU
1624
// expected-no-diagnostics
1725

1826
#ifndef HEADER
@@ -77,6 +85,12 @@ void foo1(void) {
7785
for (int i = 0; i < 100; i++)
7886
;
7987

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+
8094
#pragma omp metadirective when(implementation = {extension(match_all)} \
8195
: nothing) otherwise(parallel for)
8296
for (int i = 0; i < 16; i++)
@@ -134,8 +148,8 @@ void foo1(void) {
134148
// OMP52-NEXT: for (int i = 0; i < 16; i++) {
135149
// OMP52-NEXT: #pragma omp simd
136150
// OMP52-NEXT: for (int j = 0; j < 16; j++)
137-
// OMP52-AMDGCN: #pragma omp teams distribute parallel for
138-
// OMP52-AMDGCN-NEXT: for (int i = 0; i < 100; i++)
151+
// OMP52-GPU: #pragma omp teams distribute parallel for
152+
// OMP52-GPU-NEXT: for (int i = 0; i < 100; i++)
139153
// OMP52: for (int i = 0; i < 16; i++)
140154
// OMP52: for (int i = 0; i < 16; i++)
141155

@@ -198,6 +212,12 @@ void foo2(void) {
198212
for (int i = 0; i < 100; i++)
199213
;
200214

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+
201221
#pragma omp metadirective when(implementation = {extension(match_all)} \
202222
: nothing) default(parallel for)
203223
for (int i = 0; i < 16; i++)
@@ -266,8 +286,8 @@ void foo2(void) {
266286
// DEFAULT-NEXT: for (int i = 0; i < 16; i++) {
267287
// DEFAULT-NEXT: #pragma omp simd
268288
// DEFAULT-NEXT: for (int j = 0; j < 16; j++)
269-
// DEFAULT-AMDGCN: #pragma omp teams distribute parallel for
270-
// DEFAULT-AMDGCN-NEXT: for (int i = 0; i < 100; i++)
289+
// DEFAULT-GPU: #pragma omp teams distribute parallel for
290+
// DEFAULT-GPU-NEXT: for (int i = 0; i < 100; i++)
271291
// DEFAULT: for (int i = 0; i < 16; i++)
272292
// DEFAULT: for (int i = 0; i < 16; i++)
273293

‎clang/test/OpenMP/metadirective_device_arch_codegen.cpp‎

Lines changed: 17 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
1-
// REQUIRES: amdgpu-registered-target
2-
31
// RUN: %clang_cc1 -fopenmp -x c++ -w -std=c++11 -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host.bc
42
// RUN: %clang_cc1 -fopenmp -x c++ -w -std=c++11 -triple amdgcn-amd-amdhsa -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -target-cpu gfx906 -o - | FileCheck %s
3+
// RUN: %clang_cc1 -fopenmp -x c++ -w -std=c++11 -triple x86_64-unknown-unknown -fopenmp-targets=spirv64-intel -emit-llvm-bc %s -o %t-ppc-spirv-host.bc
4+
// RUN: %clang_cc1 -fopenmp -x c++ -w -std=c++11 -triple spirv64-intel -fopenmp-targets=spirv64-intel -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-spirv-host.bc -o - | FileCheck %s
55
// expected-no-diagnostics
66

77

@@ -16,6 +16,12 @@ Inspired from SOLLVE tests:
1616

1717
#define N 1024
1818

19+
#ifdef __AMDGPU__
20+
#define GPU "amdgcn"
21+
#else
22+
#define GPU "spirv64"
23+
#endif
24+
1925
int metadirective1() {
2026

2127
int v1[N], v2[N], v3[N];
@@ -26,7 +32,7 @@ int metadirective1() {
2632
#pragma omp target map(to:v1,v2) map(from:v3, target_device_num) device(default_device)
2733
{
2834
#pragma omp metadirective \
29-
when(device={arch("amdgcn")}: teams distribute parallel for) \
35+
when(device={arch(GPU)}: teams distribute parallel for) \
3036
default(parallel for)
3137

3238
for (int i = 0; i < N; i++) {
@@ -38,28 +44,28 @@ int metadirective1() {
3844
return errors;
3945
}
4046

41-
// CHECK: define weak_odr protected amdgpu_kernel void @[[METADIRECTIVE:.+metadirective1[a-z0-9_]+]]
47+
// CHECK: define weak_odr protected {{amdgpu|spir}}_kernel void @[[METADIRECTIVE:.+metadirective1[a-z0-9_]+]]
4248
// CHECK: entry:
43-
// CHECK: %{{[0-9]}} = call i32 @__kmpc_target_init
49+
// CHECK: %{{[0-9]}} = call{{.*}} i32 @__kmpc_target_init
4450
// CHECK: user_code.entry:
45-
// CHECK: call void @[[METADIRECTIVE]]_omp_outlined
46-
// CHECK-NOT: call void @__kmpc_parallel_51
51+
// CHECK: call{{.*}} void @[[METADIRECTIVE]]_omp_outlined
52+
// CHECK-NOT: call{{.*}} void @__kmpc_parallel_51
4753
// CHECK: ret void
4854

4955

5056
// CHECK: define internal void @[[METADIRECTIVE]]_omp_outlined
5157
// CHECK: entry:
52-
// CHECK: call void @__kmpc_distribute_static_init
58+
// CHECK: call{{.*}} void @__kmpc_distribute_static_init
5359
// CHECK: omp.loop.exit:
54-
// CHECK: call void @__kmpc_distribute_static_fini
60+
// CHECK: call{{.*}} void @__kmpc_distribute_static_fini
5561

5662

5763
// CHECK: define internal void @[[METADIRECTIVE]]_omp_outlined_omp_outlined
5864
// CHECK: entry:
59-
// CHECK: call void @__kmpc_for_static_init_4
65+
// CHECK: call{{.*}} void @__kmpc_for_static_init_4
6066
// CHECK: omp.inner.for.body:
6167
// CHECK: store atomic {{.*}} monotonic
6268
// CHECK: omp.loop.exit:
63-
// CHECK-NEXT: call void @__kmpc_for_static_fini
69+
// CHECK-NEXT: call{{.*}} void @__kmpc_for_static_fini
6470
// CHECK-NEXT: ret void
6571

‎clang/test/OpenMP/thread_limit_amdgpu.c‎

Lines changed: 0 additions & 34 deletions
This file was deleted.

‎clang/test/OpenMP/thread_limit_gpu.c‎

Lines changed: 41 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,41 @@
1+
// Test target codegen - host bc file has to be created first.
2+
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-x86-host.bc
3+
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple amdgcn-amd-amdhsa -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck -check-prefixes=CHECK,CHECK-AMDGPU %s
4+
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux-gnu -fopenmp-targets=spirv64-intel -emit-llvm-bc %s -o %t-x86-spirv-host.bc
5+
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple spirv64-intel -fopenmp-targets=spirv64-intel -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-spirv-host.bc -o - | FileCheck -check-prefixes=CHECK,CHECK-SPIRV %s
6+
// expected-no-diagnostics
7+
8+
#ifndef HEADER
9+
#define HEADER
10+
11+
void foo(int N) {
12+
#pragma omp target teams distribute parallel for simd
13+
for (int i = 0; i < N; ++i)
14+
;
15+
#pragma omp target teams distribute parallel for simd thread_limit(4)
16+
for (int i = 0; i < N; ++i)
17+
;
18+
#pragma omp target teams distribute parallel for simd ompx_attribute(__attribute__((launch_bounds(42, 42))))
19+
for (int i = 0; i < N; ++i)
20+
;
21+
#pragma omp target teams distribute parallel for simd ompx_attribute(__attribute__((launch_bounds(42, 42)))) num_threads(22)
22+
for (int i = 0; i < N; ++i)
23+
;
24+
}
25+
26+
#endif
27+
28+
// CHECK: define weak_odr protected {{amdgpu|spir}}_kernel void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+__Z3fooi_}}l12({{.*}}) #[[ATTR1:.+]] {
29+
// CHECK: define weak_odr protected {{amdgpu|spir}}_kernel void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+__Z3fooi_}}l15({{.*}}) #[[ATTR2:.+]] {
30+
// CHECK: define weak_odr protected {{amdgpu|spir}}_kernel void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+__Z3fooi_}}l18({{.*}}) #[[ATTR3:.+]] {
31+
// CHECK: define weak_odr protected {{amdgpu|spir}}_kernel void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+__Z3fooi_}}l21({{.*}}) #[[ATTR4:.+]] {
32+
33+
// CHECK-AMDGPU: attributes #[[ATTR1]] = { {{.*}} "amdgpu-flat-work-group-size"="1,256" {{.*}} }
34+
// CHECK-AMDGPU: attributes #[[ATTR2]] = { {{.*}} "amdgpu-flat-work-group-size"="1,4" {{.*}} }
35+
// CHECK-AMDGPU: attributes #[[ATTR3]] = { {{.*}} "amdgpu-flat-work-group-size"="1,42" "amdgpu-max-num-workgroups"="42,1,1"{{.*}} }
36+
// CHECK-AMDGPU: attributes #[[ATTR4]] = { {{.*}} "amdgpu-flat-work-group-size"="1,22" "amdgpu-max-num-workgroups"="42,1,1"{{.*}} }
37+
38+
// CHECK-SPIRV: attributes #[[ATTR1]] = { {{.*}} "omp_target_thread_limit"="256" {{.*}} }
39+
// CHECK-SPIRV: attributes #[[ATTR2]] = { {{.*}} "omp_target_thread_limit"="4" {{.*}} }
40+
// CHECK-SPIRV: attributes #[[ATTR3]] = { {{.*}} "omp_target_num_teams"="42" "omp_target_thread_limit"="42" {{.*}} }
41+
// CHECK-SPIRV: attributes #[[ATTR4]] = { {{.*}} "omp_target_num_teams"="42" "omp_target_thread_limit"="22" {{.*}} }

0 commit comments

Comments
 (0)