Skip to content

sycl: fix error building scheduler.cpp #9

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Closed
wants to merge 1 commit into from
Closed

sycl: fix error building scheduler.cpp #9

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
10 changes: 5 additions & 5 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -636,15 +636,15 @@ static void buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj,

CreateAndAddPrmDsc(Fld, PointerType);

FieldDecl *RangeFld = getFieldDeclByName(RecordDecl, {"__impl", "Range"});
FieldDecl *RangeFld = getFieldDeclByName(RecordDecl, {"__implx", "Range"});
assert(RangeFld &&
"The accessor must contain the Range from the __impl field");
CreateAndAddPrmDsc(RangeFld, RangeFld->getType());

FieldDecl *OffsetFld =
getFieldDeclByName(RecordDecl, {"__impl", "Offset"});
getFieldDeclByName(RecordDecl, {"__implx", "Offset"});
assert(OffsetFld &&
"The accessor must contain the Offset from the __impl field");
"The accessor must contain the Offset from the __implx field");
CreateAndAddPrmDsc(OffsetFld, OffsetFld->getType());
} else if (Util::isSyclStreamType(ArgTy)) {
// the parameter is a SYCL stream object
Expand Down Expand Up @@ -700,15 +700,15 @@ static void populateIntHeader(SYCLIntegrationHeader &H, const StringRef Name,
getAccessTarget(AccTmplTy), Offset);
// ... second descriptor (translated to range kernel parameter):
FieldDecl *RngFld =
getFieldDeclByName(AccTy, {"__impl", "Range"}, &Offset);
getFieldDeclByName(AccTy, {"__implx", "Range"}, &Offset);
uint64_t Sz = Ctx.getTypeSizeInChars(RngFld->getType()).getQuantity();
H.addParamDesc(SYCLIntegrationHeader::kind_std_layout,
static_cast<unsigned>(Sz), static_cast<unsigned>(Offset));
// ... third descriptor (translated to id kernel parameter):
// Get offset in bytes
Offset = Layout.getFieldOffset(Fld->getFieldIndex()) / 8;
FieldDecl *OffstFld =
getFieldDeclByName(AccTy, {"__impl", "Offset"}, &Offset);
getFieldDeclByName(AccTy, {"__implx", "Offset"}, &Offset);
Sz = Ctx.getTypeSizeInChars(OffstFld->getType()).getQuantity();
H.addParamDesc(SYCLIntegrationHeader::kind_std_layout,
static_cast<unsigned>(Sz), static_cast<unsigned>(Offset));
Expand Down
51 changes: 37 additions & 14 deletions sycl/include/CL/sycl/accessor.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -355,6 +355,9 @@ SYCL_ACCESSOR_SUBCLASS(accessor_common, accessor_base, true /* always */) {
template <int Dimensions = dimensions>
typename std::enable_if<(Dimensions > 0), id<Dimensions>>::type
get_offset() const { return this->__impl()->Offset; }

friend class ::cl::sycl::simple_scheduler::Node;
friend class ::cl::sycl::simple_scheduler::Scheduler;
};

SYCL_ACCESSOR_SUBCLASS(accessor_opdata_w, accessor_common,
Expand All @@ -366,6 +369,8 @@ SYCL_ACCESSOR_SUBCLASS(accessor_opdata_w, accessor_common,
operator dataT &() const {
return this->__impl()->Data[0];
}
friend class ::cl::sycl::simple_scheduler::Node;
friend class ::cl::sycl::simple_scheduler::Scheduler;
};

SYCL_ACCESSOR_SUBCLASS(accessor_subscript_wn, accessor_opdata_w,
Expand All @@ -387,6 +392,8 @@ SYCL_ACCESSOR_SUBCLASS(accessor_subscript_wn, accessor_opdata_w,
return subscript_obj<dimensions, dataT, dimensions - 1, accessMode,
accessTarget, isPlaceholder>(*this, ids);
}
friend class ::cl::sycl::simple_scheduler::Node;
friend class ::cl::sycl::simple_scheduler::Scheduler;
};

SYCL_ACCESSOR_SUBCLASS(accessor_subscript_w, accessor_subscript_wn,
Expand All @@ -408,6 +415,8 @@ SYCL_ACCESSOR_SUBCLASS(accessor_subscript_w, accessor_subscript_wn,
dataT &operator[](size_t index) const {
return this->__impl()->Data[index];
}
friend class ::cl::sycl::simple_scheduler::Node;
friend class ::cl::sycl::simple_scheduler::Scheduler;
};

SYCL_ACCESSOR_SUBCLASS(accessor_opdata_r, accessor_subscript_w,
Expand All @@ -416,6 +425,8 @@ SYCL_ACCESSOR_SUBCLASS(accessor_opdata_r, accessor_subscript_w,
operator PureType() const {
return this->__impl()->Data[0];
}
friend class ::cl::sycl::simple_scheduler::Node;
friend class ::cl::sycl::simple_scheduler::Scheduler;
};

SYCL_ACCESSOR_SUBCLASS(accessor_subscript_rn, accessor_opdata_r,
Expand All @@ -434,6 +445,8 @@ SYCL_ACCESSOR_SUBCLASS(accessor_subscript_rn, accessor_opdata_r,
return subscript_obj<dimensions, dataT, dimensions - 1, accessMode,
accessTarget, isPlaceholder>(*this, ids);
}
friend class ::cl::sycl::simple_scheduler::Node;
friend class ::cl::sycl::simple_scheduler::Scheduler;
};

SYCL_ACCESSOR_SUBCLASS(accessor_subscript_r, accessor_subscript_rn,
Expand All @@ -447,6 +460,8 @@ SYCL_ACCESSOR_SUBCLASS(accessor_subscript_r, accessor_subscript_rn,
operator[](size_t index) const {
return this->__impl()->Data[index];
}
friend class ::cl::sycl::simple_scheduler::Node;
friend class ::cl::sycl::simple_scheduler::Scheduler;
};

template <access::target accessTarget> struct getAddressSpace {
Expand All @@ -469,6 +484,8 @@ SYCL_ACCESSOR_SUBCLASS(accessor_subscript_atomic_eq0, accessor_subscript_r,
return atomic<PureType, addressSpace>(
multi_ptr<PureType, addressSpace>(&(this->__impl()->Data[0])));
}
friend class ::cl::sycl::simple_scheduler::Node;
friend class ::cl::sycl::simple_scheduler::Scheduler;
};

// Available when: accessMode == access::mode::atomic && dimensions > 0
Expand All @@ -483,6 +500,8 @@ SYCL_ACCESSOR_SUBCLASS(accessor_subscript_atomic_gt0,
multi_ptr<PureType, addressSpace>(&(this->__impl()->Data[getOffsetForId(
this->__impl()->Range, index, this->__impl()->Offset)])));
}
friend class ::cl::sycl::simple_scheduler::Node;
friend class ::cl::sycl::simple_scheduler::Scheduler;
};

// Available only when: accessMode == access::mode::atomic && dimensions == 1
Expand All @@ -496,6 +515,8 @@ SYCL_ACCESSOR_SUBCLASS(accessor_subscript_atomic_eq1,
return atomic<PureType, addressSpace>(
multi_ptr<PureType, addressSpace>(&(this->__impl()->Data[index])));
}
friend class ::cl::sycl::simple_scheduler::Node;
friend class ::cl::sycl::simple_scheduler::Scheduler;
};

// TODO:
Expand Down Expand Up @@ -535,6 +556,8 @@ SYCL_ACCESSOR_SUBCLASS(accessor_pointer, accessor_subscript_atomic_eq1, true) {
get_pointer() const {
return local_ptr<DataT>(this->__impl()->Data);
}
friend class ::cl::sycl::simple_scheduler::Node;
friend class ::cl::sycl::simple_scheduler::Scheduler;
};

} // namespace detail
Expand All @@ -557,13 +580,13 @@ class accessor
// Make sure Impl field is the first in the class, so that it is
// safe to reinterpret a pointer to accessor as a pointer to the
// implementation.
_ImplT __impl;
_ImplT __implx;

void __init(_ValueType *Ptr, range<dimensions> Range,
id<dimensions> Offset) {
__impl.Data = Ptr;
__impl.Range = Range;
__impl.Offset = Offset;
__implx.Data = Ptr;
__implx.Range = Range;
__implx.Offset = Offset;
}

public:
Expand Down Expand Up @@ -593,7 +616,7 @@ class accessor
AccessTarget == access::target::constant_buffer))) &&
Dimensions == 0),
buffer<DataT, 1>>::type &bufferRef)
: __impl(detail::getSyclObjImpl(bufferRef)->BufPtr) {
: __implx(detail::getSyclObjImpl(bufferRef)->BufPtr) {
auto BufImpl = detail::getSyclObjImpl(bufferRef);
if (AccessTarget == access::target::host_buffer) {
if (BufImpl->OpenCLInterop) {
Expand Down Expand Up @@ -633,7 +656,7 @@ class accessor
#ifdef __SYCL_DEVICE_ONLY__
; // This ctor can't be used in device code, so no need to define it.
#else // !__SYCL_DEVICE_ONLY__
: __impl(detail::getSyclObjImpl(bufferRef)->BufPtr,
: __implx(detail::getSyclObjImpl(bufferRef)->BufPtr,
detail::getSyclObjImpl(bufferRef)->Range,
&commandGroupHandlerRef) {
auto BufImpl = detail::getSyclObjImpl(bufferRef);
Expand All @@ -643,7 +666,7 @@ class accessor
"interoperability buffer");
}
commandGroupHandlerRef.AddBufDep<AccessMode, AccessTarget>(*BufImpl);
__impl.m_Buf = BufImpl.get();
__implx.m_Buf = BufImpl.get();
}
#endif // !__SYCL_DEVICE_ONLY__

Expand All @@ -669,7 +692,7 @@ class accessor
AccessTarget == access::target::constant_buffer))) &&
Dimensions > 0),
buffer<DataT, Dimensions>>::type &bufferRef)
: __impl(detail::getSyclObjImpl(bufferRef)->BufPtr,
: __implx(detail::getSyclObjImpl(bufferRef)->BufPtr,
detail::getSyclObjImpl(bufferRef)->Range) {
auto BufImpl = detail::getSyclObjImpl(bufferRef);
if (AccessTarget == access::target::host_buffer) {
Expand Down Expand Up @@ -710,7 +733,7 @@ class accessor
#ifdef __SYCL_DEVICE_ONLY__
; // This ctor can't be used in device code, so no need to define it.
#else
: __impl(detail::getSyclObjImpl(bufferRef)->BufPtr,
: __implx(detail::getSyclObjImpl(bufferRef)->BufPtr,
detail::getSyclObjImpl(bufferRef)->Range,
&commandGroupHandlerRef) {
auto BufImpl = detail::getSyclObjImpl(bufferRef);
Expand All @@ -720,7 +743,7 @@ class accessor
"interoperability buffer");
}
commandGroupHandlerRef.AddBufDep<AccessMode, AccessTarget>(*BufImpl);
__impl.m_Buf = BufImpl.get();
__implx.m_Buf = BufImpl.get();
}
#endif

Expand Down Expand Up @@ -752,7 +775,7 @@ class accessor
#ifdef __SYCL_DEVICE_ONLY__
; // This ctor can't be used in device code, so no need to define it.
#else // !__SYCL_DEVICE_ONLY__
: __impl(detail::getSyclObjImpl(bufferRef)->BufPtr, Range, Offset) {
: __implx(detail::getSyclObjImpl(bufferRef)->BufPtr, Range, Offset) {
auto BufImpl = detail::getSyclObjImpl(bufferRef);
if (AccessTarget == access::target::host_buffer) {
if (BufImpl->OpenCLInterop) {
Expand Down Expand Up @@ -796,7 +819,7 @@ class accessor
#ifdef __SYCL_DEVICE_ONLY__
; // This ctor can't be used in device code, so no need to define it.
#else // !__SYCL_DEVICE_ONLY__
: __impl(detail::getSyclObjImpl(bufferRef)->BufPtr, Range,
: __implx(detail::getSyclObjImpl(bufferRef)->BufPtr, Range,
&commandGroupHandlerRef, Offset) {
auto BufImpl = detail::getSyclObjImpl(bufferRef);
if (BufImpl->OpenCLInterop && !BufImpl->isValidAccessToMem(accessMode)) {
Expand All @@ -805,7 +828,7 @@ class accessor
"interoperability buffer");
}
commandGroupHandlerRef.AddBufDep<AccessMode, AccessTarget>(*BufImpl);
__impl.m_Buf = BufImpl.get();
__implx.m_Buf = BufImpl.get();
}
#endif // !__SYCL_DEVICE_ONLY__

Expand Down Expand Up @@ -835,7 +858,7 @@ class accessor
Dimensions > 0),
range<Dimensions>>::type allocationSize,
handler &commandGroupHandlerRef)
: __impl(allocationSize, &commandGroupHandlerRef) {}
: __implx(allocationSize, &commandGroupHandlerRef) {}
};

} // namespace sycl
Expand Down
18 changes: 6 additions & 12 deletions sycl/include/CL/sycl/detail/scheduler/scheduler.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -61,9 +61,7 @@ void Node::addAccRequirement(
accessor<dataT, dimensions, accessMode, accessTarget, isPlaceholder> &&Acc,
int argIndex) {
detail::buffer_impl<dataT, dimensions> *buf =
Acc.template accessor_base<dataT, dimensions, accessMode, accessTarget,
isPlaceholder>::__impl()
->m_Buf;
Acc.__impl()->m_Buf;
addBufRequirement<accessMode, accessTarget, dataT, dimensions>(*buf);
addInteropArg(nullptr, buf->get_size(), argIndex,
getReqForBuffer(m_Bufs, *buf));
Expand Down Expand Up @@ -128,8 +126,7 @@ template <typename T, int Dimensions, access::mode mode, access::target tgt,
access::placeholder isPlaceholder>
void Node::addExplicitMemOp(
accessor<T, Dimensions, mode, tgt, isPlaceholder> &Dest, T Src) {
auto *DestBase = Dest.template accessor_base<T, Dimensions, mode, tgt,
isPlaceholder>::__impl();
auto *DestBase = Dest.__impl();
assert(DestBase != nullptr &&
"Accessor should have an initialized accessor_base");
detail::buffer_impl<T, Dimensions> *Buf = DestBase->m_Buf;
Expand All @@ -153,13 +150,10 @@ template <typename T_src, int dim_src, access::mode mode_src,
void Node::addExplicitMemOp(
accessor<T_src, dim_src, mode_src, tgt_src, isPlaceholder_src> Src,
accessor<T_dest, dim_dest, mode_dest, tgt_dest, isPlaceholder_dest> Dest) {
auto *SrcBase = Src.template accessor_base<T_src, dim_src, mode_src, tgt_src,
isPlaceholder_src>::__impl();
auto *SrcBase = Src.__impl();
assert(SrcBase != nullptr &&
"Accessor should have an initialized accessor_base");
auto *DestBase =
Dest.template accessor_base<T_dest, dim_dest, mode_dest, tgt_dest,
isPlaceholder_dest>::__impl();
auto *DestBase = Dest.__impl();
assert(DestBase != nullptr &&
"Accessor should have an initialized accessor_base");

Expand Down Expand Up @@ -191,8 +185,8 @@ template <typename T, int Dimensions, access::mode mode, access::target tgt,
void Scheduler::updateHost(
accessor<T, Dimensions, mode, tgt, isPlaceholder> &Acc,
cl::sycl::event &Event) {
auto *AccBase = Acc.template accessor_base<T, Dimensions, mode, tgt,
isPlaceholder>::__impl();
auto *AccBase = Acc.impl();

assert(AccBase != nullptr &&
"Accessor should have an initialized accessor_base");
detail::buffer_impl<T, Dimensions> *Buf = AccBase->m_Buf;
Expand Down