Skip to content

Commit 67403d1

Browse files
vinser52vinser52
committed
Fix backward compatibility test
1 parent 4ee9054 commit 67403d1

File tree

5 files changed

+75
-29
lines changed

5 files changed

+75
-29
lines changed

‎sycl/source/detail/device_kernel_info.cpp‎

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -63,7 +63,6 @@ void DeviceKernelInfo::setCompileTimeInfoIfNeeded(
6363
const CompileTimeKernelInfoTy &Info) {
6464
if (!isCompileTimeInfoSet())
6565
CompileTimeKernelInfoTy::operator=(Info);
66-
assert(isCompileTimeInfoSet());
6766
assert(Info == *this);
6867
}
6968

‎sycl/source/detail/handler_impl.hpp‎

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -223,6 +223,15 @@ class handler_impl {
223223
// Allocation ptr to be freed asynchronously.
224224
void *MFreePtr = nullptr;
225225

226+
#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
227+
// TODO: remove in the next ABI-breaking window
228+
// Today they are used only in the handler::setKernelNameBasedCachePtr
229+
int MKernelNumArgs = 0;
230+
detail::kernel_param_desc_t (*MKernelParamDescGetter)(int) = nullptr;
231+
bool MKernelIsESIMD = false;
232+
bool MKernelHasSpecialCaptures = true;
233+
#endif
234+
226235
KernelData MKernelData;
227236
};
228237

‎sycl/source/detail/kernel_data.cpp‎

Lines changed: 27 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -103,8 +103,15 @@ static void addArgsForLocalAccessor(detail::LocalAccessorImplHost *LAcc,
103103

104104
void KernelData::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind,
105105
const int Size, const size_t Index,
106-
size_t &IndexShift,
107-
bool IsKernelCreatedFromSource) {
106+
size_t &IndexShift, bool IsKernelCreatedFromSource
107+
#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
108+
,
109+
bool IsESIMD
110+
#endif
111+
) {
112+
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
113+
bool IsESIMD = isESIMD();
114+
#endif
108115
using detail::kernel_param_kind_t;
109116
size_t GlobalSize = MNDRDesc.GlobalSize[0];
110117
for (size_t I = 1; I < MNDRDesc.Dims; ++I) {
@@ -126,14 +133,14 @@ void KernelData::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind,
126133
detail::Requirement *GBufReq = &*detail::getSyclObjImpl(*GBufBase);
127134
addArgsForGlobalAccessor(GBufReq, Index, IndexShift, Size,
128135
IsKernelCreatedFromSource, GlobalSize, MArgs,
129-
isESIMD());
136+
IsESIMD);
130137
++IndexShift;
131138
detail::AccessorBaseHost *GOffsetBase =
132139
static_cast<detail::AccessorBaseHost *>(&S->GlobalOffset);
133140
detail::Requirement *GOffsetReq = &*detail::getSyclObjImpl(*GOffsetBase);
134141
addArgsForGlobalAccessor(GOffsetReq, Index, IndexShift, Size,
135142
IsKernelCreatedFromSource, GlobalSize, MArgs,
136-
isESIMD());
143+
IsESIMD);
137144
++IndexShift;
138145
detail::AccessorBaseHost *GFlushBase =
139146
static_cast<detail::AccessorBaseHost *>(&S->GlobalFlushBuf);
@@ -152,7 +159,7 @@ void KernelData::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind,
152159
}
153160
addArgsForGlobalAccessor(GFlushReq, Index, IndexShift, Size,
154161
IsKernelCreatedFromSource, GlobalSize, MArgs,
155-
isESIMD());
162+
IsESIMD);
156163
++IndexShift;
157164
addArg(kernel_param_kind_t::kind_std_layout, &S->FlushBufferSize,
158165
sizeof(S->FlushBufferSize), Index + IndexShift);
@@ -170,15 +177,15 @@ void KernelData::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind,
170177
detail::Requirement *AccImpl = static_cast<detail::Requirement *>(Ptr);
171178
addArgsForGlobalAccessor(AccImpl, Index, IndexShift, Size,
172179
IsKernelCreatedFromSource, GlobalSize, MArgs,
173-
isESIMD());
180+
IsESIMD);
174181
break;
175182
}
176183
case access::target::local: {
177184
detail::LocalAccessorImplHost *LAccImpl =
178185
static_cast<detail::LocalAccessorImplHost *>(Ptr);
179186

180187
addArgsForLocalAccessor(LAccImpl, Index, IndexShift,
181-
IsKernelCreatedFromSource, MArgs, isESIMD());
188+
IsKernelCreatedFromSource, MArgs, IsESIMD);
182189
break;
183190
}
184191
case access::target::image:
@@ -221,7 +228,7 @@ void KernelData::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind,
221228

222229
addArgsForLocalAccessor(&DynLocalAccessorImpl->LAccImplHost, Index,
223230
IndexShift, IsKernelCreatedFromSource, MArgs,
224-
isESIMD());
231+
IsESIMD);
225232
break;
226233
}
227234
default: {
@@ -289,7 +296,12 @@ void KernelData::extractArgsAndReqs(bool IsKernelCreatedFromSource) {
289296
const detail::kernel_param_kind_t &Kind = UnPreparedArgs[I].MType;
290297
const int &Size = UnPreparedArgs[I].MSize;
291298
const int Index = UnPreparedArgs[I].MIndex;
292-
processArg(Ptr, Kind, Size, Index, IndexShift, IsKernelCreatedFromSource);
299+
processArg(Ptr, Kind, Size, Index, IndexShift, IsKernelCreatedFromSource
300+
#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
301+
,
302+
isESIMD()
303+
#endif
304+
);
293305
}
294306
}
295307

@@ -341,7 +353,12 @@ void KernelData::extractArgsAndReqsFromLambda() {
341353
}
342354

343355
processArg(Ptr, Kind, Size, I, IndexShift,
344-
/*IsKernelCreatedFromSource=*/false);
356+
/*IsKernelCreatedFromSource=*/false
357+
#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
358+
,
359+
isESIMD()
360+
#endif
361+
);
345362
}
346363
}
347364

‎sycl/source/detail/kernel_data.hpp‎

Lines changed: 6 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -169,7 +169,12 @@ class KernelData {
169169

170170
void processArg(void *Ptr, const detail::kernel_param_kind_t &Kind,
171171
const int Size, const size_t Index, size_t &IndexShift,
172-
bool IsKernelCreatedFromSource);
172+
bool IsKernelCreatedFromSource
173+
#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
174+
,
175+
bool IsESIMD
176+
#endif
177+
);
173178

174179
void extractArgsAndReqs(bool IsKernelCreatedFromSource);
175180

‎sycl/source/handler.cpp‎

Lines changed: 33 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -539,8 +539,14 @@ event handler::finalize() {
539539
}
540540

541541
if (type == detail::CGType::Kernel) {
542-
assert(impl->MKernelData.getDeviceKernelInfoPtr() != nullptr &&
543-
"DeviceKernelInfo pointer must be set in handler_impl");
542+
if (impl->MKernelData.getDeviceKernelInfoPtr() == nullptr) {
543+
// Fetch the device kernel info pointer if it hasn't been set (e.g.
544+
// in kernel bundle or free function cases).
545+
impl->MKernelData.setDeviceKernelInfoPtr(
546+
&detail::ProgramManager::getInstance().getOrCreateDeviceKernelInfo(
547+
toKernelNameStrT(MKernelName)));
548+
}
549+
assert(impl->MKernelData.getKernelName() == MKernelName);
544550

545551
// If there were uses of set_specialization_constant build the kernel_bundle
546552
detail::kernel_bundle_impl *KernelBundleImpPtr =
@@ -1049,9 +1055,8 @@ void handler::associateWithHandler(
10491055
void handler::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind,
10501056
const int Size, const size_t Index, size_t &IndexShift,
10511057
bool IsKernelCreatedFromSource, bool IsESIMD) {
1052-
(void)IsESIMD;
10531058
impl->MKernelData.processArg(Ptr, Kind, Size, Index, IndexShift,
1054-
IsKernelCreatedFromSource);
1059+
IsKernelCreatedFromSource, IsESIMD);
10551060
}
10561061
#endif
10571062

@@ -1070,6 +1075,14 @@ void handler::setArgHelper(int ArgIndex, stream &&Str) {
10701075

10711076
void handler::extractArgsAndReqs() {
10721077
assert(MKernel && "MKernel is not initialized");
1078+
#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
1079+
if (impl->MKernelData.getDeviceKernelInfoPtr() == nullptr) {
1080+
impl->MKernelData.setDeviceKernelInfoPtr(
1081+
&detail::ProgramManager::getInstance().getOrCreateDeviceKernelInfo(
1082+
toKernelNameStrT(getKernelName())));
1083+
}
1084+
#endif
1085+
assert(impl->MKernelData.getDeviceKernelInfoPtr() != nullptr);
10731086
impl->MKernelData.extractArgsAndReqs(MKernel->isCreatedFromSource());
10741087
}
10751088

@@ -1082,7 +1095,7 @@ void handler::extractArgsAndReqsFromLambda(
10821095
if (impl->MKernelData.getDeviceKernelInfoPtr() == nullptr) {
10831096
impl->MKernelData.setDeviceKernelInfoPtr(
10841097
&detail::ProgramManager::getInstance().getOrCreateDeviceKernelInfo(
1085-
toKernelNameStrT(MKernelName)));
1098+
toKernelNameStrT(getKernelName())));
10861099
}
10871100
impl->MKernelData.setKernelInfo(LambdaPtr, NumKernelParams, ParamDescGetter,
10881101
IsESIMD, true);
@@ -1094,7 +1107,6 @@ void handler::extractArgsAndReqsFromLambda(
10941107
bool IsESIMD) {
10951108
const bool IsKernelCreatedFromSource = false;
10961109
size_t IndexShift = 0;
1097-
impl->MKernelData.setESIMD(IsESIMD);
10981110

10991111
for (size_t I = 0; I < ParamDescs.size(); ++I) {
11001112
void *Ptr = LambdaPtr + ParamDescs[I].offset;
@@ -1119,7 +1131,7 @@ void handler::extractArgsAndReqsFromLambda(
11191131
}
11201132
}
11211133
impl->MKernelData.processArg(Ptr, Kind, Size, I, IndexShift,
1122-
IsKernelCreatedFromSource);
1134+
IsKernelCreatedFromSource, IsESIMD);
11231135
}
11241136
}
11251137

@@ -2314,23 +2326,27 @@ void handler::setNDRangeDescriptor(sycl::range<1> NumWorkItems,
23142326
void handler::setKernelNameBasedCachePtr(
23152327
sycl::detail::KernelNameBasedCacheT *KernelNameBasedCachePtr) {
23162328
assert(!impl->MKernelData.getDeviceKernelInfoPtr() && "Already set!");
2329+
(void)KernelNameBasedCachePtr;
2330+
CompileTimeKernelInfoTy HandlerInfo;
2331+
HandlerInfo.Name = MKernelName;
2332+
HandlerInfo.NumParams = impl->MKernelNumArgs;
2333+
HandlerInfo.ParamDescGetter = impl->MKernelParamDescGetter;
2334+
HandlerInfo.IsESIMD = impl->MKernelIsESIMD;
2335+
HandlerInfo.HasSpecialCaptures = impl->MKernelHasSpecialCaptures;
23172336
impl->MKernelData.setDeviceKernelInfoPtr(
2318-
reinterpret_cast<sycl::detail::DeviceKernelInfo *>(
2319-
KernelNameBasedCachePtr));
2337+
&detail::ProgramManager::getInstance().getOrCreateDeviceKernelInfo(
2338+
HandlerInfo));
23202339
}
23212340

23222341
void handler::setKernelInfo(
23232342
void *KernelFuncPtr, int KernelNumArgs,
23242343
detail::kernel_param_desc_t (*KernelParamDescGetter)(int),
23252344
bool KernelIsESIMD, bool KernelHasSpecialCaptures) {
2326-
if (impl->MKernelData.getDeviceKernelInfoPtr() == nullptr) {
2327-
impl->MKernelData.setDeviceKernelInfoPtr(
2328-
&detail::ProgramManager::getInstance().getOrCreateDeviceKernelInfo(
2329-
toKernelNameStrT(MKernelName)));
2330-
}
2331-
impl->MKernelData.setKernelInfo(KernelFuncPtr, KernelNumArgs,
2332-
KernelParamDescGetter, KernelIsESIMD,
2333-
KernelHasSpecialCaptures);
2345+
impl->MKernelData.setKernelInfo(KernelFuncPtr);
2346+
impl->MKernelNumArgs = KernelNumArgs;
2347+
impl->MKernelParamDescGetter = KernelParamDescGetter;
2348+
impl->MKernelIsESIMD = KernelIsESIMD;
2349+
impl->MKernelHasSpecialCaptures = KernelHasSpecialCaptures;
23342350
}
23352351
#endif
23362352

0 commit comments

Comments
 (0)