Skip to content

Commit 0824ff0

Browse files
author
iclsrc
committed
Merge from 'sycl' to 'sycl-web'
2 parents 746dd89 + 0bfd34e commit 0824ff0

22 files changed

+1004
-885
lines changed

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

+3-3
Original file line numberDiff line numberDiff line change
@@ -1908,7 +1908,7 @@ def SYCLIntelFPGAInitiationInterval : DeclOrStmtAttr {
19081908
let Subjects = SubjectList<[ForStmt, CXXForRangeStmt, WhileStmt, DoStmt, Function],
19091909
ErrorDiag,
19101910
"'for', 'while', 'do' statements, and functions">;
1911-
let Args = [ExprArgument<"IntervalExpr", /*opt*/1>];
1911+
let Args = [ExprArgument<"IntervalExpr">];
19121912
let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
19131913
let HasCustomTypeTransform = 1;
19141914
let Documentation = [SYCLIntelFPGAInitiationIntervalAttrDocs];
@@ -1955,7 +1955,7 @@ def SYCLIntelFPGAMaxInterleaving : StmtAttr {
19551955
CXX11<"intel","max_interleaving">];
19561956
let Subjects = SubjectList<[ForStmt, CXXForRangeStmt, WhileStmt, DoStmt],
19571957
ErrorDiag, "'for', 'while', and 'do' statements">;
1958-
let Args = [ExprArgument<"NExpr", /*opt*/1>];
1958+
let Args = [ExprArgument<"NExpr">];
19591959
let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
19601960
let HasCustomTypeTransform = 1;
19611961
let Documentation = [SYCLIntelFPGAMaxInterleavingAttrDocs];
@@ -1966,7 +1966,7 @@ def SYCLIntelFPGASpeculatedIterations : StmtAttr {
19661966
CXX11<"intel","speculated_iterations">];
19671967
let Subjects = SubjectList<[ForStmt, CXXForRangeStmt, WhileStmt, DoStmt],
19681968
ErrorDiag, "'for', 'while', and 'do' statements">;
1969-
let Args = [ExprArgument<"NExpr", /*opt*/1>];
1969+
let Args = [ExprArgument<"NExpr">];
19701970
let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
19711971
let HasCustomTypeTransform = 1;
19721972
let Documentation = [SYCLIntelFPGASpeculatedIterationsAttrDocs];

‎clang/include/clang/Parse/Parser.h

-12
Original file line numberDiff line numberDiff line change
@@ -2836,18 +2836,6 @@ class Parser : public CodeCompletionHandler {
28362836
void ParseBorlandTypeAttributes(ParsedAttributes &attrs);
28372837
void ParseOpenCLKernelAttributes(ParsedAttributes &attrs);
28382838
void ParseOpenCLQualifiers(ParsedAttributes &Attrs);
2839-
/// Parses opencl_unroll_hint attribute if language is OpenCL v2.0
2840-
/// or higher.
2841-
/// \return false if error happens.
2842-
bool MaybeParseOpenCLUnrollHintAttribute(ParsedAttributes &Attrs) {
2843-
if (getLangOpts().OpenCL)
2844-
return ParseOpenCLUnrollHintAttribute(Attrs);
2845-
return true;
2846-
}
2847-
/// Parses opencl_unroll_hint attribute.
2848-
/// \return false if error happens.
2849-
bool ParseOpenCLUnrollHintAttribute(ParsedAttributes &Attrs);
2850-
28512839
void ParseNullabilityTypeSpecifiers(ParsedAttributes &attrs);
28522840
VersionTuple ParseVersionTuple(SourceRange &Range);
28532841
void ParseAvailabilityAttribute(IdentifierInfo &Availability,

‎clang/lib/Frontend/CompilerInvocation.cpp

+3
Original file line numberDiff line numberDiff line change
@@ -3501,6 +3501,9 @@ void CompilerInvocation::GenerateLangArgs(const LangOptions &Opts,
35013501
if (Opts.getSignReturnAddressKey() ==
35023502
LangOptions::SignReturnAddressKeyKind::BKey)
35033503
GenerateArg(Args, OPT_msign_return_address_key_EQ, "b_key", SA);
3504+
3505+
if (Opts.DeclareSPIRVBuiltins)
3506+
GenerateArg(Args, OPT_fdeclare_spirv_builtins, SA);
35043507
}
35053508

35063509
bool CompilerInvocation::ParseLangArgs(LangOptions &Opts, ArgList &Args,

‎clang/lib/Sema/SemaStmtAttr.cpp

+5-58
Original file line numberDiff line numberDiff line change
@@ -70,40 +70,15 @@ static Attr *handleSuppressAttr(Sema &S, Stmt *St, const ParsedAttr &A,
7070

7171
template <typename FPGALoopAttrT>
7272
static Attr *handleIntelFPGALoopAttr(Sema &S, Stmt *St, const ParsedAttr &A) {
73-
if (!isa<ForStmt, CXXForRangeStmt, DoStmt, WhileStmt>(St)) {
74-
S.Diag(A.getLoc(), diag::err_attribute_wrong_decl_type_str)
75-
<< A << "'for', 'while', and 'do' statements";
76-
return nullptr;
77-
}
78-
79-
unsigned NumArgs = A.getNumArgs();
80-
if (NumArgs == 0) {
81-
if (A.getKind() == ParsedAttr::AT_SYCLIntelFPGAInitiationInterval ||
82-
A.getKind() == ParsedAttr::AT_SYCLIntelFPGAMaxConcurrency ||
83-
A.getKind() == ParsedAttr::AT_SYCLIntelFPGAMaxInterleaving ||
84-
A.getKind() == ParsedAttr::AT_SYCLIntelFPGASpeculatedIterations) {
85-
S.Diag(A.getLoc(), diag::warn_attribute_too_few_arguments) << A << 1;
86-
return nullptr;
87-
}
88-
}
89-
9073
S.CheckDeprecatedSYCLAttributeSpelling(A);
9174

9275
return S.BuildSYCLIntelFPGALoopAttr<FPGALoopAttrT>(
9376
A, A.getNumArgs() ? A.getArgAsExpr(0) : nullptr);
9477
}
9578

96-
template <>
97-
Attr *handleIntelFPGALoopAttr<SYCLIntelFPGADisableLoopPipeliningAttr>(
98-
Sema &S, Stmt *St, const ParsedAttr &A) {
99-
if (!isa<ForStmt, CXXForRangeStmt, DoStmt, WhileStmt>(St)) {
100-
S.Diag(A.getLoc(), diag::err_attribute_wrong_decl_type_str)
101-
<< A << "'for', 'while', and 'do' statements";
102-
return nullptr;
103-
}
104-
79+
static Attr *handleSYCLIntelFPGADisableLoopPipeliningAttr(Sema &S, Stmt *,
80+
const ParsedAttr &A) {
10581
S.CheckDeprecatedSYCLAttributeSpelling(A);
106-
10782
return new (S.Context) SYCLIntelFPGADisableLoopPipeliningAttr(S.Context, A);
10883
}
10984

@@ -259,12 +234,6 @@ CheckRedundantSYCLIntelFPGAIVDepAttrs(Sema &S, ArrayRef<const Attr *> Attrs) {
259234
}
260235

261236
static Attr *handleIntelFPGAIVDepAttr(Sema &S, Stmt *St, const ParsedAttr &A) {
262-
if (!isa<ForStmt, CXXForRangeStmt, DoStmt, WhileStmt>(St)) {
263-
S.Diag(A.getLoc(), diag::err_attribute_wrong_decl_type_str)
264-
<< A << "'for', 'while', and 'do' statements";
265-
return nullptr;
266-
}
267-
268237
unsigned NumArgs = A.getNumArgs();
269238

270239
S.CheckDeprecatedSYCLAttributeSpelling(A);
@@ -276,12 +245,6 @@ static Attr *handleIntelFPGAIVDepAttr(Sema &S, Stmt *St, const ParsedAttr &A) {
276245

277246
static Attr *handleIntelFPGANofusionAttr(Sema &S, Stmt *St,
278247
const ParsedAttr &A) {
279-
if (!isa<ForStmt, CXXForRangeStmt, DoStmt, WhileStmt>(St)) {
280-
S.Diag(A.getLoc(), diag::err_attribute_wrong_decl_type_str)
281-
<< A << "'for', 'while', and 'do' statements";
282-
return nullptr;
283-
}
284-
285248
return new (S.Context) SYCLIntelFPGANofusionAttr(S.Context, A);
286249
}
287250

@@ -700,28 +663,13 @@ static Attr *handleLoopUnrollHint(Sema &S, Stmt *St, const ParsedAttr &A,
700663
// determines unrolling factor) or 1 argument (the unroll factor provided
701664
// by the user).
702665

703-
if (!isa<ForStmt, CXXForRangeStmt, DoStmt, WhileStmt>(St)) {
704-
S.Diag(A.getLoc(), diag::err_attribute_wrong_decl_type_str)
705-
<< A << "'for', 'while', and 'do' statements";
706-
return nullptr;
707-
}
708-
709666
Expr *E = A.getNumArgs() ? A.getArgAsExpr(0) : nullptr;
710667
if (A.getParsedKind() == ParsedAttr::AT_OpenCLUnrollHint)
711668
return S.BuildOpenCLLoopUnrollHintAttr(A, E);
712-
else if (A.getParsedKind() == ParsedAttr::AT_LoopUnrollHint) {
713-
// FIXME: this should be hoisted up to the top level, but can't be placed
714-
// there until the opencl attribute has its parsing error converted into a
715-
// semantic error. See: Parser::ParseOpenCLUnrollHintAttribute().
716-
if (!isa<ForStmt, CXXForRangeStmt, DoStmt, WhileStmt>(St)) {
717-
S.Diag(A.getLoc(), diag::err_attribute_wrong_decl_type_str)
718-
<< A << "'for', 'while', and 'do' statements";
719-
return nullptr;
720-
}
669+
if (A.getParsedKind() == ParsedAttr::AT_LoopUnrollHint)
721670
return S.BuildLoopUnrollHintAttr(A, E);
722-
}
723671

724-
return nullptr;
672+
llvm_unreachable("Unknown loop unroll hint");
725673
}
726674

727675
static Attr *ProcessStmtAttribute(Sema &S, Stmt *St, const ParsedAttr &A,
@@ -762,8 +710,7 @@ static Attr *ProcessStmtAttribute(Sema &S, Stmt *St, const ParsedAttr &A,
762710
case ParsedAttr::AT_SYCLIntelFPGALoopCoalesce:
763711
return handleIntelFPGALoopAttr<SYCLIntelFPGALoopCoalesceAttr>(S, St, A);
764712
case ParsedAttr::AT_SYCLIntelFPGADisableLoopPipelining:
765-
return handleIntelFPGALoopAttr<SYCLIntelFPGADisableLoopPipeliningAttr>(
766-
S, St, A);
713+
return handleSYCLIntelFPGADisableLoopPipeliningAttr(S, St, A);
767714
case ParsedAttr::AT_SYCLIntelFPGAMaxInterleaving:
768715
return handleIntelFPGALoopAttr<SYCLIntelFPGAMaxInterleavingAttr>(S, St, A);
769716
case ParsedAttr::AT_SYCLIntelFPGASpeculatedIterations:

‎clang/test/SemaSYCL/initiation_interval.cpp

+1-1
Original file line numberDiff line numberDiff line change
@@ -9,7 +9,7 @@
99

1010
[[intel::initiation_interval(-1)]] void func1() {} // expected-error{{'initiation_interval' attribute requires a positive integral compile time constant expression}}
1111

12-
[[intel::initiation_interval(0, 1)]] void func2() {} // expected-error{{'initiation_interval' attribute takes no more than 1 argument}}
12+
[[intel::initiation_interval(0, 1)]] void func2() {} // expected-error{{'initiation_interval' attribute takes one argument}}
1313

1414
// Tests for Intel FPGA initiation_interval function attribute duplication.
1515
// No diagnostic is emitted because the arguments match. Duplicate attribute is silently ignored.

‎clang/test/SemaSYCL/intel-fpga-loops.cpp

+6-6
Original file line numberDiff line numberDiff line change
@@ -75,10 +75,10 @@ void boo() {
7575
// expected-error@+1 {{duplicate argument to 'ivdep'; attribute requires one or both of a safelen and array}}
7676
[[intel::ivdep(2, 2)]] for (int i = 0; i != 10; ++i)
7777
a[i] = 0;
78-
// expected-warning@+1 {{'initiation_interval' attribute takes at least 1 argument; attribute ignored}}
78+
// expected-error@+1 {{'initiation_interval' attribute takes one argument}}
7979
[[intel::initiation_interval]] for (int i = 0; i != 10; ++i)
8080
a[i] = 0;
81-
// expected-error@+1 {{'initiation_interval' attribute takes no more than 1 argument}}
81+
// expected-error@+1 {{'initiation_interval' attribute takes one argument}}
8282
[[intel::initiation_interval(2, 2)]] for (int i = 0; i != 10; ++i)
8383
a[i] = 0;
8484
// expected-error@+1 {{'max_concurrency' attribute takes one argument}}
@@ -104,16 +104,16 @@ void boo() {
104104
// expected-error@+1 {{'loop_coalesce' attribute takes no more than 1 argument}}
105105
[[intel::loop_coalesce(2, 3)]] for (int i = 0; i != 10; ++i)
106106
a[i] = 0;
107-
// expected-warning@+1 {{'max_interleaving' attribute takes at least 1 argument; attribute ignored}}
107+
// expected-error@+1 {{'max_interleaving' attribute takes one argument}}
108108
[[intel::max_interleaving]] for (int i = 0; i != 10; ++i)
109109
a[i] = 0;
110-
// expected-error@+1 {{'max_interleaving' attribute takes no more than 1 argument}}
110+
// expected-error@+1 {{'max_interleaving' attribute takes one argument}}
111111
[[intel::max_interleaving(2, 4)]] for (int i = 0; i != 10; ++i)
112112
a[i] = 0;
113-
// expected-warning@+1 {{'speculated_iterations' attribute takes at least 1 argument; attribute ignored}}
113+
// expected-error@+1 {{'speculated_iterations' attribute takes one argument}}
114114
[[intel::speculated_iterations]] for (int i = 0; i != 10; ++i)
115115
a[i] = 0;
116-
// expected-error@+1 {{'speculated_iterations' attribute takes no more than 1 argument}}
116+
// expected-error@+1 {{'speculated_iterations' attribute takes one argument}}
117117
[[intel::speculated_iterations(1, 2)]] for (int i = 0; i != 10; ++i)
118118
a[i] = 0;
119119
// expected-error@+1 {{'nofusion' attribute takes no arguments}}

‎llvm/lib/SYCLLowerIR/LowerESIMD.cpp

+13
Original file line numberDiff line numberDiff line change
@@ -813,6 +813,12 @@ static Instruction *generateVectorGenXForSpirv(ExtractElementInst *EEI,
813813
Instruction *ExtrI = ExtractElementInst::Create(
814814
IntrI, ConstantInt::get(I32Ty, ExtractIndex), ExtractName, EEI);
815815
Instruction *CastI = addCastInstIfNeeded(EEI, ExtrI);
816+
if (EEI->getDebugLoc()) {
817+
IntrI->setDebugLoc(EEI->getDebugLoc());
818+
ExtrI->setDebugLoc(EEI->getDebugLoc());
819+
// It's OK if ExtrI and CastI is the same instruction
820+
CastI->setDebugLoc(EEI->getDebugLoc());
821+
}
816822
return CastI;
817823
}
818824

@@ -839,6 +845,11 @@ static Instruction *generateGenXForSpirv(ExtractElementInst *EEI,
839845
Instruction *IntrI =
840846
IntrinsicInst::Create(NewFDecl, {}, IntrinName + Suff.str(), EEI);
841847
Instruction *CastI = addCastInstIfNeeded(EEI, IntrI);
848+
if (EEI->getDebugLoc()) {
849+
IntrI->setDebugLoc(EEI->getDebugLoc());
850+
// It's OK if IntrI and CastI is the same instruction
851+
CastI->setDebugLoc(EEI->getDebugLoc());
852+
}
842853
return CastI;
843854
}
844855

@@ -1093,6 +1104,8 @@ static void translateESIMDIntrinsicCall(CallInst &CI) {
10931104
NewFDecl, GenXArgs,
10941105
NewFDecl->getReturnType()->isVoidTy() ? "" : CI.getName() + ".esimd",
10951106
&CI);
1107+
if (CI.getDebugLoc())
1108+
NewCI->setDebugLoc(CI.getDebugLoc());
10961109
NewCI = addCastInstIfNeeded(&CI, NewCI);
10971110
CI.replaceAllUsesWith(NewCI);
10981111
CI.eraseFromParent();

‎llvm/test/SYCLLowerIR/esimd_lower_debug_info.ll

+49
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,49 @@
1+
; NOTE: Assertions have been autogenerated by utils/update_test_checks.py
2+
; RUN: opt -debugify -LowerESIMD -S < %s | FileCheck %s
3+
4+
; This test checks that debug info is preserved during lowering
5+
; ESIMD specific constructs.
6+
7+
@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
8+
9+
declare spir_func <16 x float> @_Z18__esimd_block_readIfLi16EPU3AS1fEN2cl4sycl5INTEL3gpu11vector_typeIT_XT0_EE4typeET1_j(float addrspace(1)*, i32)
10+
11+
define spir_func void @func1(float addrspace(1)* %arg1, i32 %arg2 ){
12+
; CHECK-LABEL: @func1(
13+
; CHECK-NEXT: [[TMP1:%.*]] = ptrtoint float addrspace(1)* [[ARG1:%.*]] to i32, !dbg [[DBG11:![0-9]+]]
14+
; CHECK-NEXT: [[CALL1_I_I_ESIMD:%.*]] = call <16 x float> @llvm.genx.oword.ld.unaligned.v16f32(i32 0, i32 [[TMP1]], i32 [[ARG2:%.*]]), !dbg [[DBG11]]
15+
; CHECK-NEXT: call void @llvm.dbg.value(metadata <16 x float> [[CALL1_I_I_ESIMD]], metadata [[META9:![0-9]+]], metadata !DIExpression()), !dbg [[DBG11]]
16+
; CHECK-NEXT: ret void, !dbg [[DBG12:![0-9]+]]
17+
;
18+
%call1.i.i = tail call spir_func <16 x float> @_Z18__esimd_block_readIfLi16EPU3AS1fEN2cl4sycl5INTEL3gpu11vector_typeIT_XT0_EE4typeET1_j(float addrspace(1)* %arg1, i32 %arg2)
19+
ret void
20+
}
21+
22+
define spir_func void @func2(i64 addrspace(1)* %arg1) {
23+
; CHECK-LABEL: @func2(
24+
; CHECK-NEXT: call void @llvm.dbg.value(metadata <3 x i64> undef, metadata [[META15:![0-9]+]], metadata !DIExpression()), !dbg [[DBG21:![0-9]+]]
25+
; CHECK-NEXT: [[DOTESIMD:%.*]] = call <3 x i32> @llvm.genx.local.id.v3i32(), !dbg [[DBG22:![0-9]+]]
26+
; CHECK-NEXT: [[LOCAL_ID_X:%.*]] = extractelement <3 x i32> [[DOTESIMD]], i32 0, !dbg [[DBG22]]
27+
; CHECK-NEXT: [[LOCAL_ID_X_CAST_TY:%.*]] = zext i32 [[LOCAL_ID_X]] to i64, !dbg [[DBG22]]
28+
; CHECK-NEXT: [[DOTESIMD1:%.*]] = call <3 x i32> @llvm.genx.local.size.v3i32(), !dbg [[DBG22]]
29+
; CHECK-NEXT: [[WGSIZE_X:%.*]] = extractelement <3 x i32> [[DOTESIMD1]], i32 0, !dbg [[DBG22]]
30+
; CHECK-NEXT: [[WGSIZE_X_CAST_TY:%.*]] = zext i32 [[WGSIZE_X]] to i64, !dbg [[DBG22]]
31+
; CHECK-NEXT: [[GROUP_ID_X:%.*]] = call i32 @llvm.genx.group.id.x(), !dbg [[DBG22]]
32+
; CHECK-NEXT: [[GROUP_ID_X_CAST_TY:%.*]] = zext i32 [[GROUP_ID_X]] to i64, !dbg [[DBG22]]
33+
; CHECK-NEXT: [[MUL:%.*]] = mul i64 [[WGSIZE_X_CAST_TY]], [[GROUP_ID_X_CAST_TY]]
34+
; CHECK-NEXT: [[ADD:%.*]] = add i64 [[LOCAL_ID_X_CAST_TY]], [[MUL]]
35+
; CHECK-NEXT: call void @llvm.dbg.value(metadata i64 [[ADD]], metadata [[META17:![0-9]+]], metadata !DIExpression()), !dbg [[DBG22]]
36+
; CHECK-NEXT: [[PTRIDX_I_I:%.*]] = getelementptr inbounds i64, i64 addrspace(1)* [[ARG1:%.*]], i64 2, !dbg [[DBG23:![0-9]+]]
37+
; CHECK-NEXT: call void @llvm.dbg.value(metadata i64 addrspace(1)* [[PTRIDX_I_I]], metadata [[META19:![0-9]+]], metadata !DIExpression()), !dbg [[DBG23]]
38+
; CHECK-NEXT: [[PTRIDX_ASCAST_I_I:%.*]] = addrspacecast i64 addrspace(1)* [[PTRIDX_I_I]] to i64 addrspace(4)*, !dbg [[DBG24:![0-9]+]]
39+
; CHECK-NEXT: call void @llvm.dbg.value(metadata i64 addrspace(4)* [[PTRIDX_ASCAST_I_I]], metadata [[META20:![0-9]+]], metadata !DIExpression()), !dbg [[DBG24]]
40+
; CHECK-NEXT: store i64 [[ADD]], i64 addrspace(4)* [[PTRIDX_ASCAST_I_I]], align 4, !dbg [[DBG25:![0-9]+]]
41+
; CHECK-NEXT: ret void, !dbg [[DBG26:![0-9]+]]
42+
;
43+
%1 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*)
44+
%2 = extractelement <3 x i64> %1, i64 0
45+
%ptridx.i.i = getelementptr inbounds i64, i64 addrspace(1)* %arg1, i64 2
46+
%ptridx.ascast.i.i = addrspacecast i64 addrspace(1)* %ptridx.i.i to i64 addrspace(4)*
47+
store i64 %2, i64 addrspace(4)* %ptridx.ascast.i.i
48+
ret void
49+
}

‎sycl/include/CL/sycl/INTEL/esimd/detail/esimd_host_util.hpp

+17-22
Original file line numberDiff line numberDiff line change
@@ -14,19 +14,19 @@
1414

1515
#define SIMDCF_ELEMENT_SKIP(i)
1616

17-
namespace cl {
17+
__SYCL_INLINE_NAMESPACE(cl) {
1818
namespace sycl {
19+
1920
namespace detail {
2021
namespace half_impl {
2122
class half;
2223
} // namespace half_impl
2324
} // namespace detail
24-
} // namespace sycl
25-
} // namespace cl
2625

27-
using half = cl::sycl::detail::half_impl::half;
28-
29-
namespace EsimdEmulSys {
26+
namespace INTEL {
27+
namespace gpu {
28+
namespace emu {
29+
namespace detail {
3030

3131
constexpr int sat_is_on = 1;
3232

@@ -44,14 +44,10 @@ template <typename RT> struct satur {
4444
return (RT)val;
4545
}
4646

47-
#ifdef max
48-
#undef max
49-
#endif
50-
#ifdef min
51-
#undef min
52-
#endif
53-
const RT t_max = std::numeric_limits<RT>::max();
54-
const RT t_min = std::numeric_limits<RT>::min();
47+
// min/max can be macros on Windows, so wrap them into parens to avoid their
48+
// expansion
49+
const RT t_max = (std::numeric_limits<RT>::max)();
50+
const RT t_min = (std::numeric_limits<RT>::min)();
5551

5652
if (val > t_max) {
5753
return t_max;
@@ -112,8 +108,6 @@ template <> struct SetSatur<double, true> {
112108
static unsigned int set() { return sat_is_on; }
113109
};
114110

115-
} // namespace EsimdEmulSys
116-
117111
// used for intermediate type in dp4a emulation
118112
template <typename T1, typename T2> struct restype_ex {
119113
private:
@@ -470,10 +464,11 @@ template <typename T> struct dwordtype;
470464
template <> struct dwordtype<int> { static const bool value = true; };
471465
template <> struct dwordtype<unsigned int> { static const bool value = true; };
472466

473-
template <unsigned int N1, unsigned int N2> struct ressize {
474-
static const unsigned int size = (N1 > N2) ? N1 : N2;
475-
static const bool conformable =
476-
check_true < N1 % size == 0 && N2 % size == 0 > ::value;
477-
};
467+
} // namespace detail
468+
} // namespace emu
469+
} // namespace gpu
470+
} // namespace INTEL
471+
} // namespace sycl
472+
} // __SYCL_INLINE_NAMESPACE(cl)
478473

479-
#endif
474+
#endif // #ifndef __SYCL_DEVICE_ONLY__

0 commit comments

Comments
 (0)