ABataev updated this revision to Diff 231265.
ABataev added a comment.
Address comments.
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D70739/new/
https://reviews.llvm.org/D70739
Files:
clang/include/clang/Basic/Attr.td
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/include/clang/Basic/OpenMPKinds.def
clang/include/clang/Sema/Sema.h
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/lib/Parse/ParseOpenMP.cpp
clang/lib/Sema/SemaOpenMP.cpp
clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
clang/test/OpenMP/declare_variant_ast_print.c
clang/test/OpenMP/declare_variant_ast_print.cpp
clang/test/OpenMP/declare_variant_device_isa_codegen.cpp
clang/test/OpenMP/declare_variant_messages.c
clang/test/OpenMP/declare_variant_messages.cpp
clang/test/OpenMP/declare_variant_mixed_codegen.cpp
clang/test/OpenMP/nvptx_declare_variant_device_isa_codegen.cpp
Index: clang/test/OpenMP/nvptx_declare_variant_device_isa_codegen.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/nvptx_declare_variant_device_isa_codegen.cpp
@@ -0,0 +1,154 @@
+// 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 -fopenmp-version=50
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -fopenmp-version=50 | FileCheck %s --implicit-check-not='ret i32 {{1|81|84}}'
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t -fopenmp-version=50
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -o - -fopenmp-version=50 | FileCheck %s --implicit-check-not='ret i32 {{1|81|84}}'
+// expected-no-diagnostics
+
+// CHECK-NOT: ret i32 {{1|81|84}}
+// CHECK-DAG: define {{.*}}i32 @_Z3barv()
+// CHECK-DAG: define {{.*}}i32 @_ZN16SpecSpecialFuncs6MethodEv(%struct.SpecSpecialFuncs* %{{.+}})
+// CHECK-DAG: define {{.*}}i32 @_ZN12SpecialFuncs6MethodEv(%struct.SpecialFuncs* %{{.+}})
+// CHECK-DAG: define linkonce_odr {{.*}}i32 @_ZN16SpecSpecialFuncs6methodEv(%struct.SpecSpecialFuncs* %{{.+}})
+// CHECK-DAG: define linkonce_odr {{.*}}i32 @_ZN12SpecialFuncs6methodEv(%struct.SpecialFuncs* %{{.+}})
+// CHECK-DAG: define {{.*}}i32 @_Z5prio_v()
+// CHECK-DAG: define internal i32 @_ZL6prio1_v()
+// CHECK-DAG: define {{.*}}i32 @_Z4callv()
+// CHECK-DAG: define internal i32 @_ZL9stat_usedv()
+// CHECK-DAG: define {{.*}}i32 @fn_linkage()
+// CHECK-DAG: define {{.*}}i32 @_Z11fn_linkage1v()
+
+// CHECK-DAG: ret i32 2
+// CHECK-DAG: ret i32 3
+// CHECK-DAG: ret i32 4
+// CHECK-DAG: ret i32 5
+// CHECK-DAG: ret i32 6
+// CHECK-DAG: ret i32 7
+// CHECK-DAG: ret i32 82
+// CHECK-DAG: ret i32 83
+// CHECK-DAG: ret i32 85
+// CHECK-DAG: ret i32 86
+// CHECK-DAG: ret i32 87
+
+// Outputs for function members
+// CHECK-DAG: ret i32 6
+// CHECK-DAG: ret i32 7
+// CHECK-NOT: ret i32 {{1|81|84}}
+
+#ifndef HEADER
+#define HEADER
+
+int foo() { return 2; }
+int bazzz();
+int test();
+static int stat_unused_();
+static int stat_used_();
+
+#pragma omp declare target
+
+#pragma omp declare variant(foo) match(device = {isa("nvptx64")})
+int bar() { return 1; }
+
+#pragma omp declare variant(bazzz) match(device = {isa("nvptx64")})
+int baz() { return 1; }
+
+#pragma omp declare variant(test) match(device = {isa("nvptx64")})
+int call() { return 1; }
+
+#pragma omp declare variant(stat_unused_) match(device = {isa("nvptx64")})
+static int stat_unused() { return 1; }
+
+#pragma omp declare variant(stat_used_) match(device = {isa("nvptx64")})
+static int stat_used() { return 1; }
+
+#pragma omp end declare target
+
+int main() {
+ int res;
+#pragma omp target map(from \
+ : res)
+ res = bar() + baz() + call();
+ return res;
+}
+
+int test() { return 3; }
+static int stat_unused_() { return 4; }
+static int stat_used_() { return 5; }
+
+#pragma omp declare target
+
+struct SpecialFuncs {
+ void vd() {}
+ SpecialFuncs();
+ ~SpecialFuncs();
+
+ int method_() { return 6; }
+#pragma omp declare variant(SpecialFuncs::method_) \
+ match(device = {isa("nvptx64")})
+ int method() { return 1; }
+#pragma omp declare variant(SpecialFuncs::method_) \
+ match(device = {isa("nvptx64")})
+ int Method();
+} s;
+
+int SpecialFuncs::Method() { return 1; }
+
+struct SpecSpecialFuncs {
+ void vd() {}
+ SpecSpecialFuncs();
+ ~SpecSpecialFuncs();
+
+ int method_();
+#pragma omp declare variant(SpecSpecialFuncs::method_) \
+ match(device = {isa("nvptx64")})
+ int method() { return 1; }
+#pragma omp declare variant(SpecSpecialFuncs::method_) \
+ match(device = {isa("nvptx64")})
+ int Method();
+} s1;
+
+#pragma omp end declare target
+
+int SpecSpecialFuncs::method_() { return 7; }
+int SpecSpecialFuncs::Method() { return 1; }
+
+int prio() { return 81; }
+int prio1() { return 82; }
+static int prio2() { return 83; }
+static int prio3() { return 84; }
+static int prio4() { return 84; }
+int fn_linkage_variant() { return 85; }
+extern "C" int fn_linkage_variant1() { return 86; }
+int fn_variant2() { return 1; }
+
+#pragma omp declare target
+
+void xxx() {
+ (void)s.method();
+ (void)s1.method();
+}
+
+#pragma omp declare variant(prio) match(device = {isa("nvptx64"), kind(gpu)})
+#pragma omp declare variant(prio1) match(device = {isa("nvptx64", "nvptx64")})
+int prio_() { return 1; }
+
+#pragma omp declare variant(prio4) match(device = {isa("nvptx64"), kind(gpu)})
+#pragma omp declare variant(prio2) match(device = {isa("nvptx64")})
+#pragma omp declare variant(prio3) match(device = {isa("nvptx64"), kind(gpu)})
+static int prio1_() { return 1; }
+
+int int_fn() { return prio1_(); }
+
+extern "C" {
+#pragma omp declare variant(fn_linkage_variant) match(device = {isa("nvptx64")})
+int fn_linkage() { return 1; }
+}
+
+#pragma omp declare variant(fn_linkage_variant1) match(device = {isa("nvptx64")})
+int fn_linkage1() { return 1; }
+
+#pragma omp declare variant(fn_variant2) match(device = {isa("amdgcn")})
+int fn2() { return 87; }
+
+#pragma omp end declare target
+
+#endif // HEADER
Index: clang/test/OpenMP/declare_variant_mixed_codegen.cpp
===================================================================
--- clang/test/OpenMP/declare_variant_mixed_codegen.cpp
+++ clang/test/OpenMP/declare_variant_mixed_codegen.cpp
@@ -1,6 +1,12 @@
-// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope | FileCheck %s
-// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple x86_64-unknown-linux -fexceptions -fcxx-exceptions -emit-pch -o %t -fopenmp-version=50 %s
-// RUN: %clang_cc1 -fopenmp -x c++ -triple x86_64-unknown-linux -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=50 | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope -DARCH=\"x86_64\"| FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple x86_64-unknown-linux -fexceptions -fcxx-exceptions -emit-pch -o %t -fopenmp-version=50 %s -DARCH=\"x86_64\"
+// RUN: %clang_cc1 -fopenmp -x c++ -triple x86_64-unknown-linux -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=50 -DARCH=\"x86_64\" | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple aarch64-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope -DARCH=\"aarch64\" | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple aarch64-unknown-linux -fexceptions -fcxx-exceptions -emit-pch -o %t -fopenmp-version=50 %s -DARCH=\"aarch64\"
+// RUN: %clang_cc1 -fopenmp -x c++ -triple aarch64-unknown-linux -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=50 -DARCH=\"aarch64\" | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple ppc64le-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope -DARCH=\"ppc64le\" | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple ppc64le-unknown-linux -fexceptions -fcxx-exceptions -emit-pch -o %t -fopenmp-version=50 %s -DARCH=\"ppc64le\"
+// RUN: %clang_cc1 -fopenmp -x c++ -triple ppc64le-unknown-linux -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=50 -DARCH=\"ppc64le\" | FileCheck %s
// expected-no-diagnostics
// CHECK-NOT: ret i32 {{1|4|81|84}}
@@ -34,23 +40,23 @@
int foo() { return 2; }
-#pragma omp declare variant(foo) match(implementation = {vendor(llvm)}, device={kind(cpu)})
+#pragma omp declare variant(foo) match(implementation = {vendor(llvm)}, device={kind(cpu), isa(ARCH)})
int bar() { return 1; }
int bazzz();
-#pragma omp declare variant(bazzz) match(implementation = {vendor(llvm)}, device={kind(host)})
+#pragma omp declare variant(bazzz) match(implementation = {vendor(llvm)}, device={kind(host), isa(ARCH)})
int baz() { return 1; }
int test();
-#pragma omp declare variant(test) match(implementation = {vendor(llvm)}, device={kind(cpu)})
+#pragma omp declare variant(test) match(implementation = {vendor(llvm)}, device={kind(cpu), isa(ARCH)})
int call() { return 1; }
static int stat_unused_();
-#pragma omp declare variant(stat_unused_) match(implementation = {vendor(llvm)}, device={kind(cpu)})
+#pragma omp declare variant(stat_unused_) match(implementation = {vendor(llvm)}, device={kind(cpu), isa(ARCH)})
static int stat_unused() { return 1; }
static int stat_used_();
-#pragma omp declare variant(stat_used_) match(implementation = {vendor(llvm)}, device={kind(host)})
+#pragma omp declare variant(stat_used_) match(implementation = {vendor(llvm)}, device={kind(host), isa(ARCH)})
static int stat_used() { return 1; }
int main() { return bar() + baz() + call() + stat_used(); }
@@ -66,10 +72,10 @@
int method_() { return 6; }
#pragma omp declare variant(SpecialFuncs::method_) \
- match(implementation = {vendor(llvm)}, device={kind(cpu)})
+ match(implementation = {vendor(llvm)}, device={kind(cpu), isa(ARCH)})
int method() { return 1; }
#pragma omp declare variant(SpecialFuncs::method_) \
- match(implementation = {vendor(llvm)}, device={kind(host)})
+ match(implementation = {vendor(llvm)}, device={kind(host), isa(ARCH)})
int Method();
} s;
@@ -82,10 +88,10 @@
int method_();
#pragma omp declare variant(SpecSpecialFuncs::method_) \
- match(implementation = {vendor(llvm)}, device={kind(cpu)})
+ match(implementation = {vendor(llvm)}, device={kind(cpu), isa(ARCH)})
int method() { return 1; }
#pragma omp declare variant(SpecSpecialFuncs::method_) \
- match(implementation = {vendor(llvm)}, device={kind(host)})
+ match(implementation = {vendor(llvm)}, device={kind(host), isa(ARCH)})
int Method();
} s1;
@@ -100,29 +106,29 @@
int prio() { return 81; }
int prio1() { return 82; }
-#pragma omp declare variant(prio) match(implementation = {vendor(score(2): llvm)}, device={kind(cpu,host)})
-#pragma omp declare variant(prio1) match(implementation = {vendor(score(1): llvm)}, device={kind(cpu)})
+#pragma omp declare variant(prio) match(implementation = {vendor(score(2): llvm)}, device={kind(cpu,host), isa(ARCH)})
+#pragma omp declare variant(prio1) match(implementation = {vendor(score(1): llvm)}, device={kind(cpu), isa(ARCH)})
int prio_() { return 1; }
static int prio2() { return 83; }
static int prio3() { return 84; }
static int prio4() { return 84; }
-#pragma omp declare variant(prio4) match(implementation = {vendor(score(8): llvm)},device={kind(cpu,host)})
-#pragma omp declare variant(prio2) match(implementation = {vendor(score(5): llvm)})
-#pragma omp declare variant(prio3) match(implementation = {vendor(score(7): llvm)}, device={kind(cpu)})
+#pragma omp declare variant(prio4) match(implementation = {vendor(score(8): llvm)},device={kind(cpu,host), isa(ARCH)})
+#pragma omp declare variant(prio2) match(implementation = {vendor(score(5): llvm)}, device={isa(ARCH)})
+#pragma omp declare variant(prio3) match(implementation = {vendor(score(7): llvm)}, device={kind(cpu), isa("i586")})
static int prio1_() { return 1; }
int int_fn() { return prio1_(); }
int fn_linkage_variant() { return 85; }
extern "C" {
-#pragma omp declare variant(fn_linkage_variant) match(implementation = {vendor(llvm)}, device={kind(cpu)})
+#pragma omp declare variant(fn_linkage_variant) match(implementation = {vendor(llvm)}, device={kind(cpu), isa(ARCH)})
int fn_linkage() { return 1; }
}
extern "C" int fn_linkage_variant1() { return 86; }
-#pragma omp declare variant(fn_linkage_variant1) match(implementation = {vendor(llvm)}, device={kind(host)})
+#pragma omp declare variant(fn_linkage_variant1) match(implementation = {vendor(llvm)}, device={kind(host), isa(ARCH)})
int fn_linkage1() { return 1; }
int fn_variant2() { return 1; }
@@ -132,6 +138,7 @@
#pragma omp declare variant(fn_variant2) match(implementation = {vendor(llvm)}, device={kind(cpu,nohost)})
#pragma omp declare variant(fn_variant2) match(implementation = {vendor(llvm)}, device={kind(gpu)})
#pragma omp declare variant(fn_variant2) match(implementation = {vendor(llvm)}, device={kind(fpga)})
+#pragma omp declare variant(fn_variant2) match(implementation = {vendor(llvm)}, device={isa("mips")})
int fn2() { return 87; }
#endif // HEADER
Index: clang/test/OpenMP/declare_variant_messages.cpp
===================================================================
--- clang/test/OpenMP/declare_variant_messages.cpp
+++ clang/test/OpenMP/declare_variant_messages.cpp
@@ -89,6 +89,16 @@
#pragma omp declare variant(foofoo <int>) match(device={kind(score(foofoo <int>()) ibm)}) // expected-error {{expected ')' or ',' after 'score'}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
#pragma omp declare variant(foofoo <int>) match(device={kind(score(C+5): host), kind(llvm)}) // expected-error {{context trait selector 'kind' is used already in the same 'device' context selector set of 'omp declare variant' directive}} expected-note {{previously context trait selector 'kind' used here}} expected-error {{expected ')' or ',' after 'score'}} expected-note {{to match this '('}} expected-error {{expected ')'}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}} expected-error {{unknown 'llvm' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
#pragma omp declare variant(foofoo <int>) match(device={kind(score(C+5): nohost), vendor(llvm)}) // expected-warning {{unknown context selector in 'device' context selector set of 'omp declare variant' directive, ignored}} expected-error {{expected ')' or ',' after 'score'}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
+#pragma omp declare variant(foofoo <int>) match(device={xxx}) // expected-warning {{unknown context selector in 'device' context selector set of 'omp declare variant' directive, ignored}}
+#pragma omp declare variant(foofoo <int>) match(device={isa}) // expected-error {{expected '(' after 'isa'}} expected-error {{expected ')'}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+#pragma omp declare variant(foofoo <int>) match(device={isa(}) // expected-error 2 {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+#pragma omp declare variant(foofoo <int>) match(device={isa()}) // expected-error {{expected expression}}
+#pragma omp declare variant(foofoo <int>) match(device={isa(score "x86_64")}) // expected-error {{use of undeclared identifier 'score'}} expected-error {{expected ')' or ',' after 'instruction set architecture'}}
+#pragma omp declare variant(foofoo <int>) match(device={isa(score( "powerpc64le")}) // expected-error 3 {{expected ')'}} expected-note {{to match this '('}} expected-error {{use of undeclared identifier 'score'}} expected-error {{expected expression}}
+#pragma omp declare variant(foofoo <int>) match(device={isa(score(C "nvptx")}) // expected-error 4 {{expected ')'}} expected-note 2 {{to match this '('}} expected-error {{expected expression}}
+#pragma omp declare variant(foofoo <int>) match(device={isa(score(foofoo <int>()) "arm")}) // expected-error {{expected ')'}} expected-error {{use of undeclared identifier 'score'}}
+#pragma omp declare variant(foofoo <int>) match(device={isa(score(C+5): "i586"), isa(llvm)}) // expected-error {{context trait selector 'isa' is used already in the same 'device' context selector set of 'omp declare variant' directive}} expected-note {{previously context trait selector 'isa' used here}} expected-error 3 {{expected ')'}} expected-error {{use of undeclared identifier 'score'}} expected-error {{use of undeclared identifier 'llvm'}} expected-error {{expected expression}} expected-note {{to match this '('}}
+#pragma omp declare variant(foofoo <int>) match(device={isa(score(C+5): "aarch64"), vendor(llvm)}) // expected-warning {{unknown context selector in 'device' context selector set of 'omp declare variant' directive, ignored}} expected-error 2 {{expected ')'}} expected-error {{use of undeclared identifier 'score'}} expected-error {{expected ')'}} expected-error {{expected expression}} expected-note {{to match this '('}}
template <typename T, int C>
T barbar();
Index: clang/test/OpenMP/declare_variant_messages.c
===================================================================
--- clang/test/OpenMP/declare_variant_messages.c
+++ clang/test/OpenMP/declare_variant_messages.c
@@ -1,6 +1,6 @@
-// RUN: %clang_cc1 -triple=x86_64-pc-win32 -verify -fopenmp -x c -std=c99 -fms-extensions -Wno-pragma-pack %s
+// RUN: %clang_cc1 -triple=x86_64-pc-win32 -verify -fopenmp -x c -std=c99 -fms-extensions -Wno-pragma-pack %s -Wno-implicit-function-declaration
-// RUN: %clang_cc1 -triple=x86_64-pc-win32 -verify -fopenmp-simd -x c -std=c99 -fms-extensions -Wno-pragma-pack %s
+// RUN: %clang_cc1 -triple=x86_64-pc-win32 -verify -fopenmp-simd -x c -std=c99 -fms-extensions -Wno-pragma-pack %s -Wno-implicit-function-declaration
// expected-error@+1 {{expected an OpenMP directive}}
#pragma omp declare
@@ -45,7 +45,17 @@
#pragma omp declare variant(foo) match(device={kind(score(2 gpu)}) // expected-error 2 {{expected ')'}} expected-note {{to match this '('}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
#pragma omp declare variant(foo) match(device={kind(score(foo()) ibm)}) // expected-error {{expected ')' or ',' after 'score'}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
#pragma omp declare variant(foo) match(device={kind(score(5): host), kind(llvm)}) // expected-error {{context trait selector 'kind' is used already in the same 'device' context selector set of 'omp declare variant' directive}} expected-note {{previously context trait selector 'kind' used here}} expected-error {{expected ')' or ',' after 'score'}} expected-note {{to match this '('}} expected-error {{expected ')'}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}} expected-error {{unknown 'llvm' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
-#pragma omp declare variant(foo) match(device={kind(score(5): nohost), vendor(llvm)}) // expected-warning {{unknown context selector in 'device' context selector set of 'omp declare variant' directive, ignored}} expected-error {{expected ')' or ',' after 'score'}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}}
+#pragma omp declare variant(foo) match(device={kind(score(5): nohost), vendor(llvm)}) // expected-warning {{unknown context selector in 'device' context selector set of 'omp declare variant' directive, ignored}} expected-error {{expected ')' or ',' after 'score'}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
+#pragma omp declare variant(foo) match(device={xxx}) // expected-warning {{unknown context selector in 'device' context selector set of 'omp declare variant' directive, ignored}}
+#pragma omp declare variant(foo) match(device={isa}) // expected-error {{expected '(' after 'isa'}} expected-error {{expected ')'}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+#pragma omp declare variant(foo) match(device={isa(}) // expected-error 2 {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+#pragma omp declare variant(foo) match(device={isa()}) // expected-error {{expected expression}}}
+#pragma omp declare variant(foo) match(device={isa(score "i586")}) // expected-error {{use of undeclared identifier 'score'}} expected-error {{expected ')' or ',' after 'instruction set architecture'}}
+#pragma omp declare variant(foo) match(device={isa(score( "x86_64")}) // expected-error 3 {{expected ')'}} expected-note {{to match this '('}} expected-error {{trait expression must have a string type, not 'int'}} expected-error {{trait expression must be a constant string expression}} expected-error {{expected expression}}
+#pragma omp declare variant(foo) match(device={isa(score(2 "arm")}) // expected-error 4 {{expected ')'}} expected-note 2 {{to match this '('}} expected-error {{expected expression}}
+#pragma omp declare variant(foo) match(device={isa(score(foo()) "powerpc64le")}) // expected-error {{expected ')'}} expected-error {{trait expression must have a string type, not 'int'}} expected-error {{trait expression must be a constant string expression}}
+#pragma omp declare variant(foo) match(device={isa(score(5): "x86"), isa(llvm)}) // expected-error {{context trait selector 'isa' is used already in the same 'device' context selector set of 'omp declare variant' directive}} expected-note {{previously context trait selector 'isa' used here}} expected-note {{to match this '('}} expected-error 3 {{expected ')'}} expected-error {{use of undeclared identifier 'llvm'}} expected-error {{trait expression must have a string type, not 'int'}} expected-error {{trait expression must be a constant string expression}} expected-error {{expected expression}}
+#pragma omp declare variant(foo) match(device={isa(score(5): "x86"), vendor(llvm)}) // expected-warning {{unknown context selector in 'device' context selector set of 'omp declare variant' directive, ignored}} expected-error 3 {{expected ')'}} expected-note {{to match this '('}} expected-error {{trait expression must have a string type, not 'int'}} expected-error {{trait expression must be a constant string expression}} expected-error {{expected expression}}
int bar(void);
// expected-error@+2 {{'#pragma omp declare variant' can only be applied to functions}}
Index: clang/test/OpenMP/declare_variant_device_isa_codegen.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/declare_variant_device_isa_codegen.cpp
@@ -0,0 +1,152 @@
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope -DCORRECT=\"x86_64\" | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple x86_64-unknown-linux -fexceptions -fcxx-exceptions -emit-pch -o %t -fopenmp-version=50 %s -DCORRECT=\"x86_64\"
+// RUN: %clang_cc1 -fopenmp -x c++ -triple x86_64-unknown-linux -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=50 -DCORRECT=\"x86_64\" | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple aarch64-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope -DCORRECT=\"aarch64\" | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple aarch64-unknown-linux -fexceptions -fcxx-exceptions -emit-pch -o %t -fopenmp-version=50 %s -DCORRECT=\"aarch64\"
+// RUN: %clang_cc1 -fopenmp -x c++ -triple aarch64-unknown-linux -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=50 -DCORRECT=\"aarch64\" | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple ppc64le-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope -DCORRECT=\"powerpc64le\" | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple ppc64le-unknown-linux -fexceptions -fcxx-exceptions -emit-pch -o %t -fopenmp-version=50 %s -DCORRECT=\"ppc64le\"
+// RUN: %clang_cc1 -fopenmp -x c++ -triple ppc64le-unknown-linux -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=50 -DCORRECT=\"ppc64le\" | FileCheck %s
+
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -fopenmp-targets=x86_64-unknown-linux -emit-llvm-bc %s -o %t-host.bc -DCORRECT=\"x86_64\"
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -o - -DCORRECT=\"x86_64\" | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -emit-pch -o %t -DCORRECT=\"x86_64\"
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -include-pch %t -o - -DCORRECT=\"x86_64\" | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple ppc64le-unknown-linux -fopenmp-targets=ppc64le-unknown-linux -emit-llvm-bc %s -o %t-host.bc -DCORRECT=\"ppc64le\"
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple ppc64le-unknown-linux -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -o - -DCORRECT=\"ppc64le\" | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple ppc64le-unknown-linux -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -emit-pch -o %t -DCORRECT=\"ppc64le\"
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple ppc64le-unknown-linux -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -include-pch %t -o - -DCORRECT=\"ppc64le\" | FileCheck %s
+
+// expected-no-diagnostics
+
+// CHECK-NOT: ret i32 {{1|4|81|84}}
+// CHECK-DAG: @_Z3barv = {{.*}}alias i32 (), i32 ()* @_Z3foov
+// CHECK-DAG: @_ZN16SpecSpecialFuncs6MethodEv = {{.*}}alias i32 (%struct.SpecSpecialFuncs*), i32 (%struct.SpecSpecialFuncs*)* @_ZN16SpecSpecialFuncs7method_Ev
+// CHECK-DAG: @_ZN16SpecSpecialFuncs6methodEv = linkonce_odr {{.*}}alias i32 (%struct.SpecSpecialFuncs*), i32 (%struct.SpecSpecialFuncs*)* @_ZN16SpecSpecialFuncs7method_Ev
+// CHECK-DAG: @_ZN12SpecialFuncs6methodEv = linkonce_odr {{.*}}alias i32 (%struct.SpecialFuncs*), i32 (%struct.SpecialFuncs*)* @_ZN12SpecialFuncs7method_Ev
+// CHECK-DAG: @_Z5prio_v = {{.*}}alias i32 (), i32 ()* @_Z5prio1v
+// CHECK-DAG: @_ZL6prio1_v = internal alias i32 (), i32 ()* @_ZL5prio2v
+// CHECK-DAG: @_Z4callv = {{.*}}alias i32 (), i32 ()* @_Z4testv
+// CHECK-DAG: @_ZL9stat_usedv = internal alias i32 (), i32 ()* @_ZL10stat_used_v
+// CHECK-DAG: @_ZN12SpecialFuncs6MethodEv = {{.*}}alias i32 (%struct.SpecialFuncs*), i32 (%struct.SpecialFuncs*)* @_ZN12SpecialFuncs7method_Ev
+// CHECK-DAG: @fn_linkage = {{.*}}alias i32 (), i32 ()* @_Z18fn_linkage_variantv
+// CHECK-DAG: @_Z11fn_linkage1v = {{.*}}alias i32 (), i32 ()* @fn_linkage_variant1
+// CHECK-DAG: declare {{.*}}i32 @_Z5bazzzv()
+// CHECK-DAG: declare {{.*}}i32 @_Z3bazv()
+// CHECK-DAG: ret i32 2
+// CHECK-DAG: ret i32 3
+// CHECK-DAG: ret i32 5
+// CHECK-DAG: ret i32 6
+// CHECK-DAG: ret i32 7
+// CHECK-DAG: ret i32 82
+// CHECK-DAG: ret i32 83
+// CHECK-DAG: ret i32 85
+// CHECK-DAG: ret i32 86
+// CHECK-DAG: ret i32 87
+// CHECK-NOT: ret i32 {{1|4|81|84}}
+
+#ifndef HEADER
+#define HEADER
+
+#pragma omp declare target
+#define WRONG "mips"
+
+int foo() { return 2; }
+
+#pragma omp declare variant(foo) match(device = {isa(CORRECT)})
+int bar() { return 1; }
+
+int bazzz();
+#pragma omp declare variant(bazzz) match(device = {isa(CORRECT)})
+int baz() { return 1; }
+
+int test();
+#pragma omp declare variant(test) match(device = {isa(CORRECT)})
+int call() { return 1; }
+
+static int stat_unused_();
+#pragma omp declare variant(stat_unused_) match(device = {isa(CORRECT)})
+static int stat_unused() { return 1; }
+
+static int stat_used_();
+#pragma omp declare variant(stat_used_) match(device = {isa(CORRECT)})
+static int stat_used() { return 1; }
+
+int main() { return bar() + baz() + call() + stat_used(); }
+
+int test() { return 3; }
+static int stat_unused_() { return 4; }
+static int stat_used_() { return 5; }
+
+struct SpecialFuncs {
+ void vd() {}
+ SpecialFuncs();
+ ~SpecialFuncs();
+
+ int method_() { return 6; }
+#pragma omp declare variant(SpecialFuncs::method_) \
+ match(device = {isa(CORRECT)})
+ int method() { return 1; }
+#pragma omp declare variant(SpecialFuncs::method_) \
+ match(device = {isa(CORRECT)})
+ int Method();
+} s;
+
+int SpecialFuncs::Method() { return 1; }
+
+struct SpecSpecialFuncs {
+ void vd() {}
+ SpecSpecialFuncs();
+ ~SpecSpecialFuncs();
+
+ int method_();
+#pragma omp declare variant(SpecSpecialFuncs::method_) \
+ match(device = {isa(CORRECT)})
+ int method() { return 1; }
+#pragma omp declare variant(SpecSpecialFuncs::method_) \
+ match(device = {isa(CORRECT)})
+ int Method();
+} s1;
+
+int SpecSpecialFuncs::method_() { return 7; }
+int SpecSpecialFuncs::Method() { return 1; }
+
+void xxx() {
+ (void)s.method();
+ (void)s1.method();
+}
+
+int prio() { return 81; }
+int prio1() { return 82; }
+
+#pragma omp declare variant(prio) match(device = {isa(CORRECT), kind(cpu)})
+#pragma omp declare variant(prio1) match(device = {isa(CORRECT)})
+int prio_() { return 1; }
+
+static int prio2() { return 83; }
+static int prio3() { return 84; }
+static int prio4() { return 84; }
+
+#pragma omp declare variant(prio4) match(device = {isa(CORRECT),kind(cpu)})
+#pragma omp declare variant(prio2) match(device = {isa(CORRECT)})
+#pragma omp declare variant(prio3) match(device = {isa(CORRECT), kind(nohost)})
+static int prio1_() { return 1; }
+
+int int_fn() { return prio1_(); }
+
+int fn_linkage_variant() { return 85; }
+extern "C" {
+#pragma omp declare variant(fn_linkage_variant) match(device = {isa(CORRECT)})
+int fn_linkage() { return 1; }
+}
+
+extern "C" int fn_linkage_variant1() { return 86; }
+#pragma omp declare variant(fn_linkage_variant1) match(device = {isa(CORRECT)})
+int fn_linkage1() { return 1; }
+
+int fn_variant2() { return 1; }
+#pragma omp declare variant(fn_variant2) match(device = {isa(WRONG)})
+int fn2() { return 87; }
+
+#pragma omp end declare target
+#endif // HEADER
Index: clang/test/OpenMP/declare_variant_ast_print.cpp
===================================================================
--- clang/test/OpenMP/declare_variant_ast_print.cpp
+++ clang/test/OpenMP/declare_variant_ast_print.cpp
@@ -17,20 +17,22 @@
// CHECK-NEXT: return int();
// CHECK-NEXT: }
-// CHECK: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(5):ibm)},device={kind(fpga)})
+// CHECK: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(5):ibm)},device={kind(fpga),isa("ppc64le")})
// CHECK-NEXT: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(0):unknown)})
-// CHECK-NEXT: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(0):llvm)},device={kind(cpu)})
+// CHECK-NEXT: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(0):llvm)},device={kind(cpu),isa("x86_64", "arm")})
// CHECK-NEXT: int bar();
#pragma omp declare variant(foofoo <int>) match(xxx = {})
#pragma omp declare variant(foofoo <int>) match(xxx = {vvv})
-#pragma omp declare variant(foofoo <int>) match(implementation={vendor(llvm), xxx}, device={kind(cpu)})
+#pragma omp declare variant(foofoo <int>) match(implementation={vendor(llvm), xxx}, device={kind(cpu), isa("x86""_64", "arm")})
#pragma omp declare variant(foofoo <int>) match(implementation={vendor(unknown)})
-#pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(5): ibm)}, device={kind(fpga)})
+#pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(5): ibm)}, device={kind(fpga), isa("ppc64le")})
int bar();
-// CHECK: #pragma omp declare variant(foofoo<T>) match(implementation={vendor(score(C + 5):ibm, xxx)},device={kind(cpu, host)})
+constexpr auto Arch = "x86""_64";
+
+// CHECK: #pragma omp declare variant(foofoo<T>) match(implementation={vendor(score(C + 5):ibm, xxx)},device={kind(cpu, host),isa("ppc64le")})
// CHECK-NEXT: #pragma omp declare variant(foofoo<T>) match(implementation={vendor(score(0):unknown)})
-// CHECK-NEXT: #pragma omp declare variant(foofoo<T>) match(implementation={vendor(score(0):llvm)},device={kind(cpu)})
+// CHECK-NEXT: #pragma omp declare variant(foofoo<T>) match(implementation={vendor(score(0):llvm)},device={kind(cpu),isa("x86_64", "arm")})
// CHECK-NEXT: template <typename T, int C> T barbar();
#pragma omp declare variant(foofoo <T>) match(xxx = {})
#pragma omp declare variant(foofoo <T>) match(xxx = {vvv})
@@ -38,15 +40,15 @@
#pragma omp declare variant(foofoo <T>) match(user = {score(<expr>) : condition(<expr>)})
#pragma omp declare variant(foofoo <T>) match(user = {condition(<expr>)})
#pragma omp declare variant(foofoo <T>) match(user = {condition(<expr>)})
-#pragma omp declare variant(foofoo <T>) match(implementation={vendor(llvm)},device={kind(cpu)})
+#pragma omp declare variant(foofoo <T>) match(implementation={vendor(llvm)},device={kind(cpu), isa(Arch, "arm")})
#pragma omp declare variant(foofoo <T>) match(implementation={vendor(unknown)})
-#pragma omp declare variant(foofoo <T>) match(implementation={vendor(score(C+5): ibm, xxx, ibm)},device={kind(cpu,host)})
+#pragma omp declare variant(foofoo <T>) match(implementation={vendor(score(C+5): ibm, xxx, ibm)},device={kind(cpu,host), isa("ppc64le")})
template <typename T, int C>
T barbar();
-// CHECK: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(3 + 5):ibm, xxx)},device={kind(cpu, host)})
+// CHECK: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(3 + 5):ibm, xxx)},device={kind(cpu, host),isa("ppc64le")})
// CHECK-NEXT: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(0):unknown)})
-// CHECK-NEXT: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(0):llvm)},device={kind(cpu)})
+// CHECK-NEXT: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(0):llvm)},device={kind(cpu),isa("x86_64", "arm")})
// CHECK-NEXT: template<> int barbar<int, 3>();
// CHECK-NEXT: int baz() {
@@ -66,19 +68,19 @@
void h_ref(C *hp, C *hp2, C *hq, C *lin) {
}
-// CHECK: #pragma omp declare variant(h_ref<C>) match(implementation={vendor(score(0):unknown)},device={kind(nohost)})
-// CHECK-NEXT: #pragma omp declare variant(h_ref<C>) match(implementation={vendor(score(0):llvm)},device={kind(gpu)})
+// CHECK: #pragma omp declare variant(h_ref<C>) match(implementation={vendor(score(0):unknown)},device={kind(nohost),isa("x86", "aarch64")})
+// CHECK-NEXT: #pragma omp declare variant(h_ref<C>) match(implementation={vendor(score(0):llvm)},device={kind(gpu),isa("ppc64")})
// CHECK-NEXT: template <class C> void h(C *hp, C *hp2, C *hq, C *lin) {
// CHECK-NEXT: }
#pragma omp declare variant(h_ref <C>) match(xxx = {})
-#pragma omp declare variant(h_ref <C>) match(implementation={vendor(llvm)}, device={kind(gpu)})
-#pragma omp declare variant(h_ref <C>) match(implementation={vendor(unknown)},device={kind(nohost)})
+#pragma omp declare variant(h_ref <C>) match(implementation={vendor(llvm)}, device={kind(gpu),isa("ppc64")})
+#pragma omp declare variant(h_ref <C>) match(implementation={vendor(unknown)},device={kind(nohost), isa("x86","aarch64")})
template <class C>
void h(C *hp, C *hp2, C *hq, C *lin) {
}
-// CHECK: #pragma omp declare variant(h_ref<float>) match(implementation={vendor(score(0):unknown)},device={kind(nohost)})
-// CHECK-NEXT: #pragma omp declare variant(h_ref<float>) match(implementation={vendor(score(0):llvm)},device={kind(gpu)})
+// CHECK: #pragma omp declare variant(h_ref<float>) match(implementation={vendor(score(0):unknown)},device={kind(nohost),isa("x86", "aarch64")})
+// CHECK-NEXT: #pragma omp declare variant(h_ref<float>) match(implementation={vendor(score(0):llvm)},device={kind(gpu),isa("ppc64")})
// CHECK-NEXT: template<> void h<float>(float *hp, float *hp2, float *hq, float *lin) {
// CHECK-NEXT: }
@@ -86,7 +88,7 @@
// CHECK-NEXT: h((float *)hp, (float *)hp2, (float *)hq, (float *)lin);
// CHECK-NEXT: }
#pragma omp declare variant(h_ref <double>) match(xxx = {})
-#pragma omp declare variant(h_ref <double>) match(implementation={vendor(ibm)},device={kind(cpu,gpu)})
+#pragma omp declare variant(h_ref <double>) match(implementation={vendor(ibm)},device={kind(cpu,gpu),isa("x86", "aarch64")})
#pragma omp declare variant(h_ref <double>) match(implementation={vendor(unknown)})
template <>
void h(double *hp, double *hp2, double *hq, double *lin) {
@@ -97,36 +99,36 @@
int fn();
// CHECK: int fn(int);
int fn(int);
-// CHECK: #pragma omp declare variant(fn) match(implementation={vendor(score(0):unknown)},device={kind(cpu, gpu)})
+// CHECK: #pragma omp declare variant(fn) match(implementation={vendor(score(0):unknown)},device={kind(cpu, gpu),isa("x86", "aarch64")})
// CHECK-NEXT: #pragma omp declare variant(fn) match(implementation={vendor(score(0):llvm)})
// CHECK-NEXT: int overload();
#pragma omp declare variant(fn) match(xxx = {})
#pragma omp declare variant(fn) match(implementation={vendor(llvm)})
-#pragma omp declare variant(fn) match(implementation={vendor(unknown)},device={kind(cpu,gpu)})
+#pragma omp declare variant(fn) match(implementation={vendor(unknown)},device={kind(cpu,gpu),isa("x86", "aarch64")})
int overload(void);
// CHECK: int fn_deduced_variant() {
// CHECK-NEXT: return 0;
// CHECK-NEXT: }
auto fn_deduced_variant() { return 0; }
-// CHECK: #pragma omp declare variant(fn_deduced_variant) match(implementation={vendor(score(0):unknown)},device={kind(gpu, nohost)})
-// CHECK-NEXT: #pragma omp declare variant(fn_deduced_variant) match(implementation={vendor(score(0):llvm)},device={kind(cpu, host)})
+// CHECK: #pragma omp declare variant(fn_deduced_variant) match(implementation={vendor(score(0):unknown)},device={kind(gpu, nohost),isa("ppc64le", "aarch64")})
+// CHECK-NEXT: #pragma omp declare variant(fn_deduced_variant) match(implementation={vendor(score(0):llvm)},device={kind(cpu, host),isa("x86", "aarch64")})
// CHECK-NEXT: int fn_deduced();
#pragma omp declare variant(fn_deduced_variant) match(xxx = {})
-#pragma omp declare variant(fn_deduced_variant) match(implementation={vendor(llvm)},device={kind(cpu,host)})
-#pragma omp declare variant(fn_deduced_variant) match(implementation={vendor(unknown)},device={kind(gpu,nohost)})
+#pragma omp declare variant(fn_deduced_variant) match(implementation={vendor(llvm)},device={kind(cpu,host),isa("x86", "aarch64")})
+#pragma omp declare variant(fn_deduced_variant) match(implementation={vendor(unknown)},device={kind(gpu,nohost),isa("ppc64le", "aarch64")})
int fn_deduced();
// CHECK: int fn_deduced_variant1();
int fn_deduced_variant1();
-// CHECK: #pragma omp declare variant(fn_deduced_variant1) match(implementation={vendor(score(0):unknown)},device={kind(cpu, host)})
-// CHECK-NEXT: #pragma omp declare variant(fn_deduced_variant1) match(implementation={vendor(score(0):ibm)},device={kind(gpu, nohost)})
+// CHECK: #pragma omp declare variant(fn_deduced_variant1) match(implementation={vendor(score(0):unknown)},device={kind(cpu, host),isa("ppc64le", "aarch64")})
+// CHECK-NEXT: #pragma omp declare variant(fn_deduced_variant1) match(implementation={vendor(score(0):ibm)},device={kind(gpu, nohost),isa("x86", "aarch64")})
// CHECK-NEXT: int fn_deduced1() {
// CHECK-NEXT: return 0;
// CHECK-NEXT: }
#pragma omp declare variant(fn_deduced_variant1) match(xxx = {})
-#pragma omp declare variant(fn_deduced_variant1) match(implementation={vendor(ibm)},device={kind(gpu,nohost)})
-#pragma omp declare variant(fn_deduced_variant1) match(implementation={vendor(unknown)},device={kind(cpu,host)})
+#pragma omp declare variant(fn_deduced_variant1) match(implementation={vendor(ibm)},device={kind(gpu,nohost),isa("x86", "aarch64")})
+#pragma omp declare variant(fn_deduced_variant1) match(implementation={vendor(unknown)},device={kind(cpu,host),isa("ppc64le", "aarch64")})
auto fn_deduced1() { return 0; }
// CHECK: struct SpecialFuncs {
@@ -140,11 +142,11 @@
// CHECK-NEXT: }
// CHECK-NEXT: void bar(int) {
// CHECK-NEXT: }
-// CHECK-NEXT: #pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(score(0):unknown)},device={kind(nohost)})
-// CHECK-NEXT: #pragma omp declare variant(SpecialFuncs::bar) match(implementation={vendor(score(0):ibm)},device={kind(cpu)})
+// CHECK-NEXT: #pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(score(0):unknown)},device={kind(nohost),isa("ppc64le", "aarch64")})
+// CHECK-NEXT: #pragma omp declare variant(SpecialFuncs::bar) match(implementation={vendor(score(0):ibm)},device={kind(cpu),isa("x86", "aarch64")})
// CHECK-NEXT: void foo1() {
// CHECK-NEXT: }
-// CHECK-NEXT: #pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(score(0):unknown)},device={kind(cpu, host)})
+// CHECK-NEXT: #pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(score(0):unknown)},device={kind(cpu, host),isa("x86_64")})
// CHECK-NEXT: void xxx();
// CHECK-NEXT: } s;
struct SpecialFuncs {
@@ -157,14 +159,14 @@
void bar(int) {}
#pragma omp declare variant(SpecialFuncs::baz) match(xxx = {})
#pragma omp declare variant(SpecialFuncs::bar) match(xxx = {})
-#pragma omp declare variant(SpecialFuncs::bar) match(implementation={vendor(ibm)},device={kind(cpu)})
-#pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(unknown)},device={kind(nohost)})
+#pragma omp declare variant(SpecialFuncs::bar) match(implementation={vendor(ibm)},device={kind(cpu),isa("x86", "aarch64")})
+#pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(unknown)},device={kind(nohost),isa("ppc64le", "aarch64")})
void foo1() {}
-#pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(unknown)},device={kind(cpu, host)})
+#pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(unknown)},device={kind(cpu, host),isa("x86_64")})
void xxx();
} s;
-// CHECK: #pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(score(0):unknown)},device={kind(cpu, host)})
+// CHECK: #pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(score(0):unknown)},device={kind(cpu, host),isa("x86_64")})
// CHECK-NEXT: void SpecialFuncs::xxx() {
// CHECK-NEXT: }
void SpecialFuncs::xxx() {}
@@ -173,11 +175,11 @@
// CHECK-NEXT: }
static void static_f_variant() {}
// CHECK: #pragma omp declare variant(static_f_variant) match(implementation={vendor(score(0):unknown)})
-// CHECK-NEXT: #pragma omp declare variant(static_f_variant) match(implementation={vendor(score(0):llvm)},device={kind(fpga)})
+// CHECK-NEXT: #pragma omp declare variant(static_f_variant) match(implementation={vendor(score(0):llvm)},device={kind(fpga),isa("aarch64")})
// CHECK-NEXT: static void static_f() {
// CHECK-NEXT: }
#pragma omp declare variant(static_f_variant) match(xxx = {})
-#pragma omp declare variant(static_f_variant) match(implementation={vendor(llvm)},device={kind(fpga)})
+#pragma omp declare variant(static_f_variant) match(implementation={vendor(llvm)},device={kind(fpga),isa("aarch64")})
#pragma omp declare variant(static_f_variant) match(implementation={vendor(unknown)})
static void static_f() {}
@@ -192,19 +194,19 @@
// CHECK: int fn_linkage_variant();
// CHECK: extern "C" {
-// CHECK: #pragma omp declare variant(fn_linkage_variant) match(implementation={vendor(score(0):xxx)},device={kind(cpu, host)})
+// CHECK: #pragma omp declare variant(fn_linkage_variant) match(implementation={vendor(score(0):xxx)},device={kind(cpu, host),isa("x86", "aarch64")})
// CHECK: int fn_linkage();
// CHECK: }
int fn_linkage_variant();
extern "C" {
-#pragma omp declare variant(fn_linkage_variant) match(implementation = {vendor(xxx)},device={kind(cpu,host)})
+#pragma omp declare variant(fn_linkage_variant) match(implementation = {vendor(xxx)},device={kind(cpu,host),isa("x86", "aarch64")})
int fn_linkage();
}
// CHECK: extern "C" int fn_linkage_variant1()
-// CHECK: #pragma omp declare variant(fn_linkage_variant1) match(implementation={vendor(score(0):xxx)},device={kind(cpu, host)})
+// CHECK: #pragma omp declare variant(fn_linkage_variant1) match(implementation={vendor(score(0):xxx)},device={kind(cpu, host),isa("x86", "aarch64")})
// CHECK: int fn_linkage1();
extern "C" int fn_linkage_variant1();
-#pragma omp declare variant(fn_linkage_variant1) match(implementation = {vendor(xxx)},device={kind(cpu,host)})
+#pragma omp declare variant(fn_linkage_variant1) match(implementation = {vendor(xxx)},device={kind(cpu,host),isa("x86", "aarch64")})
int fn_linkage1();
Index: clang/test/OpenMP/declare_variant_ast_print.c
===================================================================
--- clang/test/OpenMP/declare_variant_ast_print.c
+++ clang/test/OpenMP/declare_variant_ast_print.c
@@ -8,10 +8,10 @@
#pragma omp declare variant(foo) match(xxx={}, yyy={ccc})
#pragma omp declare variant(foo) match(xxx={vvv})
-#pragma omp declare variant(foo) match(implementation={vendor(llvm)}, device={kind(fpga)})
+#pragma omp declare variant(foo) match(implementation={vendor(llvm)}, device={kind(fpga), isa("i586")})
#pragma omp declare variant(foo) match(implementation={vendor(llvm), xxx})
-#pragma omp declare variant(foo) match(implementation={vendor(unknown)}, device={kind(gpu)})
-#pragma omp declare variant(foo) match(implementation={vendor(score(5): ibm, xxx, ibm)}, device={kind(cpu, nohost)})
+#pragma omp declare variant(foo) match(implementation={vendor(unknown)}, device={kind(gpu), isa("x86""_64", "arm")})
+#pragma omp declare variant(foo) match(implementation={vendor(score(5): ibm, xxx, ibm)}, device={kind(cpu, nohost), isa("ppc64le")})
#pragma omp declare variant(foo) match(device={kind(host)})
#pragma omp declare variant(foo) match(device={kind(nohost), xxx})
int bar(void);
@@ -19,8 +19,8 @@
// CHECK: int foo();
// CHECK-NEXT: #pragma omp declare variant(foo) match(device={kind(nohost)})
// CHECK-NEXT: #pragma omp declare variant(foo) match(device={kind(host)})
-// CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={vendor(score(5):ibm, xxx)},device={kind(cpu, nohost)})
-// CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={vendor(score(0):unknown)},device={kind(gpu)})
+// CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={vendor(score(5):ibm, xxx)},device={kind(cpu, nohost),isa("ppc64le")})
+// CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={vendor(score(0):unknown)},device={kind(gpu),isa("x86_64", "arm")})
// CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={vendor(score(0):llvm)})
-// CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={vendor(score(0):llvm)},device={kind(fpga)})
+// CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={vendor(score(0):llvm)},device={kind(fpga),isa("i586")})
// CHECK-NEXT: int bar();
Index: clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
===================================================================
--- clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
@@ -23,6 +23,7 @@
#include "clang/Sema/Lookup.h"
#include "clang/Sema/Template.h"
#include "clang/Sema/TemplateInstCallback.h"
+#include "llvm/ADT/Any.h"
#include "llvm/Support/TimeProfiler.h"
using namespace clang;
@@ -404,30 +405,33 @@
*std::next(Attr.ctxSelectorSets_begin(), I));
auto Ctx = static_cast<OpenMPContextSelectorKind>(
*std::next(Attr.ctxSelectors_begin(), I));
- switch (CtxSet) {
- case OMP_CTX_SET_implementation:
- switch (Ctx) {
- case OMP_CTX_vendor:
- Data.emplace_back(CtxSet, Ctx, Score, Attr.implVendors());
- break;
- case OMP_CTX_kind:
- case OMP_CTX_unknown:
- llvm_unreachable("Unexpected context selector kind.");
- }
+ SmallVector<llvm::Any, 4> Values;
+ switch (Ctx) {
+ case OMP_CTX_vendor:
+ assert(CtxSet == OMP_CTX_SET_implementation &&
+ "Expected implementation context selector set.");
+ for (StringRef Vendor : Attr.implVendors())
+ Values.push_back(Sema::OMPCtxStringType(Vendor));
break;
- case OMP_CTX_SET_device:
- switch (Ctx) {
- case OMP_CTX_kind:
- Data.emplace_back(CtxSet, Ctx, Score, Attr.deviceKinds());
- break;
- case OMP_CTX_vendor:
- case OMP_CTX_unknown:
- llvm_unreachable("Unexpected context selector kind.");
+ case OMP_CTX_kind:
+ assert(CtxSet == OMP_CTX_SET_device &&
+ "Expected device context selector set.");
+ for (StringRef Kind : Attr.deviceKinds())
+ Values.push_back(Sema::OMPCtxStringType(Kind));
+ break;
+ case OMP_CTX_isa:
+ assert(CtxSet == OMP_CTX_SET_device &&
+ "Expected device context selector set.");
+ for (Expr *ISA : Attr.deviceISAs()) {
+ ExprResult ISARes = Subst(ISA);
+ if (ISARes.isUsable())
+ Values.push_back(ISARes);
}
break;
- case OMP_CTX_SET_unknown:
- llvm_unreachable("Unexpected context selector set kind.");
+ case OMP_CTX_unknown:
+ llvm_unreachable("Unexpected context selector kind.");
}
+ Data.emplace_back(CtxSet, Ctx, Score, Values);
}
S.ActOnOpenMPDeclareVariantDirective(DeclVarData.getValue().first,
DeclVarData.getValue().second,
Index: clang/lib/Sema/SemaOpenMP.cpp
===================================================================
--- clang/lib/Sema/SemaOpenMP.cpp
+++ clang/lib/Sema/SemaOpenMP.cpp
@@ -28,6 +28,7 @@
#include "clang/Sema/Scope.h"
#include "clang/Sema/ScopeInfo.h"
#include "clang/Sema/SemaInternal.h"
+#include "llvm/ADT/Any.h"
#include "llvm/ADT/PointerEmbeddedInt.h"
using namespace clang;
@@ -5364,6 +5365,92 @@
return std::make_pair(FD, cast<Expr>(DRE));
}
+static ExprResult performOpenMPImplicitStringConversion(Sema &S, Expr *E) {
+ if (E->isTypeDependent() || E->isValueDependent() ||
+ E->isInstantiationDependent())
+ return E;
+ E = S.DefaultFunctionArrayLvalueConversion(E).get();
+ if (!E)
+ return ExprError();
+ class StringConverter : public Sema::ContextualImplicitConverter {
+ public:
+ StringConverter()
+ : Sema::ContextualImplicitConverter(/*Suppress=*/false,
+ /*SuppressConversion=*/true) {}
+
+ bool match(QualType ConvType) override {
+ if (ConvType->isArrayType() || ConvType->isPointerType()) {
+ const Type *ElTy = ConvType->getPointeeOrArrayElementType();
+ return ElTy->isWideCharType() || ElTy->isCharType() ||
+ ElTy->isChar8Type() || ElTy->isChar16Type();
+ }
+ return false;
+ }
+
+ Sema::SemaDiagnosticBuilder diagnoseNoMatch(Sema &S, SourceLocation Loc,
+ QualType T) override {
+ return S.Diag(Loc, diag::err_omp_trait_not_string) << T;
+ }
+
+ Sema::SemaDiagnosticBuilder diagnoseIncomplete(Sema &S, SourceLocation Loc,
+ QualType T) override {
+ return S.Diag(Loc, diag::err_omp_incomplete_type) << T;
+ }
+
+ Sema::SemaDiagnosticBuilder diagnoseExplicitConv(Sema &S,
+ SourceLocation Loc,
+ QualType T,
+ QualType ConvTy) override {
+ return S.Diag(Loc, diag::err_omp_explicit_conversion) << T << ConvTy;
+ }
+
+ Sema::SemaDiagnosticBuilder noteExplicitConv(Sema &S,
+ CXXConversionDecl *Conv,
+ QualType ConvTy) override {
+ return S.Diag(Conv->getLocation(), diag::note_omp_trait_conversion_here)
+ << ConvTy;
+ }
+
+ Sema::SemaDiagnosticBuilder diagnoseAmbiguous(Sema &S, SourceLocation Loc,
+ QualType T) override {
+ return S.Diag(Loc, diag::err_omp_trait_ambiguous_conversion) << T;
+ }
+
+ Sema::SemaDiagnosticBuilder noteAmbiguous(Sema &S, CXXConversionDecl *Conv,
+ QualType ConvTy) override {
+ return S.Diag(Conv->getLocation(), diag::note_omp_trait_conversion_here)
+ << ConvTy;
+ }
+
+ Sema::SemaDiagnosticBuilder diagnoseConversion(Sema &S, SourceLocation Loc,
+ QualType T,
+ QualType ConvTy) override {
+ llvm_unreachable("conversion functions are permitted");
+ }
+ } Converter;
+
+ ExprResult ER =
+ S.PerformContextualImplicitConversion(E->getBeginLoc(), E, Converter);
+ if (!ER.isUsable())
+ return ExprError();
+ E = ER.get();
+ Expr::EvalResult Res;
+ if (!E->EvaluateAsRValue(Res, S.Context) || !Res.Val.isLValue()) {
+ S.Diag(E->getExprLoc(), diag::err_omp_trait_not_constant_string)
+ << E->getSourceRange();
+ return ExprError();
+ }
+ const auto *SL = dyn_cast_or_null<StringLiteral>(
+ Res.Val.getLValueBase().dyn_cast<const Expr *>());
+ if (!SL) {
+ S.Diag(E->getExprLoc(), diag::err_omp_trait_not_constant_string)
+ << E->getSourceRange();
+ return ExprError();
+ }
+ return StringLiteral::Create(S.Context, SL->getBytes(), SL->getKind(),
+ SL->isPascal(), SL->getType(), E->getExprLoc());
+}
+
void Sema::ActOnOpenMPDeclareVariantDirective(
FunctionDecl *FD, Expr *VariantRef, SourceRange SR,
ArrayRef<OMPCtxSelectorData> Data) {
@@ -5373,6 +5460,7 @@
SmallVector<unsigned, 4> CtxSets;
SmallVector<unsigned, 4> Ctxs;
SmallVector<StringRef, 4> ImplVendors, DeviceKinds;
+ SmallVector<Expr *, 4> DeviceISAs;
bool IsError = false;
for (const OMPCtxSelectorData &D : Data) {
OpenMPContextSelectorSetKind CtxSet = D.CtxSet;
@@ -5396,29 +5484,45 @@
// The kind, arch, and isa selectors are given the values 2^l, 2^(l+1) and
// 2^(l+2), respectively, where l is the number of traits in the construct
// set.
- // TODO: implement correct logic for isa and arch traits.
+ // TODO: implement correct logic for arch traits.
// TODO: take the construct context set into account when it is
// implemented.
int L = 0; // Currently set the number of traits in construct set to 0,
// since the construct trait set in not supported yet.
if (CtxSet == OMP_CTX_SET_device && Ctx == OMP_CTX_kind)
Score = ActOnIntegerConstant(SourceLocation(), std::pow(2, L)).get();
+ else if (CtxSet == OMP_CTX_SET_device && Ctx == OMP_CTX_isa)
+ Score =
+ ActOnIntegerConstant(SourceLocation(), std::pow(2, L + 1)).get();
else
Score = ActOnIntegerConstant(SourceLocation(), 0).get();
}
- switch (Ctx) {
- case OMP_CTX_vendor:
- assert(CtxSet == OMP_CTX_SET_implementation &&
- "Expected implementation context selector set.");
- ImplVendors.append(D.Names.begin(), D.Names.end());
- break;
- case OMP_CTX_kind:
- assert(CtxSet == OMP_CTX_SET_device &&
- "Expected device context selector set.");
- DeviceKinds.append(D.Names.begin(), D.Names.end());
- break;
- case OMP_CTX_unknown:
- llvm_unreachable("Unknown context selector kind.");
+ for (const llvm::Any &Val : D.Names) {
+ switch (Ctx) {
+ case OMP_CTX_vendor:
+ assert(CtxSet == OMP_CTX_SET_implementation &&
+ "Expected implementation context selector set.");
+ ImplVendors.push_back(*llvm::any_cast<OMPCtxStringType>(&Val));
+ break;
+ case OMP_CTX_kind:
+ assert(CtxSet == OMP_CTX_SET_device &&
+ "Expected device context selector set.");
+ DeviceKinds.push_back(*llvm::any_cast<OMPCtxStringType>(&Val));
+ break;
+ case OMP_CTX_isa: {
+ assert(CtxSet == OMP_CTX_SET_device &&
+ "Expected device context selector set.");
+ auto ER = llvm::any_cast<ExprResult>(Val);
+ ER = performOpenMPImplicitStringConversion(*this, ER.get());
+ if (ER.isUsable())
+ DeviceISAs.push_back(ER.get());
+ else
+ IsError = true;
+ break;
+ }
+ case OMP_CTX_unknown:
+ llvm_unreachable("Unknown context selector kind.");
+ }
}
IsError = IsError || !Score;
CtxSets.push_back(CtxSet);
@@ -5430,7 +5534,7 @@
Context, VariantRef, CtxScores.begin(), CtxScores.size(),
CtxSets.begin(), CtxSets.size(), Ctxs.begin(), Ctxs.size(),
ImplVendors.begin(), ImplVendors.size(), DeviceKinds.begin(),
- DeviceKinds.size(), SR);
+ DeviceKinds.size(), DeviceISAs.begin(), DeviceISAs.size(), SR);
FD->addAttr(NewAttr);
}
}
Index: clang/lib/Parse/ParseOpenMP.cpp
===================================================================
--- clang/lib/Parse/ParseOpenMP.cpp
+++ clang/lib/Parse/ParseOpenMP.cpp
@@ -16,6 +16,7 @@
#include "clang/Parse/Parser.h"
#include "clang/Parse/RAIIObjectsForParser.h"
#include "clang/Sema/Scope.h"
+#include "llvm/ADT/Any.h"
#include "llvm/ADT/PointerIntPair.h"
#include "llvm/ADT/UniqueVector.h"
@@ -875,6 +876,7 @@
break;
}
case OMP_CTX_kind:
+ case OMP_CTX_isa:
case OMP_CTX_unknown:
P.Diag(Tok.getLocation(), diag::warn_omp_declare_variant_cs_name_expected)
<< "implementation";
@@ -888,6 +890,7 @@
/// Parse context selector for 'device' selector set:
/// 'kind' '(' <kind> { ',' <kind> } ')'
+/// 'isa' '(' <ISA> { ',' <ISA> } ')'
static void
parseDeviceSelector(Parser &P, SourceLocation Loc,
llvm::StringMap<SourceLocation> &UsedCtx,
@@ -959,6 +962,30 @@
Data.emplace_back(OMP_CTX_SET_device, CSKind, ExprResult(), Kinds);
break;
}
+ case OMP_CTX_isa: {
+ // Parse '('.
+ BalancedDelimiterTracker T(P, tok::l_paren, tok::annot_pragma_openmp_end);
+ (void)T.expectAndConsume(diag::err_expected_lparen_after,
+ CtxSelectorName.data());
+ SmallVector<ExprResult, 4> ISAs;
+ SourceLocation PrevLoc;
+ do {
+ PrevLoc = Tok.getLocation();
+ // Parse <ISA>, which is constant expression of string type.
+ ExprResult ISA = P.ParseConstantExpression();
+ if (ISA.isUsable())
+ ISAs.push_back(ISA);
+ if (!P.TryConsumeToken(tok::comma) && Tok.isNot(tok::r_paren))
+ P.Diag(Tok, diag::err_expected_punc) << "instruction set architecture";
+ } while (Tok.isNot(tok::r_paren) &&
+ Tok.isNot(tok::annot_pragma_openmp_end) &&
+ PrevLoc.getRawEncoding() != Tok.getLocation().getRawEncoding());
+ // Parse ')'.
+ (void)T.consumeClose();
+ if (!ISAs.empty())
+ Data.emplace_back(OMP_CTX_SET_device, CSKind, ExprResult(), ISAs);
+ break;
+ }
case OMP_CTX_vendor:
case OMP_CTX_unknown:
P.Diag(Tok.getLocation(), diag::warn_omp_declare_variant_cs_name_expected)
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -21,10 +21,12 @@
#include "clang/Basic/BitmaskEnum.h"
#include "llvm/ADT/ArrayRef.h"
#include "llvm/ADT/SetOperations.h"
+#include "llvm/ADT/UniqueVector.h"
#include "llvm/Bitcode/BitcodeReader.h"
#include "llvm/IR/DerivedTypes.h"
#include "llvm/IR/GlobalValue.h"
#include "llvm/IR/Value.h"
+#include "llvm/Support/ConvertUTF.h"
#include "llvm/Support/Format.h"
#include "llvm/Support/raw_ostream.h"
#include <cassert>
@@ -11028,7 +11030,7 @@
namespace {
using OMPContextSelectorData =
- OpenMPCtxSelectorData<ArrayRef<StringRef>, llvm::APSInt>;
+ OpenMPCtxSelectorData<SmallVector<SmallString<8>, 4>, llvm::APSInt>;
using CompleteOMPContextSelectorData = SmallVector<OMPContextSelectorData, 4>;
} // anonymous namespace
@@ -11135,6 +11137,24 @@
return true;
}
+/// Checks for device={isa(<isa>)} context selector.
+/// \returns true if <isa>=current isa or false, otherwise.
+template <>
+bool checkContext<OMP_CTX_SET_device, OMP_CTX_isa, CodeGenModule &>(
+ const OMPContextSelectorData &Data, CodeGenModule &CGM) {
+ for (StringRef Name : Data.Names) {
+ /// Consider arch with vendor/os/env as invalid.
+ if (Name.contains('-'))
+ return false;
+ llvm::Triple CustomTriple(Name.lower());
+ if (CustomTriple.getArch() != CGM.getTarget().getTriple().getArch() ||
+ (CustomTriple.getSubArch() != llvm::Triple::NoSubArch &&
+ CustomTriple.getSubArch() != CGM.getTarget().getTriple().getSubArch()))
+ return false;
+ }
+ return true;
+}
+
bool matchesContext(CodeGenModule &CGM,
const CompleteOMPContextSelectorData &ContextData) {
for (const OMPContextSelectorData &Data : ContextData) {
@@ -11152,6 +11172,13 @@
CGM))
return false;
break;
+ case OMP_CTX_isa:
+ assert(Data.CtxSet == OMP_CTX_SET_device &&
+ "Expected device context selector set.");
+ if (!checkContext<OMP_CTX_SET_device, OMP_CTX_isa, CodeGenModule &>(Data,
+ CGM))
+ return false;
+ break;
case OMP_CTX_unknown:
llvm_unreachable("Unknown context selector kind.");
}
@@ -11177,15 +11204,47 @@
case OMP_CTX_vendor:
assert(CtxSet == OMP_CTX_SET_implementation &&
"Expected implementation context selector set.");
- Data.back().Names =
- llvm::makeArrayRef(A->implVendors_begin(), A->implVendors_end());
+ Data.back().Names.append(A->implVendors_begin(), A->implVendors_end());
break;
case OMP_CTX_kind:
assert(CtxSet == OMP_CTX_SET_device &&
"Expected device context selector set.");
- Data.back().Names =
- llvm::makeArrayRef(A->deviceKinds_begin(), A->deviceKinds_end());
+ Data.back().Names.append(A->deviceKinds_begin(), A->deviceKinds_end());
break;
+ case OMP_CTX_isa: {
+ assert(CtxSet == OMP_CTX_SET_device &&
+ "Expected device context selector set.");
+ llvm::UniqueVector<std::string> ISAs;
+ for (const Expr *ISA : A->deviceISAs()) {
+ const auto *SL = cast<StringLiteral>(ISA);
+ std::string ISAName;
+ switch(SL->getKind()) {
+ case StringLiteral::Ascii:
+ case StringLiteral::UTF8:
+ ISAName = SL->getString();
+ break;
+ case StringLiteral::UTF16:
+ if (!llvm::convertUTF16ToUTF8String(
+ llvm::makeArrayRef(SL->getBytes().begin(), SL->getBytes().end()),
+ ISAName))
+ return CompleteOMPContextSelectorData();
+ break;
+ case StringLiteral::UTF32:
+ llvm_unreachable("UTF32 is unsupported.");
+ case StringLiteral::Wide: {
+ std::wstring WideString;
+ for (int I = 0, E = SL->getLength(); I < E; ++I)
+ WideString += static_cast<wchar_t>(SL->getCodeUnit(I));
+ if (!llvm::convertWideToUTF8(WideString, ISAName))
+ return CompleteOMPContextSelectorData();
+ break;
+ }
+ }
+ ISAs.insert(ISAName);
+ }
+ Data.back().Names.append(ISAs.begin(), ISAs.end());
+ break;
+ }
case OMP_CTX_unknown:
llvm_unreachable("Unknown context selector kind.");
}
Index: clang/include/clang/Sema/Sema.h
===================================================================
--- clang/include/clang/Sema/Sema.h
+++ clang/include/clang/Sema/Sema.h
@@ -61,6 +61,7 @@
#include <vector>
namespace llvm {
+ class Any;
class APSInt;
template <typename ValueT> struct DenseMapInfo;
template <typename ValueT, typename ValueInfoT> class DenseSet;
@@ -9318,7 +9319,7 @@
/// Struct to store the context selectors info for declare variant directive.
using OMPCtxStringType = SmallString<8>;
using OMPCtxSelectorData =
- OpenMPCtxSelectorData<SmallVector<OMPCtxStringType, 4>, ExprResult>;
+ OpenMPCtxSelectorData<SmallVector<llvm::Any, 4>, ExprResult>;
/// Checks if the variant/multiversion functions are compatible.
bool areMultiversionVariantFunctionsCompatible(
Index: clang/include/clang/Basic/OpenMPKinds.def
===================================================================
--- clang/include/clang/Basic/OpenMPKinds.def
+++ clang/include/clang/Basic/OpenMPKinds.def
@@ -226,6 +226,7 @@
// OpenMP context selectors.
OPENMP_CONTEXT_SELECTOR(vendor)
OPENMP_CONTEXT_SELECTOR(kind)
+OPENMP_CONTEXT_SELECTOR(isa)
// OpenMP directives.
OPENMP_DIRECTIVE(threadprivate)
Index: clang/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -9634,6 +9634,14 @@
def note_omp_marked_declare_variant_here : Note<"marked as 'declare variant' here">;
def err_omp_one_defaultmap_each_category: Error<
"at most one defaultmap clause for each variable-category can appear on the directive">;
+def err_omp_trait_not_string : Error<
+ "trait expression must have a string type, not %0">;
+def note_omp_trait_conversion_here : Note<
+ "conversion to a string type %0 declared here">;
+def err_omp_trait_ambiguous_conversion : Error<
+ "ambiguous conversion from type %0 to a string type">;
+def err_omp_trait_not_constant_string : Error<
+ "trait expression must be a constant string expression">;
} // end of OpenMP category
let CategoryName = "Related Result Type Issue" in {
Index: clang/include/clang/Basic/Attr.td
===================================================================
--- clang/include/clang/Basic/Attr.td
+++ clang/include/clang/Basic/Attr.td
@@ -3316,7 +3316,8 @@
VariadicUnsignedArgument<"CtxSelectorSets">,
VariadicUnsignedArgument<"CtxSelectors">,
VariadicStringArgument<"ImplVendors">,
- VariadicStringArgument<"DeviceKinds">
+ VariadicStringArgument<"DeviceKinds">,
+ VariadicExprArgument<"DeviceISAs">
];
let AdditionalMembers = [{
void printScore(raw_ostream & OS, const PrintingPolicy &Policy, unsigned I) const {
@@ -3387,6 +3388,19 @@
}
OS << ")";
break;
+ case OMP_CTX_isa:
+ assert(CtxSet == OMP_CTX_SET_device &&
+ "Expected device context selector set.");
+ OS << "isa(";
+ if (deviceISAs_size() > 0) {
+ (*deviceISAs().begin())->printPretty(OS, nullptr, Policy);
+ for (const Expr *ISA : llvm::drop_begin(deviceISAs(), 1)) {
+ OS << ", ";
+ ISA->printPretty(OS, nullptr, Policy);
+ }
+ }
+ OS << ")";
+ break;
case OMP_CTX_unknown:
llvm_unreachable("Unknown context selector.");
}
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
