Skip to content

Commit 44e9a85

Browse files
alexey-bataevalexey-bataev
committed
Rebase
Created using spr 1.3.5
2 parents 0a4edb5 + 4a1a697 commit 44e9a85

File tree

133 files changed

+1834
-928
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

133 files changed

+1834
-928
lines changed

‎clang/include/clang/Basic/BuiltinsAMDGPU.def‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -504,6 +504,8 @@ TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_b128_v4i16, "V4sV4s*1", "nc", "gf
504504
TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_b128_v4f16, "V4hV4h*1", "nc", "gfx12-insts,wavefrontsize64")
505505
TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_b128_v4bf16, "V4yV4y*1", "nc", "gfx12-insts,wavefrontsize64")
506506

507+
TARGET_BUILTIN(__builtin_amdgcn_ds_bpermute_fi_b32, "iii", "nc", "gfx12-insts")
508+
507509
//===----------------------------------------------------------------------===//
508510
// WMMA builtins.
509511
// Postfix w32 indicates the builtin requires wavefront size of 32.

‎clang/lib/Frontend/CompilerInvocation.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -4215,7 +4215,7 @@ bool CompilerInvocation::ParseLangArgs(LangOptions &Opts, ArgList &Args,
42154215
}
42164216
}
42174217

4218-
// Check if -fopenmp is specified and set default version to 5.0.
4218+
// Check if -fopenmp is specified and set default version to 5.1.
42194219
Opts.OpenMP = Args.hasArg(OPT_fopenmp) ? 51 : 0;
42204220
// Check if -fopenmp-simd is specified.
42214221
bool IsSimdSpecified =

‎clang/lib/Frontend/InitPreprocessor.cpp‎

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1466,9 +1466,15 @@ static void InitializePredefinedMacros(const TargetInfo &TI,
14661466
case 50:
14671467
Builder.defineMacro("_OPENMP", "201811");
14681468
break;
1469+
case 51:
1470+
Builder.defineMacro("_OPENMP", "202011");
1471+
break;
14691472
case 52:
14701473
Builder.defineMacro("_OPENMP", "202111");
14711474
break;
1475+
case 60:
1476+
Builder.defineMacro("_OPENMP", "202411");
1477+
break;
14721478
default: // case 51:
14731479
// Default version is OpenMP 5.1
14741480
Builder.defineMacro("_OPENMP", "202011");

‎clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11-err.cl‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,7 @@
22

33
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -verify -emit-llvm -o - %s
44

5-
void test_s_sleep_var(int d)
6-
{
7-
__builtin_amdgcn_s_sleep_var(d); // expected-error {{'__builtin_amdgcn_s_sleep_var' needs target feature gfx12-insts}}
5+
void builtin_test_unsupported(int a, int b) {
6+
__builtin_amdgcn_s_sleep_var(a); // expected-error {{'__builtin_amdgcn_s_sleep_var' needs target feature gfx12-insts}}
7+
b = __builtin_amdgcn_ds_bpermute_fi_b32(a, b); // expected-error {{'__builtin_amdgcn_ds_bpermute_fi_b32' needs target feature gfx12-insts}}
88
}

‎clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl‎

Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -296,3 +296,26 @@ void test_s_buffer_prefetch_data(__amdgpu_buffer_rsrc_t rsrc, unsigned int len)
296296
__builtin_amdgcn_s_buffer_prefetch_data(rsrc, 128, len);
297297
__builtin_amdgcn_s_buffer_prefetch_data(rsrc, 0, 31);
298298
}
299+
300+
// CHECK-LABEL: @test_ds_bpermute_fi_b32(
301+
// CHECK-NEXT: entry:
302+
// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
303+
// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
304+
// CHECK-NEXT: [[B_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
305+
// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr
306+
// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr
307+
// CHECK-NEXT: [[B_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[B_ADDR]] to ptr
308+
// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr [[OUT_ADDR_ASCAST]], align 8
309+
// CHECK-NEXT: store i32 [[A:%.*]], ptr [[A_ADDR_ASCAST]], align 4
310+
// CHECK-NEXT: store i32 [[B:%.*]], ptr [[B_ADDR_ASCAST]], align 4
311+
// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR_ASCAST]], align 4
312+
// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[B_ADDR_ASCAST]], align 4
313+
// CHECK-NEXT: [[TMP2:%.*]] = call i32 @llvm.amdgcn.ds.bpermute.fi.b32(i32 [[TMP0]], i32 [[TMP1]])
314+
// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(1), ptr [[OUT_ADDR_ASCAST]], align 8
315+
// CHECK-NEXT: store i32 [[TMP2]], ptr addrspace(1) [[TMP3]], align 4
316+
// CHECK-NEXT: ret void
317+
//
318+
void test_ds_bpermute_fi_b32(global int* out, int a, int b)
319+
{
320+
*out = __builtin_amdgcn_ds_bpermute_fi_b32(a, b);
321+
}

‎clang/test/OpenMP/declare_mapper_messages.c‎

Lines changed: 23 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -1,10 +1,12 @@
11
// RUN: %clang_cc1 -verify=omp50,expected -fopenmp -fopenmp-version=50 -ferror-limit 100 -DOMP50 %s
22
// RUN: %clang_cc1 -verify=omp51,expected -fopenmp -ferror-limit 100 %s
33
// RUN: %clang_cc1 -verify=expected,omp52 -fopenmp -fopenmp-version=52 -ferror-limit 100 -DOMP52 %s
4+
// RUN: %clang_cc1 -verify=expected,omp60 -fopenmp -fopenmp-version=60 -ferror-limit 100 -DOMP60 %s
45

56
// RUN: %clang_cc1 -verify=omp50,expected -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 -DOMP50 %s
67
// RUN: %clang_cc1 -verify=omp51-simd,expected -fopenmp-simd -ferror-limit 100 %s
78
// RUN: %clang_cc1 -verify=expected,omp52 -fopenmp-simd -fopenmp-version=52 -ferror-limit 100 -DOMP52 %s
9+
// RUN: %clang_cc1 -verify=expected,omp60-simd -fopenmp-simd -fopenmp-version=60 -ferror-limit 100 -DOMP60 %s
810

911
int temp; // expected-note {{'temp' declared here}}
1012

@@ -32,11 +34,11 @@ struct vec { // expec
3234
#pragma omp declare mapper(struct vec v) map(v.len) // expected-error {{redefinition of user-defined mapper for type 'struct vec' with name 'default'}}
3335
#pragma omp declare mapper(int v) map(v) // expected-error {{mapper type must be of struct, union or class type}}
3436

35-
#ifndef OMP52
36-
// omp51-simd-error@+6 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present', 'ompx_hold'}}
37-
// omp50-error@+5 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'ompx_hold'}}
38-
// omp51-error@+4 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present', 'ompx_hold'}}
39-
// expected-error@+3 {{only variable 'vvec' is allowed in map clauses of this 'omp declare mapper' directive}}
37+
#if !defined(OMP52) && !defined(OMP60)
38+
// omp51-simd-error@+6 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present', 'ompx_hold'}}
39+
// omp50-error@+5 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'ompx_hold'}}
40+
// omp51-error@+4 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present', 'ompx_hold'}}
41+
// expected-error@+3 {{only variable 'vvec' is allowed in map clauses of this 'omp declare mapper' directive}}
4042
// expected-error@+2 {{expected at least one clause on '#pragma omp declare mapper' directive}}
4143
// expected-note@+1 {{'it' declared here}}
4244
#pragma omp declare mapper(id2: struct vec vvec) map(iterator(it=0:vvec.len:2), tofrom:vvec.data[it])
@@ -68,15 +70,28 @@ int fun(int arg) {
6870
{}
6971
#pragma omp target map(mapper(aa :vv) // expected-error {{use of undeclared identifier 'aa'}} expected-error {{expected ')'}} expected-error {{call to undeclared function 'mapper'}} expected-note {{to match this '('}}
7072
{}
73+
#ifndef OMP60
7174
#pragma omp target map(mapper(ab) :vv) // expected-error {{missing map type}} expected-error {{cannot find a valid user-defined mapper for type 'struct vec' with name 'ab'}}
75+
#endif
7276
{}
77+
#ifndef OMP60
7378
#pragma omp target map(mapper(ab) :arr[0:2]) // expected-error {{missing map type}} expected-error {{cannot find a valid user-defined mapper for type 'struct vec' with name 'ab'}}
79+
#endif
7480
{}
75-
#pragma omp target map(mapper(aa) :vv) // expected-error {{missing map type}}
81+
#ifndef OMP60
82+
#pragma omp target map(mapper(aa) :vv) // omp50-error {{missing map type}} omp51-error {{missing map type}} omp52-error {{missing map type}} omp51-simd-error {{missing map type}}
7683
{}
77-
#pragma omp target map(mapper(aa) to:d) // expected-error {{mapper type must be of struct, union or class type}} omp52-error{{missing ',' after map type modifier}}
84+
#endif
85+
// expected-error@+4 {{mapper type must be of struct, union or class type}}
86+
// omp52-error@+3 {{missing ',' after map type modifier}}
87+
// omp60-error@+2 {{missing ',' after map type modifier}}
88+
// omp60-simd-error@+1 {{missing ',' after map type modifier}}
89+
#pragma omp target map(mapper(aa) to:d)
7890
{}
79-
#pragma omp target map(mapper(aa) to:vv) map(close mapper(aa) from:v1) map(mapper(aa) to:arr[0]) // omp52-error 4 {{missing ',' after map type modifier}}
91+
// omp52-error@+3 4 {{missing ',' after map type modifier}}
92+
// omp60-error@+2 4 {{missing ',' after map type modifier}}
93+
// omp60-simd-error@+1 4 {{missing ',' after map type modifier}}
94+
#pragma omp target map(mapper(aa) to:vv) map(close mapper(aa) from:v1) map(mapper(aa) to:arr[0])
8095
{}
8196

8297
#pragma omp target update to(mapper) // expected-error {{expected '(' after 'mapper'}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}

‎clang/test/OpenMP/declare_target_ast_print.cpp‎

Lines changed: 5 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -4,6 +4,8 @@
44
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -I %S/Inputs -ast-print %s | FileCheck %s --check-prefix=CHECK --check-prefix=OMP50
55
// RUN: %clang_cc1 -verify -fopenmp -I %S/Inputs -ast-print %s | FileCheck %s --check-prefix=CHECK --check-prefix=OMP51
66
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=52 -I %S/Inputs -ast-print %s | FileCheck %s --check-prefix=CHECK --check-prefix=OMP52
7+
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=60 -I %S/Inputs -ast-print %s | FileCheck %s --check-prefix=CHECK
8+
79
// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -I %S/Inputs -emit-pch -o %t %s
810
// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -std=c++11 -include-pch %t -I %S/Inputs -verify %s -ast-print | FileCheck %s --check-prefix=CHECK --check-prefix=OMP50
911
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -I %S/Inputs -emit-pch -o %t %s
@@ -229,7 +231,7 @@ void f1() {
229231
int b1, b2, b3;
230232
void f2() {
231233
}
232-
#if _OPENMP == 202111
234+
#if _OPENMP >= 202111
233235
#pragma omp declare target enter(b1) enter(b2), enter(b3, f2)
234236
#else
235237
#pragma omp declare target to(b1) to(b2), to(b3, f2)
@@ -336,7 +338,7 @@ int baz() { return 1; }
336338

337339
#pragma omp declare target
338340
int abc1() { return 1; }
339-
#if _OPENMP == 202111
341+
#if _OPENMP >= 202111
340342
#pragma omp declare target enter(abc1) device_type(nohost)
341343
#else
342344
#pragma omp declare target to(abc1) device_type(nohost)
@@ -379,7 +381,7 @@ int main (int argc, char **argv) {
379381
baz<float>();
380382
baz<int>();
381383

382-
#if _OPENMP == 202111
384+
#if _OPENMP >= 202111
383385
#pragma omp declare target enter(foo2)
384386
#else
385387
#pragma omp declare target to (foo2)

0 commit comments

Comments
 (0)