Skip to content

Commit 70e4c84

Browse files
author
iclsrc
committed
Merge from 'sycl' to 'sycl-web' (#5)
2 parents 586bea3 + 50e7abd commit 70e4c84

File tree

20 files changed

+1113
-764
lines changed

20 files changed

+1113
-764
lines changed

‎clang/include/clang/AST/PrettyPrinter.h

Lines changed: 6 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -171,10 +171,13 @@ struct PrintingPolicy {
171171
/// When true, suppress printing of lifetime qualifier in ARC.
172172
unsigned SuppressLifetimeQualifiers : 1;
173173

174-
/// When true prints a canonical type instead of an alias. E.g.
174+
/// When true prints a canonical type instead of an alias.
175+
/// Also removes preceeding keywords if there is one. E.g.
175176
/// \code
176-
/// using SizeT = int;
177-
/// template<SizeT N> class C;
177+
/// namespace NS {
178+
/// using SizeT = int;
179+
/// }
180+
/// template<typename NS::SizeT N> class C;
178181
/// \endcode
179182
/// will be printed as
180183
/// \code

‎clang/include/clang/Basic/Attr.td

Lines changed: 11 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -1191,17 +1191,22 @@ def OpenCLUnrollHint : InheritableAttr {
11911191

11921192
def LoopUnrollHint : InheritableAttr {
11931193
let Spellings = [CXX11<"clang","loop_unroll">];
1194-
let Args = [UnsignedArgument<"UnrollHint">];
1194+
let Args = [ExprArgument<"UnrollHintExpr">];
11951195
let LangOpts = [SYCLIsDevice, SYCLIsHost];
1196+
let HasCustomTypeTransform = 1;
11961197
let AdditionalMembers = [{
11971198
static const char *getName() {
11981199
return "loop_unroll";
11991200
}
1200-
std::string getDiagnosticName() const {
1201-
std::string Value = "";
1202-
if (getUnrollHint())
1203-
Value = "(" + std::to_string(getUnrollHint()) + ")";
1204-
return "[[clang::loop_unroll" + Value + "]]";
1201+
std::string getDiagnosticName(const PrintingPolicy &Policy) const {
1202+
std::string ValueName;
1203+
llvm::raw_string_ostream OS(ValueName);
1204+
if (auto *E = getUnrollHintExpr()) {
1205+
OS << "(";
1206+
E->printPretty(OS, nullptr, Policy);
1207+
OS << ")";
1208+
}
1209+
return "[[clang::loop_unroll" + OS.str() + "]]";
12051210
}
12061211
}];
12071212
let Documentation = [LoopUnrollHintDocs];

‎clang/include/clang/Sema/Sema.h

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1706,6 +1706,11 @@ class Sema final {
17061706
FPGALoopAttrT *BuildSYCLIntelFPGALoopAttr(const AttributeCommonInfo &A,
17071707
Expr *E);
17081708

1709+
LoopUnrollHintAttr *BuildLoopUnrollHintAttr(const AttributeCommonInfo &A,
1710+
Expr *E);
1711+
OpenCLUnrollHintAttr *
1712+
BuildOpenCLLoopUnrollHintAttr(const AttributeCommonInfo &A, Expr *E);
1713+
17091714
bool CheckQualifiedFunctionForTypeId(QualType T, SourceLocation Loc);
17101715

17111716
bool CheckFunctionReturnType(QualType T, SourceLocation Loc);

‎clang/lib/AST/TypePrinter.cpp

Lines changed: 8 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1332,9 +1332,14 @@ void TypePrinter::printElaboratedBefore(const ElaboratedType *T,
13321332
// The tag definition will take care of these.
13331333
if (!Policy.IncludeTagDefinition)
13341334
{
1335-
OS << TypeWithKeyword::getKeywordName(T->getKeyword());
1336-
if (T->getKeyword() != ETK_None)
1337-
OS << " ";
1335+
// When removing aliases don't print keywords to avoid having things
1336+
// like 'typename int'
1337+
if (!Policy.SuppressTypedefs)
1338+
{
1339+
OS << TypeWithKeyword::getKeywordName(T->getKeyword());
1340+
if (T->getKeyword() != ETK_None)
1341+
OS << " ";
1342+
}
13381343
NestedNameSpecifier *Qualifier = T->getQualifier();
13391344
if (Qualifier && !(Policy.SuppressTypedefs &&
13401345
T->getNamedType()->getTypeClass() == Type::Typedef))

‎clang/lib/CodeGen/CGLoopInfo.cpp

Lines changed: 6 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -706,8 +706,12 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx,
706706
// 1 - disable unroll.
707707
// other positive integer n - unroll by n.
708708
if (OpenCLHint || UnrollHint) {
709-
ValueInt = OpenCLHint ? OpenCLHint->getUnrollHint()
710-
: UnrollHint->getUnrollHint();
709+
ValueInt = 0;
710+
if (OpenCLHint)
711+
ValueInt = OpenCLHint->getUnrollHint();
712+
else if (Expr *E = UnrollHint->getUnrollHintExpr())
713+
ValueInt = E->EvaluateKnownConstInt(Ctx).getSExtValue();
714+
711715
if (ValueInt == 0) {
712716
State = LoopHintAttr::Enable;
713717
} else if (ValueInt != 1) {

‎clang/lib/Sema/SemaStmtAttr.cpp

Lines changed: 45 additions & 27 deletions
Original file line numberDiff line numberDiff line change
@@ -554,11 +554,47 @@ void CheckForIncompatibleUnrollHintAttributes(
554554
SourceLocation Loc = Range.getBegin();
555555
S.Diag(Loc, diag::err_loop_unroll_compatibility)
556556
<< PragmaUnroll->getDiagnosticName(Policy)
557-
<< AttrUnroll->getDiagnosticName();
557+
<< AttrUnroll->getDiagnosticName(Policy);
558558
}
559559
}
560560

561-
template <typename LoopUnrollAttrT>
561+
static bool CheckLoopUnrollAttrExpr(Sema &S, Expr *E,
562+
const AttributeCommonInfo &A,
563+
unsigned *UnrollFactor = nullptr) {
564+
if (E && !E->isInstantiationDependent()) {
565+
llvm::APSInt ArgVal(32);
566+
567+
if (!E->isIntegerConstantExpr(ArgVal, S.Context))
568+
return S.Diag(E->getExprLoc(), diag::err_attribute_argument_type)
569+
<< A.getAttrName() << AANT_ArgumentIntegerConstant
570+
<< E->getSourceRange();
571+
572+
if (ArgVal.isNonPositive())
573+
return S.Diag(E->getExprLoc(),
574+
diag::err_attribute_requires_positive_integer)
575+
<< A.getAttrName() << /* positive */ 0;
576+
577+
if (UnrollFactor)
578+
*UnrollFactor = ArgVal.getZExtValue();
579+
}
580+
return false;
581+
}
582+
583+
LoopUnrollHintAttr *Sema::BuildLoopUnrollHintAttr(const AttributeCommonInfo &A,
584+
Expr *E) {
585+
return !CheckLoopUnrollAttrExpr(*this, E, A)
586+
? new (Context) LoopUnrollHintAttr(Context, A, E)
587+
: nullptr;
588+
}
589+
590+
OpenCLUnrollHintAttr *
591+
Sema::BuildOpenCLLoopUnrollHintAttr(const AttributeCommonInfo &A, Expr *E) {
592+
unsigned UnrollFactor = 0;
593+
return !CheckLoopUnrollAttrExpr(*this, E, A, &UnrollFactor)
594+
? new (Context) OpenCLUnrollHintAttr(Context, A, UnrollFactor)
595+
: nullptr;
596+
}
597+
562598
static Attr *handleLoopUnrollHint(Sema &S, Stmt *St, const ParsedAttr &A,
563599
SourceRange Range) {
564600
// Although the feature was introduced only in OpenCL C v2.0 s6.11.5, it's
@@ -574,30 +610,13 @@ static Attr *handleLoopUnrollHint(Sema &S, Stmt *St, const ParsedAttr &A,
574610
return nullptr;
575611
}
576612

577-
unsigned UnrollFactor = 0;
578-
579-
if (NumArgs == 1) {
580-
Expr *E = A.getArgAsExpr(0);
581-
llvm::APSInt ArgVal(32);
582-
583-
if (!E->isIntegerConstantExpr(ArgVal, S.Context)) {
584-
S.Diag(A.getLoc(), diag::err_attribute_argument_type)
585-
<< A << AANT_ArgumentIntegerConstant << E->getSourceRange();
586-
return nullptr;
587-
}
588-
589-
int Val = ArgVal.getSExtValue();
590-
591-
if (Val <= 0) {
592-
S.Diag(A.getRange().getBegin(),
593-
diag::err_attribute_requires_positive_integer)
594-
<< A << /* positive */ 0;
595-
return nullptr;
596-
}
597-
UnrollFactor = Val;
598-
}
613+
Expr *E = NumArgs ? A.getArgAsExpr(0) : nullptr;
614+
if (A.getParsedKind() == ParsedAttr::AT_OpenCLUnrollHint)
615+
return S.BuildOpenCLLoopUnrollHintAttr(A, E);
616+
else if (A.getParsedKind() == ParsedAttr::AT_LoopUnrollHint)
617+
return S.BuildLoopUnrollHintAttr(A, E);
599618

600-
return LoopUnrollAttrT::CreateImplicit(S.Context, UnrollFactor);
619+
return nullptr;
601620
}
602621

603622
static Attr *ProcessStmtAttribute(Sema &S, Stmt *St, const ParsedAttr &A,
@@ -620,9 +639,8 @@ static Attr *ProcessStmtAttribute(Sema &S, Stmt *St, const ParsedAttr &A,
620639
case ParsedAttr::AT_SYCLIntelFPGAMaxConcurrency:
621640
return handleIntelFPGALoopAttr<SYCLIntelFPGAMaxConcurrencyAttr>(S, A);
622641
case ParsedAttr::AT_OpenCLUnrollHint:
623-
return handleLoopUnrollHint<OpenCLUnrollHintAttr>(S, St, A, Range);
624642
case ParsedAttr::AT_LoopUnrollHint:
625-
return handleLoopUnrollHint<LoopUnrollHintAttr>(S, St, A, Range);
643+
return handleLoopUnrollHint(S, St, A, Range);
626644
case ParsedAttr::AT_Suppress:
627645
return handleSuppressAttr(S, St, A, Range);
628646
default:

‎clang/lib/Sema/SemaTemplateInstantiate.cpp

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1055,6 +1055,8 @@ namespace {
10551055
const SYCLIntelFPGAMaxConcurrencyAttr *
10561056
TransformSYCLIntelFPGAMaxConcurrencyAttr(
10571057
const SYCLIntelFPGAMaxConcurrencyAttr *MC);
1058+
const LoopUnrollHintAttr *
1059+
TransformLoopUnrollHintAttr(const LoopUnrollHintAttr *LU);
10581060

10591061
ExprResult TransformPredefinedExpr(PredefinedExpr *E);
10601062
ExprResult TransformDeclRefExpr(DeclRefExpr *E);
@@ -1519,6 +1521,13 @@ TemplateInstantiator::TransformSYCLIntelFPGAMaxConcurrencyAttr(
15191521
*MC, TransformedExpr);
15201522
}
15211523

1524+
const LoopUnrollHintAttr *TemplateInstantiator::TransformLoopUnrollHintAttr(
1525+
const LoopUnrollHintAttr *LU) {
1526+
Expr *TransformedExpr =
1527+
getDerived().TransformExpr(LU->getUnrollHintExpr()).get();
1528+
return getSema().BuildLoopUnrollHintAttr(*LU, TransformedExpr);
1529+
}
1530+
15221531
ExprResult TemplateInstantiator::transformNonTypeTemplateParmRef(
15231532
NonTypeTemplateParmDecl *parm,
15241533
SourceLocation loc,

‎clang/test/CodeGenSYCL/kernel_name_with_typedefs.cpp

Lines changed: 17 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -85,6 +85,18 @@ struct kernel_name2<space::clong_t, volatile space::culong_t> {};
8585
template <>
8686
struct kernel_name2<space::a_t, volatile space::b_t> {};
8787

88+
// CHECK: template <long T> struct kernel_name3;
89+
template <typename space::long_t T>
90+
struct kernel_name3;
91+
92+
struct foo {
93+
using type = long;
94+
};
95+
96+
// CHECK: template <long T> struct kernel_name4;
97+
template <typename foo::type T>
98+
struct kernel_name4;
99+
88100
int main() {
89101
dummy_functor f;
90102
// non-type template arguments
@@ -116,5 +128,10 @@ int main() {
116128
single_task<kernel_name2<space::clong_t, volatile space::culong_t>>(f);
117129
// CHECK: template <> struct KernelInfo<::kernel_name2< ::A, volatile ::space::B>> {
118130
single_task<kernel_name2<space::a_t, volatile space::b_t>>(f);
131+
// CHECK: template <> struct KernelInfo<::kernel_name3<1>> {
132+
single_task<kernel_name3<1>>(f);
133+
// CHECK: template <> struct KernelInfo<::kernel_name4<1>> {
134+
single_task<kernel_name4<1>>(f);
135+
119136
return 0;
120137
}

‎clang/test/CodeGenSYCL/loop_unroll.cpp

Lines changed: 30 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -1,42 +1,54 @@
11
// RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -fsycl-is-device -emit-llvm %s -o - | FileCheck %s
22

3-
// CHECK: br label %for.cond, !llvm.loop ![[COUNT:[0-9]+]]
4-
// CHECK: br label %while.cond, !llvm.loop ![[DISABLE:[0-9]+]]
5-
// CHECK: br i1 %{{.*}}, label %do.body, label %do.end, !llvm.loop ![[ENABLE:[0-9]+]]
63

7-
// CHECK: ![[COUNT]] = distinct !{![[COUNT]], ![[COUNT_A:[0-9]+]]}
8-
// CHECK-NEXT: ![[COUNT_A]] = !{!"llvm.loop.unroll.count", i32 8}
4+
void enable() {
5+
int i = 1000;
6+
// CHECK: br i1 %{{.*}}, label %do.body, label %do.end, !llvm.loop ![[ENABLE:[0-9]+]]
7+
[[clang::loop_unroll]]
8+
do {} while (i--);
9+
}
10+
11+
template <int A>
912
void count() {
13+
// CHECK: br label %for.cond, !llvm.loop ![[COUNT:[0-9]+]]
1014
[[clang::loop_unroll(8)]]
1115
for (int i = 0; i < 1000; ++i);
16+
// CHECK: br label %for.cond2, !llvm.loop ![[COUNT_TEMPLATE:[0-9]+]]
17+
[[clang::loop_unroll(A)]]
18+
for (int i = 0; i < 1000; ++i);
1219
}
1320

14-
// CHECK: ![[DISABLE]] = distinct !{![[DISABLE]], ![[DISABLE_A:[0-9]+]]}
15-
// CHECK-NEXT: ![[DISABLE_A]] = !{!"llvm.loop.unroll.disable"}
21+
template <int A>
1622
void disable() {
17-
int i = 1000;
23+
int i = 1000, j = 100;
24+
// CHECK: br label %while.cond, !llvm.loop ![[DISABLE:[0-9]+]]
1825
[[clang::loop_unroll(1)]]
26+
while (j--);
27+
// CHECK: br label %while.cond1, !llvm.loop ![[DISABLE_TEMPLATE:[0-9]+]]
28+
[[clang::loop_unroll(A)]]
1929
while (i--);
2030
}
2131

22-
// CHECK: ![[ENABLE]] = distinct !{![[ENABLE]], ![[ENABLE_A:[0-9]+]]}
23-
// CHECK-NEXT: ![[ENABLE_A]] = !{!"llvm.loop.unroll.enable"}
24-
void enable() {
25-
int i = 1000;
26-
[[clang::loop_unroll]]
27-
do {} while (i--);
28-
}
29-
3032
template <typename name, typename Func>
3133
__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
3234
kernelFunc();
3335
}
3436

3537
int main() {
3638
kernel_single_task<class kernel_function>([]() {
37-
count();
38-
disable();
39+
count<4>();
40+
disable<1>();
3941
enable();
4042
});
4143
return 0;
4244
}
45+
46+
// CHECK: ![[ENABLE]] = distinct !{![[ENABLE]], ![[ENABLE_A:[0-9]+]]}
47+
// CHECK-NEXT: ![[ENABLE_A]] = !{!"llvm.loop.unroll.enable"}
48+
// CHECK: ![[COUNT]] = distinct !{![[COUNT]], ![[COUNT_A:[0-9]+]]}
49+
// CHECK-NEXT: ![[COUNT_A]] = !{!"llvm.loop.unroll.count", i32 8}
50+
// CHECK: ![[COUNT_TEMPLATE]] = distinct !{![[COUNT_TEMPLATE]], ![[COUNT_TEMPLATE_A:[0-9]+]]}
51+
// CHECK-NEXT: ![[COUNT_TEMPLATE_A]] = !{!"llvm.loop.unroll.count", i32 4}
52+
// CHECK: ![[DISABLE]] = distinct !{![[DISABLE]], ![[DISABLE_A:[0-9]+]]}
53+
// CHECK-NEXT: ![[DISABLE_A]] = !{!"llvm.loop.unroll.disable"}
54+
// CHECKL ![[DISABLE_TEMPLATE]] = distinct !{!![[DISABLE_TEMPLATE]], ![[DISABLE_A]]}

‎clang/test/CodeGenSYCL/loop_unroll_host.cpp

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2,6 +2,14 @@
22
// CHECK: br label %{{.*}}, !llvm.loop ![[COUNT:[0-9]+]]
33
// CHECK: br label %{{.*}}, !llvm.loop ![[DISABLE:[0-9]+]]
44
// CHECK: br i1 %{{.*}}, label %{{.*}}, label %{{.*}}, !llvm.loop ![[ENABLE:[0-9]+]]
5+
// CHECK: br label %{{.*}}, !llvm.loop ![[COUNT_TEMPLATE:[0-9]+]]
6+
// CHECK: br label %{{.*}}, !llvm.loop ![[DISABLE_TEMPLATE:[0-9]+]]
7+
8+
template <int A>
9+
void unroll() {
10+
[[clang::loop_unroll(A)]]
11+
for (int i = 0; i < 1000; ++i);
12+
}
513

614
int main() {
715
// CHECK: ![[COUNT]] = distinct !{![[COUNT]], ![[COUNT_A:[0-9]+]]}
@@ -18,5 +26,11 @@ int main() {
1826
i = 1000;
1927
[[clang::loop_unroll]]
2028
do {} while (i--);
29+
30+
// CHECK: ![[COUNT_TEMPLATE]] = distinct !{![[COUNT_TEMPLATE]], ![[COUNT_TEMPLATE_A:[0-9]+]]}
31+
// CHECK-NEXT: ![[COUNT_TEMPLATE_A]] = !{!"llvm.loop.unroll.count", i32 8}
32+
unroll<8>();
33+
// CHECK: ![[DISABLE_TEMPLATE]] = distinct !{![[DISABLE_TEMPLATE]], ![[DISABLE_A]]}
34+
unroll<1>();
2135
return 0;
2236
}

‎clang/test/SemaSYCL/loop_unroll.cpp

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,12 @@
11
// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify -pedantic %s
22

3+
template <int A>
4+
void bar() {
5+
// expected-error@+1 {{'loop_unroll' attribute requires a positive integral compile time constant expression}}
6+
[[clang::loop_unroll(A)]]
7+
for (int i = 0; i < 10; ++i);
8+
}
9+
310
void foo() {
411
// expected-error@+1 {{clang loop attributes must be applied to for, while, or do statements}}
512
[[clang::loop_unroll(8)]] int a[10];
@@ -44,6 +51,12 @@ void foo() {
4451
constexpr int c = 4;
4552
[[clang::loop_unroll(c)]]
4653
for (int i = 0; i < 10; ++i);
54+
55+
// expected-note@+1 {{in instantiation of function template specialization}}
56+
bar<-4>();
57+
58+
// no error expected
59+
bar<c>();
4760
}
4861

4962
template <typename name, typename Func>

‎sycl/include/CL/sycl/detail/type_traits.hpp

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -157,6 +157,10 @@ template <typename T, int N, template <typename> class S>
157157
using is_gen_based_on_type_sizeof =
158158
bool_constant<S<T>::value && (sizeof(vector_element_t<T>) == N)>;
159159

160+
template <typename> struct is_vec : std::false_type {};
161+
template <typename T, std::size_t N>
162+
struct is_vec<cl::sycl::vec<T, N>> : std::true_type {};
163+
160164
// is_integral
161165
template <typename T>
162166
struct is_integral : std::is_integral<vector_element_t<T>> {};

‎sycl/include/CL/sycl/event.hpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -9,6 +9,7 @@
99
#pragma once
1010

1111
#include <CL/sycl/detail/common.hpp>
12+
#include <CL/sycl/info/info_desc.hpp>
1213
#include <CL/sycl/stl.hpp>
1314

1415
#include <memory>

0 commit comments

Comments
 (0)