Skip to content

Commit 4bb7473

Browse files
ldorauldorau
committed
Revert "Revert "[UR][SYCL] Introduce UR api to set kernel args + launch in one call." (#19661)"
This reverts commit d25d6d6. Signed-off-by: Lukasz Dorau <lukasz.dorau@intel.com>
1 parent a963e89 commit 4bb7473

File tree

86 files changed

+3393
-451
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

86 files changed

+3393
-451
lines changed

‎sycl/source/detail/queue_impl.cpp‎

Lines changed: 8 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -467,13 +467,14 @@ EventImplPtr queue_impl::submit_kernel_scheduler_bypass(
467467
BinImage = detail::retrieveKernelBinary(*this, KData.getKernelName());
468468
assert(BinImage && "Failed to obtain a binary image.");
469469
}
470-
enqueueImpKernel(*this, KData.getNDRDesc(), KData.getArgs(),
471-
KernelBundleImpPtr, KernelImplPtr,
472-
*KData.getDeviceKernelInfoPtr(), RawEvents,
473-
ResultEvent.get(), nullptr, KData.getKernelCacheConfig(),
474-
KData.isCooperative(), KData.usesClusterLaunch(),
475-
KData.getKernelWorkGroupMemorySize(), BinImage,
476-
KData.getKernelFuncPtr());
470+
enqueueImpKernel(
471+
*this, KData.getNDRDesc(), KData.getArgs(), KernelBundleImpPtr,
472+
KernelImplPtr, *KData.getDeviceKernelInfoPtr(), RawEvents,
473+
ResultEvent.get(), nullptr, KData.getKernelCacheConfig(),
474+
KData.isCooperative(), KData.usesClusterLaunch(),
475+
KData.getKernelWorkGroupMemorySize(), BinImage,
476+
KData.getKernelFuncPtr(), KData.getKernelNumArgs(),
477+
KData.getKernelParamDescGetter(), KData.hasSpecialCaptures());
477478
#ifdef XPTI_ENABLE_INSTRUMENTATION
478479
if (xptiEnabled) {
479480
// Emit signal only when event is created

‎sycl/source/detail/scheduler/commands.cpp‎

Lines changed: 177 additions & 66 deletions
Large diffs are not rendered by default.

‎sycl/source/detail/scheduler/commands.hpp‎

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -633,7 +633,9 @@ void enqueueImpKernel(
633633
ur_kernel_cache_config_t KernelCacheConfig, bool KernelIsCooperative,
634634
const bool KernelUsesClusterLaunch, const size_t WorkGroupMemorySize,
635635
const RTDeviceBinaryImage *BinImage = nullptr,
636-
void *KernelFuncPtr = nullptr);
636+
void *KernelFuncPtr = nullptr, int KernelNumArgs = 0,
637+
detail::kernel_param_desc_t (*KernelParamDescGetter)(int) = nullptr,
638+
bool KernelHasSpecialCaptures = true);
637639

638640
/// The exec CG command enqueues execution of kernel or explicit memory
639641
/// operation.

‎sycl/test-e2e/Adapters/level_zero/batch_barrier.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -24,7 +24,7 @@ int main(int argc, char *argv[]) {
2424
queue q;
2525

2626
submit_kernel(q); // starts a batch
27-
// CHECK: ---> urEnqueueKernelLaunch
27+
// CHECK: ---> urEnqueueKernelLaunchWithArgsExp
2828
// CHECK-NOT: zeCommandQueueExecuteCommandLists
2929

3030
// Initializing Level Zero driver is required if this test is linked
@@ -42,7 +42,7 @@ int main(int argc, char *argv[]) {
4242
// CHECK-NOT: zeCommandQueueExecuteCommandLists
4343

4444
submit_kernel(q);
45-
// CHECK: ---> urEnqueueKernelLaunch
45+
// CHECK: ---> urEnqueueKernelLaunchWithArgsExp
4646
// CHECK-NOT: zeCommandQueueExecuteCommandLists
4747

4848
// interop should close the batch

‎sycl/test-e2e/Adapters/level_zero/batch_test.cpp‎

Lines changed: 32 additions & 33 deletions
Original file line numberDiff line numberDiff line change
@@ -55,55 +55,54 @@
5555
// variable SYCL_PI_LEVEL_ZEOR+BATCH_SIZE=N.
5656
// This test enqueues 8 kernels and then does a wait. And it does this 3 times.
5757
// Expected output is that for batching =1 you will see zeCommandListClose,
58-
// and zeCommandQueueExecuteCommandLists after every urEnqueueKernelLaunch.
59-
// For batching=3 you will see that after 3rd and 6th enqueues, and then after
60-
// urQueueFinish. For 5, after 5th urEnqueue, and then after urQueueFinish. For
61-
// 4 you will see these after 4th and 8th Enqueue, and for 8, only after the
62-
// 8th enqueue. And lastly for 9, you will see the Close and Execute calls
63-
// only after the urQueueFinish.
64-
// Since the test does this 3 times, this pattern will repeat 2 more times,
65-
// and then the test will print Test Passed 8 times, once for each kernel
66-
// validation check.
58+
// and zeCommandQueueExecuteCommandLists after every
59+
// urEnqueueKernelLaunchWithArgsExp. For batching=3 you will see that after 3rd
60+
// and 6th enqueues, and then after urQueueFinish. For 5, after 5th urEnqueue,
61+
// and then after urQueueFinish. For 4 you will see these after 4th and 8th
62+
// Enqueue, and for 8, only after the 8th enqueue. And lastly for 9, you will
63+
// see the Close and Execute calls only after the urQueueFinish. Since the test
64+
// does this 3 times, this pattern will repeat 2 more times, and then the test
65+
// will print Test Passed 8 times, once for each kernel validation check.
6766
// Pattern starts first set of kernel executions.
68-
// CKALL: ---> urEnqueueKernelLaunch
67+
// CKALL: ---> urEnqueueKernelLaunchWithArgsExp
6968
// CKALL: zeCommandListAppendLaunchKernel(
7069
// CKB1: zeCommandListClose(
7170
// CKB1: zeCommandQueueExecuteCommandLists(
72-
// CKALL: ---> urEnqueueKernelLaunch
71+
// CKALL: ---> urEnqueueKernelLaunchWithArgsExp
7372
// CKALL: zeCommandListAppendLaunchKernel(
7473
// CKB1: zeCommandListClose(
7574
// CKB1: zeCommandQueueExecuteCommandLists(
76-
// CKALL: ---> urEnqueueKernelLaunch
75+
// CKALL: ---> urEnqueueKernelLaunchWithArgsExp
7776
// CKALL: zeCommandListAppendLaunchKernel(
7877
// CKB1: zeCommandListClose(
7978
// CKB1: zeCommandQueueExecuteCommandLists(
8079
// CKB3: zeCommandListClose(
8180
// CKB3: zeCommandQueueExecuteCommandLists(
82-
// CKALL: ---> urEnqueueKernelLaunch
81+
// CKALL: ---> urEnqueueKernelLaunchWithArgsExp
8382
// CKALL: zeCommandListAppendLaunchKernel(
8483
// CKB1: zeCommandListClose(
8584
// CKB1: zeCommandQueueExecuteCommandLists(
8685
// CKB4: zeCommandListClose(
8786
// CKB4: zeCommandQueueExecuteCommandLists(
88-
// CKALL: ---> urEnqueueKernelLaunch
87+
// CKALL: ---> urEnqueueKernelLaunchWithArgsExp
8988
// CKALL: zeCommandListAppendLaunchKernel(
9089
// CKB1: zeCommandListClose(
9190
// CKB1: zeCommandQueueExecuteCommandLists(
9291
// CKB5: zeCommandListClose(
9392
// CKB5: zeCommandQueueExecuteCommandLists(
94-
// CKALL: ---> urEnqueueKernelLaunch
93+
// CKALL: ---> urEnqueueKernelLaunchWithArgsExp
9594
// CKALL: zeCommandListAppendLaunchKernel(
9695
// CKB1: zeCommandListClose(
9796
// CKB1: zeCommandQueueExecuteCommandLists(
9897
// CKB3: zeCommandListClose(
9998
// CKB3: zeCommandQueueExecuteCommandLists(
100-
// CKALL: ---> urEnqueueKernelLaunch
99+
// CKALL: ---> urEnqueueKernelLaunchWithArgsExp
101100
// CKALL: zeCommandListAppendLaunchKernel(
102101
// CKB1: zeCommandListClose(
103102
// CKB1: zeCommandQueueExecuteCommandLists(
104103
// CKB7: zeCommandListClose(
105104
// CKB7: zeCommandQueueExecuteCommandLists(
106-
// CKALL: ---> urEnqueueKernelLaunch
105+
// CKALL: ---> urEnqueueKernelLaunchWithArgsExp
107106
// CKALL: zeCommandListAppendLaunchKernel(
108107
// CKB1: zeCommandListClose(
109108
// CKB1: zeCommandQueueExecuteCommandLists(
@@ -121,45 +120,45 @@
121120
// CKB9: zeCommandListClose(
122121
// CKB9: zeCommandQueueExecuteCommandLists(
123122
// Pattern starts 2nd set of kernel executions
124-
// CKALL: ---> urEnqueueKernelLaunch
123+
// CKALL: ---> urEnqueueKernelLaunchWithArgsExp
125124
// CKALL: zeCommandListAppendLaunchKernel(
126125
// CKB1: zeCommandListClose(
127126
// CKB1: zeCommandQueueExecuteCommandLists(
128-
// CKALL: ---> urEnqueueKernelLaunch
127+
// CKALL: ---> urEnqueueKernelLaunchWithArgsExp
129128
// CKALL: zeCommandListAppendLaunchKernel(
130129
// CKB1: zeCommandListClose(
131130
// CKB1: zeCommandQueueExecuteCommandLists(
132-
// CKALL: ---> urEnqueueKernelLaunch
131+
// CKALL: ---> urEnqueueKernelLaunchWithArgsExp
133132
// CKALL: zeCommandListAppendLaunchKernel(
134133
// CKB1: zeCommandListClose(
135134
// CKB1: zeCommandQueueExecuteCommandLists(
136135
// CKB3: zeCommandListClose(
137136
// CKB3: zeCommandQueueExecuteCommandLists(
138-
// CKALL: ---> urEnqueueKernelLaunch
137+
// CKALL: ---> urEnqueueKernelLaunchWithArgsExp
139138
// CKALL: zeCommandListAppendLaunchKernel(
140139
// CKB1: zeCommandListClose(
141140
// CKB1: zeCommandQueueExecuteCommandLists(
142141
// CKB4: zeCommandListClose(
143142
// CKB4: zeCommandQueueExecuteCommandLists(
144-
// CKALL: ---> urEnqueueKernelLaunch
143+
// CKALL: ---> urEnqueueKernelLaunchWithArgsExp
145144
// CKALL: zeCommandListAppendLaunchKernel(
146145
// CKB1: zeCommandListClose(
147146
// CKB1: zeCommandQueueExecuteCommandLists(
148147
// CKB5: zeCommandListClose(
149148
// CKB5: zeCommandQueueExecuteCommandLists(
150-
// CKALL: ---> urEnqueueKernelLaunch
149+
// CKALL: ---> urEnqueueKernelLaunchWithArgsExp
151150
// CKALL: zeCommandListAppendLaunchKernel(
152151
// CKB1: zeCommandListClose(
153152
// CKB1: zeCommandQueueExecuteCommandLists(
154153
// CKB3: zeCommandListClose(
155154
// CKB3: zeCommandQueueExecuteCommandLists(
156-
// CKALL: ---> urEnqueueKernelLaunch
155+
// CKALL: ---> urEnqueueKernelLaunchWithArgsExp
157156
// CKALL: zeCommandListAppendLaunchKernel(
158157
// CKB1: zeCommandListClose(
159158
// CKB1: zeCommandQueueExecuteCommandLists(
160159
// CKB7: zeCommandListClose(
161160
// CKB7: zeCommandQueueExecuteCommandLists(
162-
// CKALL: ---> urEnqueueKernelLaunch
161+
// CKALL: ---> urEnqueueKernelLaunchWithArgsExp
163162
// CKALL: zeCommandListAppendLaunchKernel(
164163
// CKB1: zeCommandListClose(
165164
// CKB1: zeCommandQueueExecuteCommandLists(
@@ -177,45 +176,45 @@
177176
// CKB9: zeCommandListClose(
178177
// CKB9: zeCommandQueueExecuteCommandLists(
179178
// Pattern starts 3rd set of kernel executions
180-
// CKALL: ---> urEnqueueKernelLaunch
179+
// CKALL: ---> urEnqueueKernelLaunchWithArgsExp
181180
// CKALL: zeCommandListAppendLaunchKernel(
182181
// CKB1: zeCommandListClose(
183182
// CKB1: zeCommandQueueExecuteCommandLists(
184-
// CKALL: ---> urEnqueueKernelLaunch
183+
// CKALL: ---> urEnqueueKernelLaunchWithArgsExp
185184
// CKALL: zeCommandListAppendLaunchKernel(
186185
// CKB1: zeCommandListClose(
187186
// CKB1: zeCommandQueueExecuteCommandLists(
188-
// CKALL: ---> urEnqueueKernelLaunch
187+
// CKALL: ---> urEnqueueKernelLaunchWithArgsExp
189188
// CKALL: zeCommandListAppendLaunchKernel(
190189
// CKB1: zeCommandListClose(
191190
// CKB1: zeCommandQueueExecuteCommandLists(
192191
// CKB3: zeCommandListClose(
193192
// CKB3: zeCommandQueueExecuteCommandLists(
194-
// CKALL: ---> urEnqueueKernelLaunch
193+
// CKALL: ---> urEnqueueKernelLaunchWithArgsExp
195194
// CKALL: zeCommandListAppendLaunchKernel(
196195
// CKB1: zeCommandListClose(
197196
// CKB1: zeCommandQueueExecuteCommandLists(
198197
// CKB4: zeCommandListClose(
199198
// CKB4: zeCommandQueueExecuteCommandLists(
200-
// CKALL: ---> urEnqueueKernelLaunch
199+
// CKALL: ---> urEnqueueKernelLaunchWithArgsExp
201200
// CKALL: zeCommandListAppendLaunchKernel(
202201
// CKB1: zeCommandListClose(
203202
// CKB1: zeCommandQueueExecuteCommandLists(
204203
// CKB5: zeCommandListClose(
205204
// CKB5: zeCommandQueueExecuteCommandLists(
206-
// CKALL: ---> urEnqueueKernelLaunch
205+
// CKALL: ---> urEnqueueKernelLaunchWithArgsExp
207206
// CKALL: zeCommandListAppendLaunchKernel(
208207
// CKB1: zeCommandListClose(
209208
// CKB1: zeCommandQueueExecuteCommandLists(
210209
// CKB3: zeCommandListClose(
211210
// CKB3: zeCommandQueueExecuteCommandLists(
212-
// CKALL: ---> urEnqueueKernelLaunch
211+
// CKALL: ---> urEnqueueKernelLaunchWithArgsExp
213212
// CKALL: zeCommandListAppendLaunchKernel(
214213
// CKB1: zeCommandListClose(
215214
// CKB1: zeCommandQueueExecuteCommandLists(
216215
// CKB7: zeCommandListClose(
217216
// CKB7: zeCommandQueueExecuteCommandLists(
218-
// CKALL: ---> urEnqueueKernelLaunch
217+
// CKALL: ---> urEnqueueKernelLaunchWithArgsExp
219218
// CKALL: zeCommandListAppendLaunchKernel(
220219
// CKB1: zeCommandListClose(
221220
// CKB1: zeCommandQueueExecuteCommandLists(

0 commit comments

Comments
 (0)