1 /*
2 * Copyright (C) 2018-2021 Intel Corporation
3 *
4 * SPDX-License-Identifier: MIT
5 *
6 */
7
8 #include "shared/source/command_stream/command_stream_receiver.h"
9 #include "shared/source/os_interface/os_context.h"
10 #include "shared/test/common/helpers/debug_manager_state_restore.h"
11 #include "shared/test/common/helpers/unit_test_helper.h"
12 #include "shared/test/common/libult/ult_command_stream_receiver.h"
13 #include "shared/test/common/mocks/mock_allocation_properties.h"
14 #include "shared/test/common/test_macros/test.h"
15 #include "shared/test/common/test_macros/test_checks_shared.h"
16
17 #include "opencl/source/event/user_event.h"
18 #include "opencl/source/helpers/cl_memory_properties_helpers.h"
19 #include "opencl/test/unit_test/command_queue/command_enqueue_fixture.h"
20 #include "opencl/test/unit_test/command_queue/command_queue_fixture.h"
21 #include "opencl/test/unit_test/fixtures/cl_device_fixture.h"
22 #include "opencl/test/unit_test/fixtures/image_fixture.h"
23 #include "opencl/test/unit_test/mocks/mock_context.h"
24 #include "opencl/test/unit_test/mocks/mock_kernel.h"
25
26 using namespace NEO;
27
28 struct EnqueueMapImageTest : public ClDeviceFixture,
29 public CommandQueueHwFixture,
30 public ::testing::Test {
31 typedef CommandQueueHwFixture CommandQueueFixture;
32
EnqueueMapImageTestEnqueueMapImageTest33 EnqueueMapImageTest() {
34 }
35
SetUpEnqueueMapImageTest36 void SetUp() override {
37 REQUIRE_IMAGES_OR_SKIP(defaultHwInfo);
38
39 ClDeviceFixture::SetUp();
40 CommandQueueFixture::SetUp(pClDevice, 0);
41 context = new MockContext(pClDevice);
42 image = ImageHelper<ImageUseHostPtr<Image2dDefaults>>::create(context);
43 }
44
TearDownEnqueueMapImageTest45 void TearDown() override {
46 if (IsSkipped()) {
47 return;
48 }
49 delete image;
50 context->release();
51 CommandQueueFixture::TearDown();
52 ClDeviceFixture::TearDown();
53 }
54
55 MockContext *context;
56 cl_int retVal = CL_INVALID_VALUE;
57 Image *image = nullptr;
58 char srcMemory[128];
59 };
60
61 struct EnqueueMapImageParamsTest : public EnqueueMapImageTest,
62 public ::testing::WithParamInterface<uint32_t> {
63 };
64
TEST_F(EnqueueMapImageTest,GivenTiledImageWhenMappingImageThenPointerIsReused)65 TEST_F(EnqueueMapImageTest, GivenTiledImageWhenMappingImageThenPointerIsReused) {
66 auto mapFlags = CL_MAP_READ;
67 const size_t origin[3] = {0, 0, 0};
68 const size_t region[3] = {1, 1, 1};
69
70 auto mapAllocation = image->getMapAllocation(pClDevice->getRootDeviceIndex());
71 EXPECT_NE(nullptr, mapAllocation);
72
73 auto ptr1 = pCmdQ->enqueueMapImage(
74 image, true, mapFlags, origin,
75 region, nullptr, nullptr, 0,
76 nullptr, nullptr, retVal);
77 EXPECT_EQ(CL_SUCCESS, retVal);
78 EXPECT_NE(nullptr, image->getHostPtr());
79 mapAllocation = image->getMapAllocation(pClDevice->getRootDeviceIndex());
80 EXPECT_NE(nullptr, mapAllocation);
81
82 auto ptr2 = pCmdQ->enqueueMapImage(
83 image, true, mapFlags, origin,
84 region, nullptr, nullptr, 0,
85 nullptr, nullptr, retVal);
86 EXPECT_EQ(CL_SUCCESS, retVal);
87
88 EXPECT_EQ(ptr1, ptr2);
89
90 retVal = pCmdQ->enqueueUnmapMemObject(image, ptr1, 0, nullptr, nullptr);
91 EXPECT_EQ(CL_SUCCESS, retVal);
92 }
93
HWTEST_F(EnqueueMapImageTest,givenAllocatedMapPtrAndMapWithDifferentOriginIsCalledThenReturnDifferentPointers)94 HWTEST_F(EnqueueMapImageTest, givenAllocatedMapPtrAndMapWithDifferentOriginIsCalledThenReturnDifferentPointers) {
95 if (!UnitTestHelper<FamilyType>::tiledImagesSupported) {
96 GTEST_SKIP();
97 }
98 std::unique_ptr<Image> img(Image2dHelper<Image2dDefaults>::create(context));
99 auto mapFlags = CL_MAP_READ;
100 const size_t origin1[3] = {0, 0, 0};
101 const size_t origin2[3] = {2, 2, 0};
102 const size_t region[3] = {1, 1, 1};
103
104 auto ptr1 = pCmdQ->enqueueMapImage(img.get(), true, mapFlags, origin1,
105 region, nullptr, nullptr, 0,
106 nullptr, nullptr, retVal);
107 EXPECT_EQ(CL_SUCCESS, retVal);
108
109 auto ptr2 = pCmdQ->enqueueMapImage(img.get(), true, mapFlags, origin2,
110 region, nullptr, nullptr, 0,
111 nullptr, nullptr, retVal);
112 EXPECT_EQ(CL_SUCCESS, retVal);
113
114 EXPECT_NE(ptr1, ptr2);
115 EXPECT_NE(nullptr, img->getAllocatedMapPtr());
116
117 size_t mapOffset = img->getSurfaceFormatInfo().surfaceFormat.ImageElementSizeInBytes * origin2[0] +
118 img->getHostPtrRowPitch() * origin2[1];
119 EXPECT_EQ(ptr2, ptrOffset(ptr1, mapOffset));
120 }
121
122 typedef EnqueueMapImageParamsTest MipMapMapImageParamsTest;
123
TEST_P(MipMapMapImageParamsTest,givenAllocatedMapPtrWhenMapsWithDifferentMipMapsAreCalledThenReturnDifferentPointers)124 TEST_P(MipMapMapImageParamsTest, givenAllocatedMapPtrWhenMapsWithDifferentMipMapsAreCalledThenReturnDifferentPointers) {
125 auto image_type = (cl_mem_object_type)GetParam();
126 cl_int retVal = CL_SUCCESS;
127 cl_image_desc imageDesc = {};
128 imageDesc.image_type = image_type;
129 imageDesc.num_mip_levels = 10;
130 imageDesc.image_width = 4;
131 imageDesc.image_height = 1;
132 imageDesc.image_depth = 1;
133 const size_t origin1[4] = {0, 0, 0, 0};
134 size_t origin2[4] = {0, 0, 0, 0};
135 std::unique_ptr<Image> image;
136 size_t mapOffset = 16u;
137 switch (image_type) {
138 case CL_MEM_OBJECT_IMAGE1D:
139 origin2[1] = 1;
140 image = std::unique_ptr<Image>(ImageHelper<Image1dDefaults>::create(context, &imageDesc));
141 break;
142 case CL_MEM_OBJECT_IMAGE1D_ARRAY:
143 origin2[2] = 1;
144 imageDesc.image_array_size = 2;
145 image = std::unique_ptr<Image>(ImageHelper<Image1dArrayDefaults>::create(context, &imageDesc));
146 break;
147 case CL_MEM_OBJECT_IMAGE2D:
148 origin2[2] = 1;
149 image = std::unique_ptr<Image>(ImageHelper<Image2dDefaults>::create(context, &imageDesc));
150 break;
151 case CL_MEM_OBJECT_IMAGE2D_ARRAY:
152 origin2[3] = 1;
153 imageDesc.image_array_size = 2;
154 image = std::unique_ptr<Image>(ImageHelper<Image2dArrayDefaults>::create(context, &imageDesc));
155 break;
156 case CL_MEM_OBJECT_IMAGE3D:
157 origin2[3] = 1;
158 image = std::unique_ptr<Image>(ImageHelper<Image3dDefaults>::create(context, &imageDesc));
159 break;
160 }
161 EXPECT_NE(nullptr, image.get());
162
163 auto mapFlags = CL_MAP_READ;
164 const size_t region[3] = {1, 1, 1};
165
166 auto ptr1 = pCmdQ->enqueueMapImage(image.get(), true, mapFlags, origin1,
167 region, nullptr, nullptr, 0,
168 nullptr, nullptr, retVal);
169 EXPECT_EQ(CL_SUCCESS, retVal);
170
171 auto ptr2 = pCmdQ->enqueueMapImage(image.get(), true, mapFlags, origin2,
172 region, nullptr, nullptr, 0,
173 nullptr, nullptr, retVal);
174 EXPECT_EQ(CL_SUCCESS, retVal);
175
176 EXPECT_NE(ptr1, ptr2);
177 if (image->mappingOnCpuAllowed() == false) {
178 EXPECT_NE(nullptr, image->getAllocatedMapPtr());
179 }
180
181 EXPECT_EQ(ptr2, ptrOffset(ptr1, mapOffset));
182 }
183
184 INSTANTIATE_TEST_CASE_P(MipMapMapImageParamsTest_givenAllocatedMapPtrAndMapWithDifferentMipMapsIsCalledThenReturnDifferentPointers,
185 MipMapMapImageParamsTest, ::testing::Values(CL_MEM_OBJECT_IMAGE1D, CL_MEM_OBJECT_IMAGE1D_ARRAY, CL_MEM_OBJECT_IMAGE2D, CL_MEM_OBJECT_IMAGE2D_ARRAY, CL_MEM_OBJECT_IMAGE3D));
186
187 template <typename GfxFamily>
188 struct mockedImage : public ImageHw<GfxFamily> {
189 using ImageHw<GfxFamily>::ImageHw;
setAllocatedMapPtrmockedImage190 void setAllocatedMapPtr(void *allocatedMapPtr) override {
191 ownershipTaken = this->hasOwnership();
192 MemObj::setAllocatedMapPtr(allocatedMapPtr);
193 }
194 bool ownershipTaken = false;
195 };
196
HWTEST_F(EnqueueMapImageTest,givenTiledImageWhenMapImageIsCalledThenStorageIsSetWithImageMutexTaken)197 HWTEST_F(EnqueueMapImageTest, givenTiledImageWhenMapImageIsCalledThenStorageIsSetWithImageMutexTaken) {
198 if (!UnitTestHelper<FamilyType>::tiledImagesSupported) {
199 GTEST_SKIP();
200 }
201 auto imageFormat = image->getImageFormat();
202 auto imageDesc = image->getImageDesc();
203 auto graphicsAllocation = image->getGraphicsAllocation(pClDevice->getRootDeviceIndex());
204 auto surfaceFormatInfo = image->getSurfaceFormatInfo();
205
206 mockedImage<FamilyType> mockImage(context,
207 {},
208 0,
209 0,
210 4096u,
211 nullptr,
212 nullptr,
213 imageFormat,
214 imageDesc,
215 false,
216 GraphicsAllocationHelper::toMultiGraphicsAllocation(graphicsAllocation),
217 true,
218 0,
219 0,
220 surfaceFormatInfo,
221 nullptr);
222
223 mockImage.createFunction = image->createFunction;
224
225 auto mapAllocation = mockImage.getMapAllocation(pClDevice->getRootDeviceIndex());
226 EXPECT_EQ(nullptr, mapAllocation);
227 EXPECT_EQ(nullptr, mockImage.getHostPtr());
228
229 auto mapFlags = CL_MAP_READ;
230 const size_t origin[3] = {0, 0, 0};
231 const size_t region[3] = {1, 1, 1};
232
233 auto apiMapPtr = pCmdQ->enqueueMapImage(
234 &mockImage, true, mapFlags, origin,
235 region, nullptr, nullptr, 0,
236 nullptr, nullptr, retVal);
237 EXPECT_TRUE(mockImage.ownershipTaken);
238
239 auto mapPtr = mockImage.getAllocatedMapPtr();
240 EXPECT_EQ(apiMapPtr, mapPtr);
241 mapAllocation = mockImage.getMapAllocation(pClDevice->getRootDeviceIndex());
242 EXPECT_NE(nullptr, mapAllocation);
243 EXPECT_EQ(apiMapPtr, mapAllocation->getUnderlyingBuffer());
244
245 auto osContextId = pCmdQ->getGpgpuCommandStreamReceiver().getOsContext().getContextId();
246 auto expectedTaskCount = pCmdQ->getGpgpuCommandStreamReceiver().peekTaskCount();
247 auto actualMapAllocationTaskCount = mapAllocation->getTaskCount(osContextId);
248 EXPECT_EQ(expectedTaskCount, actualMapAllocationTaskCount);
249
250 pDevice->getMemoryManager()->freeGraphicsMemory(mockImage.getMapAllocation(pClDevice->getRootDeviceIndex()));
251 mockImage.releaseAllocatedMapPtr();
252 }
253
TEST_F(EnqueueMapImageTest,WhenMappingImageThenCpuAndGpuAddressAreEqualWhenZeroCopyIsUsed)254 TEST_F(EnqueueMapImageTest, WhenMappingImageThenCpuAndGpuAddressAreEqualWhenZeroCopyIsUsed) {
255 auto mapFlags = CL_MAP_READ;
256 const size_t origin[3] = {0, 0, 0};
257 const size_t region[3] = {1, 1, 1};
258 size_t imageRowPitch = 0;
259 size_t imageSlicePitch = 0;
260 auto ptr = pCmdQ->enqueueMapImage(
261 image,
262 true,
263 mapFlags,
264 origin,
265 region,
266 &imageRowPitch,
267 &imageSlicePitch,
268 0,
269 nullptr,
270 nullptr,
271 retVal);
272 if (image->isMemObjZeroCopy()) {
273 EXPECT_EQ(image->getCpuAddress(), ptr);
274 } else {
275 EXPECT_NE(image->getCpuAddress(), ptr);
276 }
277 size_t imageRowPitchRef = 0;
278 image->getImageInfo(CL_IMAGE_ROW_PITCH, sizeof(imageRowPitchRef), &imageRowPitchRef, nullptr);
279 EXPECT_EQ(imageRowPitch, imageRowPitchRef);
280
281 size_t imageSlicePitchRef = 0;
282 image->getImageInfo(CL_IMAGE_SLICE_PITCH, sizeof(imageSlicePitchRef), &imageSlicePitchRef, nullptr);
283 EXPECT_EQ(imageSlicePitch, imageSlicePitchRef);
284 }
285
TEST_F(EnqueueMapImageTest,GivenCmdqAndValidArgsWhenMappingImageThenSuccessIsReturned)286 TEST_F(EnqueueMapImageTest, GivenCmdqAndValidArgsWhenMappingImageThenSuccessIsReturned) {
287 auto mapFlags = CL_MAP_READ;
288 const size_t origin[3] = {0, 0, 0};
289 const size_t region[3] = {1, 1, 1};
290 size_t imageRowPitch = 0;
291 size_t imageSlicePitch = 0;
292 auto ptr = pCmdQ->enqueueMapImage(
293 image,
294 true,
295 mapFlags,
296 origin,
297 region,
298 &imageRowPitch,
299 &imageSlicePitch,
300 0,
301 nullptr,
302 nullptr,
303 retVal);
304 EXPECT_NE(nullptr, ptr);
305 EXPECT_EQ(CL_SUCCESS, retVal);
306
307 size_t imageRowPitchRef = 0;
308 image->getImageInfo(CL_IMAGE_ROW_PITCH, sizeof(imageRowPitchRef), &imageRowPitchRef, nullptr);
309 EXPECT_EQ(imageRowPitch, imageRowPitchRef);
310
311 size_t imageSlicePitchRef = 0;
312 image->getImageInfo(CL_IMAGE_SLICE_PITCH, sizeof(imageSlicePitchRef), &imageSlicePitchRef, nullptr);
313 EXPECT_EQ(imageSlicePitch, imageSlicePitchRef);
314 }
315
HWTEST_F(EnqueueMapImageTest,givenNonReadOnlyMapWithOutEventWhenMappedThenSetEventAndIncraseTaskCountFromWriteImage)316 HWTEST_F(EnqueueMapImageTest, givenNonReadOnlyMapWithOutEventWhenMappedThenSetEventAndIncraseTaskCountFromWriteImage) {
317 if (!UnitTestHelper<FamilyType>::tiledImagesSupported) {
318 GTEST_SKIP();
319 }
320 DebugManagerStateRestore dbgRestore;
321 DebugManager.flags.EnableAsyncEventsHandler.set(false);
322 cl_event mapEventReturned = nullptr;
323 cl_event unmapEventReturned = nullptr;
324 uint32_t tagHW = 0;
325 auto mapFlags = CL_MAP_WRITE;
326 const size_t origin[3] = {0, 0, 0};
327 const size_t region[3] = {1, 1, 1};
328 size_t imageRowPitch = 0;
329 size_t imageSlicePitch = 0;
330 size_t GWS = 1;
331
332 MockKernelWithInternals kernel(*pClDevice);
333 *pTagMemory = tagHW;
334 auto &commandStreamReceiver = pCmdQ->getGpgpuCommandStreamReceiver();
335 auto tag_address = commandStreamReceiver.getTagAddress();
336 EXPECT_TRUE(pTagMemory == tag_address);
337
338 struct E2Clb {
339 static void CL_CALLBACK SignalEv2(cl_event e, cl_int status, void *data) {
340 uint32_t *pTagMem = static_cast<uint32_t *>(data);
341 *pTagMem = 4;
342 }
343 };
344
345 uint32_t taskCount = commandStreamReceiver.peekTaskCount();
346 EXPECT_EQ(1u, taskCount);
347
348 // enqueue something that can be finished...
349 retVal = clEnqueueNDRangeKernel(pCmdQ, kernel.mockMultiDeviceKernel, 1, 0, &GWS, nullptr, 0, nullptr, nullptr);
350 EXPECT_EQ(retVal, CL_SUCCESS);
351
352 *pTagMemory = tagHW += 3;
353 auto ptr = pCmdQ->enqueueMapImage(
354 image,
355 false,
356 mapFlags,
357 origin,
358 region,
359 &imageRowPitch,
360 &imageSlicePitch,
361 0,
362 nullptr,
363 &mapEventReturned,
364 retVal);
365
366 EXPECT_NE(nullptr, ptr);
367 EXPECT_EQ(CL_SUCCESS, retVal);
368 auto mapEvent = castToObject<Event>(mapEventReturned);
369 EXPECT_TRUE(CL_COMMAND_MAP_IMAGE == mapEvent->getCommandType());
370
371 taskCount = commandStreamReceiver.peekTaskCount();
372 EXPECT_EQ(3u, taskCount);
373
374 clSetEventCallback(mapEventReturned, CL_COMPLETE, E2Clb::SignalEv2, (void *)pTagMemory);
375
376 retVal = clWaitForEvents(1, &mapEventReturned);
377 EXPECT_EQ(CL_SUCCESS, retVal);
378 EXPECT_EQ(4u, *pTagMemory);
379 taskCount = commandStreamReceiver.peekTaskCount();
380 EXPECT_EQ(3u, taskCount);
381
382 (*pTagMemory)++;
383 retVal = clEnqueueUnmapMemObject(
384 pCmdQ,
385 image,
386 ptr,
387 0,
388 nullptr,
389 &unmapEventReturned);
390 EXPECT_EQ(CL_SUCCESS, retVal);
391 auto unmapEvent = castToObject<Event>(unmapEventReturned);
392 EXPECT_TRUE(CL_COMMAND_UNMAP_MEM_OBJECT == unmapEvent->getCommandType());
393
394 retVal = clWaitForEvents(1, &unmapEventReturned);
395
396 taskCount = commandStreamReceiver.peekTaskCount();
397 EXPECT_EQ(4u, taskCount);
398
399 clReleaseEvent(mapEventReturned);
400 clReleaseEvent(unmapEventReturned);
401 }
402
HWTEST_F(EnqueueMapImageTest,givenReadOnlyMapWithOutEventWhenMappedThenSetEventAndDontIncraseTaskCountFromWriteImage)403 HWTEST_F(EnqueueMapImageTest, givenReadOnlyMapWithOutEventWhenMappedThenSetEventAndDontIncraseTaskCountFromWriteImage) {
404 if (!UnitTestHelper<FamilyType>::tiledImagesSupported) {
405 GTEST_SKIP();
406 }
407 DebugManagerStateRestore dbgRestore;
408 DebugManager.flags.EnableAsyncEventsHandler.set(false);
409 cl_event mapEventReturned = nullptr;
410 cl_event unmapEventReturned = nullptr;
411 auto mapFlags = CL_MAP_READ;
412 const size_t origin[3] = {0, 0, 0};
413 const size_t region[3] = {1, 1, 1};
414 *pTagMemory = 5;
415
416 auto &commandStreamReceiver = pCmdQ->getGpgpuCommandStreamReceiver();
417
418 EXPECT_EQ(1u, commandStreamReceiver.peekTaskCount());
419
420 auto ptr = pCmdQ->enqueueMapImage(image, false, mapFlags, origin, region, nullptr, nullptr, 0,
421 nullptr, &mapEventReturned, retVal);
422
423 EXPECT_NE(nullptr, ptr);
424 EXPECT_EQ(CL_SUCCESS, retVal);
425 EXPECT_EQ(2u, commandStreamReceiver.peekTaskCount());
426
427 auto mapEvent = castToObject<Event>(mapEventReturned);
428 EXPECT_TRUE(CL_COMMAND_MAP_IMAGE == mapEvent->getCommandType());
429
430 retVal = clWaitForEvents(1, &mapEventReturned);
431 EXPECT_EQ(CL_SUCCESS, retVal);
432
433 retVal = clEnqueueUnmapMemObject(pCmdQ, image, ptr, 0, nullptr, &unmapEventReturned);
434 EXPECT_EQ(CL_SUCCESS, retVal);
435
436 EXPECT_EQ(2u, commandStreamReceiver.peekTaskCount());
437
438 auto unmapEvent = castToObject<Event>(unmapEventReturned);
439 EXPECT_TRUE(CL_COMMAND_UNMAP_MEM_OBJECT == unmapEvent->getCommandType());
440
441 retVal = clWaitForEvents(1, &unmapEventReturned);
442 EXPECT_EQ(CL_SUCCESS, retVal);
443
444 clReleaseEvent(mapEventReturned);
445 clReleaseEvent(unmapEventReturned);
446 }
447
HWTEST_F(EnqueueMapImageTest,GivenPtrToReturnEventWhenMappingImageThenEventIsNotNull)448 HWTEST_F(EnqueueMapImageTest, GivenPtrToReturnEventWhenMappingImageThenEventIsNotNull) {
449 if (!UnitTestHelper<FamilyType>::tiledImagesSupported) {
450 GTEST_SKIP();
451 }
452 cl_event eventReturned = nullptr;
453 auto mapFlags = CL_MAP_READ;
454 const size_t origin[3] = {0, 0, 0};
455 const size_t region[3] = {1, 1, 1};
456 size_t imageRowPitch = 0;
457 size_t imageSlicePitch = 0;
458 uint32_t forceTaskCount = 100;
459
460 auto &commandStreamReceiver = pDevice->getUltCommandStreamReceiver<FamilyType>();
461 commandStreamReceiver.taskCount = forceTaskCount;
462
463 auto ptr = clEnqueueMapImage(
464 pCmdQ,
465 image,
466 CL_FALSE,
467 mapFlags,
468 origin,
469 region,
470 &imageRowPitch,
471 &imageSlicePitch,
472 0,
473 nullptr,
474 &eventReturned,
475 &retVal);
476
477 EXPECT_NE(nullptr, ptr);
478 EXPECT_EQ(CL_SUCCESS, retVal);
479 EXPECT_NE(nullptr, eventReturned);
480
481 auto eventObject = castToObject<Event>(eventReturned);
482 EXPECT_EQ(forceTaskCount + 1, eventObject->peekTaskCount());
483 EXPECT_TRUE(eventObject->updateStatusAndCheckCompletion());
484
485 retVal = clEnqueueUnmapMemObject(
486 pCmdQ,
487 image,
488 ptr,
489 0,
490 nullptr,
491 nullptr);
492 EXPECT_EQ(CL_SUCCESS, retVal);
493
494 clReleaseEvent(eventReturned);
495 }
496
HWTEST_F(EnqueueMapImageTest,givenZeroCopyImageWhenItIsMappedAndReturnsEventThenEventHasCorrectProperties)497 HWTEST_F(EnqueueMapImageTest, givenZeroCopyImageWhenItIsMappedAndReturnsEventThenEventHasCorrectProperties) {
498 cl_event eventReturned = nullptr;
499 auto mapFlags = CL_MAP_READ;
500 const size_t origin[3] = {0, 0, 0};
501 const size_t region[3] = {1, 1, 1};
502 size_t imageRowPitch = 0;
503 size_t imageSlicePitch = 0;
504 uint32_t forceTaskCount = 100;
505
506 auto &commandStreamReceiver = pDevice->getUltCommandStreamReceiver<FamilyType>();
507 commandStreamReceiver.taskCount = forceTaskCount;
508
509 std::unique_ptr<Image> zero_copy_image(ImageHelper<ImageWriteOnly<Image1dDefaults>>::create(context));
510
511 ASSERT_TRUE(zero_copy_image->isMemObjZeroCopy());
512 pCmdQ->taskCount = 40u;
513
514 auto ptr = clEnqueueMapImage(
515 pCmdQ,
516 zero_copy_image.get(),
517 CL_FALSE,
518 mapFlags,
519 origin,
520 region,
521 &imageRowPitch,
522 &imageSlicePitch,
523 0,
524 nullptr,
525 &eventReturned,
526 &retVal);
527
528 EXPECT_NE(nullptr, ptr);
529 EXPECT_EQ(CL_SUCCESS, retVal);
530 EXPECT_NE(nullptr, eventReturned);
531
532 EXPECT_EQ(ptr, zero_copy_image->getCpuAddressForMemoryTransfer());
533
534 auto eventObject = castToObject<Event>(eventReturned);
535 EXPECT_EQ(pCmdQ->taskCount, eventObject->peekTaskCount());
536 EXPECT_TRUE(eventObject->updateStatusAndCheckCompletion());
537
538 retVal = clEnqueueUnmapMemObject(
539 pCmdQ,
540 zero_copy_image.get(),
541 ptr,
542 0,
543 nullptr,
544 nullptr);
545 EXPECT_EQ(CL_SUCCESS, retVal);
546
547 clReleaseEvent(eventReturned);
548 }
549
TEST_F(EnqueueMapImageTest,GivenNonZeroCopyImageWhenMappedWithOffsetThenCorrectPointerIsReturned)550 TEST_F(EnqueueMapImageTest, GivenNonZeroCopyImageWhenMappedWithOffsetThenCorrectPointerIsReturned) {
551 auto mapFlags = CL_MAP_WRITE;
552 const size_t origin[3] = {1, 0, 0};
553 const size_t region[3] = {1, 1, 1};
554 size_t imageRowPitch = 0;
555 size_t imageSlicePitch = 0;
556
557 Image *nonZeroCopyImage = ImageHelper<ImageUseHostPtr<Image1dDefaults>>::create(context);
558
559 EXPECT_FALSE(nonZeroCopyImage->isMemObjZeroCopy());
560
561 auto ptr = clEnqueueMapImage(
562 pCmdQ,
563 nonZeroCopyImage,
564 CL_TRUE,
565 mapFlags,
566 origin,
567 region,
568 &imageRowPitch,
569 &imageSlicePitch,
570 0,
571 nullptr,
572 nullptr,
573 &retVal);
574
575 float *HostPtrOffseted = (float *)Image1dDefaults::hostPtr + 1; //
576
577 EXPECT_NE(nullptr, ptr);
578
579 if (!image->isTiledAllocation()) {
580 EXPECT_EQ(HostPtrOffseted, ptr); // Returned pointer should be offseted
581 }
582
583 EXPECT_EQ(CL_SUCCESS, retVal);
584
585 retVal = clEnqueueUnmapMemObject(
586 pCmdQ,
587 nonZeroCopyImage,
588 ptr,
589 0,
590 nullptr,
591 nullptr);
592
593 EXPECT_EQ(CL_SUCCESS, retVal);
594
595 delete nonZeroCopyImage;
596 }
597
HWTEST_F(EnqueueMapImageTest,givenSharingHandlerWhenNonReadOnlyMapAndUnmapOnNonTiledImageIsCalledThenMakeGpuCopy)598 HWTEST_F(EnqueueMapImageTest, givenSharingHandlerWhenNonReadOnlyMapAndUnmapOnNonTiledImageIsCalledThenMakeGpuCopy) {
599 std::unique_ptr<Image> image(ImageHelper<ImageUseHostPtr<Image1dDefaults>>::create(context));
600 ASSERT_NE(nullptr, image);
601 image->setSharingHandler(new SharingHandler());
602 EXPECT_FALSE(image->isTiledAllocation());
603
604 auto &csr = pDevice->getUltCommandStreamReceiver<FamilyType>();
605 csr.taskCount = 1;
606 csr.taskLevel = 1;
607 pCmdQ->taskCount = 1;
608 pCmdQ->taskLevel = 1;
609
610 size_t origin[] = {0, 0, 0};
611 size_t region[] = {1, 1, 1};
612 void *data = clEnqueueMapImage(pCmdQ, image.get(), CL_TRUE, CL_MAP_WRITE, origin, region, nullptr, nullptr, 0, NULL, NULL, &retVal);
613 EXPECT_NE(nullptr, data);
614 EXPECT_EQ(CL_SUCCESS, retVal);
615 EXPECT_EQ(2u, pCmdQ->taskCount);
616 EXPECT_EQ(2u, pCmdQ->taskLevel);
617
618 retVal = clEnqueueUnmapMemObject(pCmdQ, image.get(), data, 0, NULL, NULL);
619 EXPECT_EQ(3u, pCmdQ->taskCount);
620 EXPECT_EQ(3u, pCmdQ->taskLevel);
621 }
622
HWTEST_F(EnqueueMapImageTest,givenSharingHandlerWhenReadOnlyMapAndUnmapOnNonTiledImageIsCalledThenMakeGpuCopy)623 HWTEST_F(EnqueueMapImageTest, givenSharingHandlerWhenReadOnlyMapAndUnmapOnNonTiledImageIsCalledThenMakeGpuCopy) {
624 std::unique_ptr<Image> image(ImageHelper<ImageUseHostPtr<Image1dDefaults>>::create(context));
625 ASSERT_NE(nullptr, image);
626 image->setSharingHandler(new SharingHandler());
627 EXPECT_FALSE(image->isTiledAllocation());
628
629 auto &csr = pDevice->getUltCommandStreamReceiver<FamilyType>();
630 csr.taskCount = 1;
631 csr.taskLevel = 1;
632 pCmdQ->taskCount = 1;
633 pCmdQ->taskLevel = 1;
634
635 size_t origin[] = {0, 0, 0};
636 size_t region[] = {1, 1, 1};
637 void *data = clEnqueueMapImage(pCmdQ, image.get(), CL_TRUE, CL_MAP_READ, origin, region, nullptr, nullptr, 0, NULL, NULL, &retVal);
638 EXPECT_NE(nullptr, data);
639 EXPECT_EQ(CL_SUCCESS, retVal);
640 EXPECT_EQ(2u, pCmdQ->taskCount);
641 EXPECT_EQ(2u, pCmdQ->taskLevel);
642
643 retVal = clEnqueueUnmapMemObject(pCmdQ, image.get(), data, 0, NULL, NULL);
644 EXPECT_EQ(2u, pCmdQ->taskCount);
645 EXPECT_EQ(2u, pCmdQ->taskLevel);
646 }
647
HWTEST_F(EnqueueMapImageTest,givenImageWithouUsetHostPtrFlagWhenMappedOnCpuThenSetAllMapProperties)648 HWTEST_F(EnqueueMapImageTest, givenImageWithouUsetHostPtrFlagWhenMappedOnCpuThenSetAllMapProperties) {
649 std::unique_ptr<Image> image(ImageHelper<Image1dDefaults>::create(context));
650 ASSERT_NE(nullptr, image);
651 EXPECT_TRUE(image->mappingOnCpuAllowed());
652
653 size_t origin[] = {2, 0, 0};
654 size_t region[] = {2, 1, 1};
655 void *mappedPtr = clEnqueueMapImage(pCmdQ, image.get(), CL_TRUE, CL_MAP_READ, origin, region, nullptr, nullptr, 0, NULL, NULL, &retVal);
656 EXPECT_NE(nullptr, mappedPtr);
657
658 MapInfo mappedInfo;
659 auto success = image->findMappedPtr(mappedPtr, mappedInfo);
660 EXPECT_TRUE(success);
661 EXPECT_NE(nullptr, mappedInfo.ptr);
662
663 EXPECT_EQ(origin[0], mappedInfo.offset[0]);
664 EXPECT_EQ(origin[1], mappedInfo.offset[1]);
665 EXPECT_EQ(origin[2], mappedInfo.offset[2]);
666
667 EXPECT_EQ(region[0], mappedInfo.size[0]);
668 EXPECT_EQ(region[1], mappedInfo.size[1]);
669 EXPECT_EQ(region[2], mappedInfo.size[2]);
670
671 auto expectedPtr = ptrOffset(image->getCpuAddressForMapping(), image->calculateOffsetForMapping(mappedInfo.offset));
672
673 EXPECT_EQ(mappedPtr, expectedPtr);
674 }
675
HWTEST_F(EnqueueMapImageTest,givenImageWithUseHostPtrFlagWhenMappedOnCpuThenSetAllMapProperties)676 HWTEST_F(EnqueueMapImageTest, givenImageWithUseHostPtrFlagWhenMappedOnCpuThenSetAllMapProperties) {
677 std::unique_ptr<Image> image(ImageHelper<ImageUseHostPtr<Image1dDefaults>>::create(context));
678 ASSERT_NE(nullptr, image);
679 EXPECT_TRUE(image->mappingOnCpuAllowed());
680
681 size_t origin[] = {2, 0, 0};
682 size_t region[] = {2, 1, 1};
683 void *mappedPtr = clEnqueueMapImage(pCmdQ, image.get(), CL_TRUE, CL_MAP_READ, origin, region, nullptr, nullptr, 0, NULL, NULL, &retVal);
684 EXPECT_NE(nullptr, mappedPtr);
685
686 MapInfo mappedInfo;
687 auto success = image->findMappedPtr(mappedPtr, mappedInfo);
688 EXPECT_TRUE(success);
689 EXPECT_NE(nullptr, mappedInfo.ptr);
690
691 EXPECT_EQ(origin[0], mappedInfo.offset[0]);
692 EXPECT_EQ(origin[1], mappedInfo.offset[1]);
693 EXPECT_EQ(origin[2], mappedInfo.offset[2]);
694
695 EXPECT_EQ(region[0], mappedInfo.size[0]);
696 EXPECT_EQ(region[1], mappedInfo.size[1]);
697 EXPECT_EQ(region[2], mappedInfo.size[2]);
698
699 auto expectedPtr = ptrOffset(image->getCpuAddressForMapping(), image->calculateOffsetForMapping(mappedInfo.offset));
700
701 EXPECT_EQ(mappedPtr, expectedPtr);
702 }
703
TEST_F(EnqueueMapImageTest,givenBlockedCommandQueueWhenBlockingMapWith2DImageIsEnqueuedAndEventAsynchrounouslyCompletedThenEnqueueFinishesWithoutStall)704 TEST_F(EnqueueMapImageTest, givenBlockedCommandQueueWhenBlockingMapWith2DImageIsEnqueuedAndEventAsynchrounouslyCompletedThenEnqueueFinishesWithoutStall) {
705 auto mapFlags = CL_MAP_READ;
706 const size_t origin[3] = {0, 0, 0};
707 const size_t region[3] = {1, 1, 1};
708 size_t imageRowPitch = 0;
709 size_t imageSlicePitch = 0;
710
711 class MockEventWithSetCompleteOnUpdate : public Event {
712 public:
713 MockEventWithSetCompleteOnUpdate(CommandQueue *cmdQueue, cl_command_type cmdType,
714 uint32_t taskLevel, uint32_t taskCount) : Event(cmdQueue, cmdType, taskLevel, taskCount) {
715 }
716 void updateExecutionStatus() override {
717 setStatus(CL_COMPLETE);
718 }
719 };
720 MockEventWithSetCompleteOnUpdate blockingEvent(pCmdQ, CL_COMMAND_NDRANGE_KERNEL, 0, 1);
721
722 cl_event blockingClEvent = &blockingEvent;
723
724 int32_t initialRefCountCmdQ = pCmdQ->getRefInternalCount();
725
726 auto ptr = pCmdQ->enqueueMapImage(
727 image,
728 true,
729 mapFlags,
730 origin,
731 region,
732 &imageRowPitch,
733 &imageSlicePitch,
734 1,
735 &blockingClEvent,
736 nullptr,
737 retVal);
738 EXPECT_NE(nullptr, ptr);
739 EXPECT_EQ(CL_SUCCESS, retVal);
740 EXPECT_EQ(initialRefCountCmdQ, pCmdQ->getRefInternalCount());
741 }
742
TEST_F(EnqueueMapImageTest,givenBlockedCommandQueueWhenBlockingMapWith1DImageIsEnqueuedAndEventAsynchrounouslyCompletedThenEnqueueFinishesWithoutStall)743 TEST_F(EnqueueMapImageTest, givenBlockedCommandQueueWhenBlockingMapWith1DImageIsEnqueuedAndEventAsynchrounouslyCompletedThenEnqueueFinishesWithoutStall) {
744 auto mapFlags = CL_MAP_READ;
745 const size_t origin[3] = {0, 0, 0};
746 const size_t region[3] = {1, 1, 1};
747 size_t imageRowPitch = 0;
748 size_t imageSlicePitch = 0;
749
750 Image *image1D = ImageHelper<ImageUseHostPtr<Image1dDefaults>>::create(context);
751
752 ASSERT_NE(nullptr, image1D);
753 class MockEventWithSetCompleteOnUpdate : public Event {
754 public:
755 MockEventWithSetCompleteOnUpdate(CommandQueue *cmdQueue, cl_command_type cmdType,
756 uint32_t taskLevel, uint32_t taskCount) : Event(cmdQueue, cmdType, taskLevel, taskCount) {
757 }
758 void updateExecutionStatus() override {
759 setStatus(CL_COMPLETE);
760 }
761 };
762 MockEventWithSetCompleteOnUpdate blockingEvent(pCmdQ, CL_COMMAND_NDRANGE_KERNEL, 0, 1);
763
764 cl_event blockingClEvent = &blockingEvent;
765
766 int32_t initialRefCountCmdQ = pCmdQ->getRefInternalCount();
767
768 auto ptr = pCmdQ->enqueueMapImage(
769 image1D,
770 true,
771 mapFlags,
772 origin,
773 region,
774 &imageRowPitch,
775 &imageSlicePitch,
776 1,
777 &blockingClEvent,
778 nullptr,
779 retVal);
780 EXPECT_NE(nullptr, ptr);
781 EXPECT_EQ(CL_SUCCESS, retVal);
782
783 EXPECT_EQ(initialRefCountCmdQ, pCmdQ->getRefInternalCount());
784 delete image1D;
785 }
786
TEST_F(EnqueueMapImageTest,givenBlockedCommandQueueWhenBlockingCpuMapIsCalledThenReturnRowPitchAndSlicePitch)787 TEST_F(EnqueueMapImageTest, givenBlockedCommandQueueWhenBlockingCpuMapIsCalledThenReturnRowPitchAndSlicePitch) {
788 const size_t origin[3] = {0, 0, 0};
789 const size_t region[3] = {1, 1, 1};
790 size_t retImageRowPitch = 0;
791 size_t retImageSlicePitch = 0;
792
793 struct MyMockUserEvent : public UserEvent {
794 MyMockUserEvent() : UserEvent(nullptr) {}
795 void updateExecutionStatus() override {
796 setStatus(CL_COMPLETE);
797 }
798 };
799
800 std::unique_ptr<Image> image(ImageHelper<Image1dArrayDefaults>::create(context));
801 EXPECT_TRUE(image->mappingOnCpuAllowed());
802
803 MyMockUserEvent blockingEvent;
804 cl_event blockingClEvent = &blockingEvent;
805
806 pCmdQ->enqueueMapImage(image.get(), true, CL_MAP_READ, origin, region,
807 &retImageRowPitch, &retImageSlicePitch,
808 1, &blockingClEvent, nullptr, retVal);
809 EXPECT_EQ(CL_SUCCESS, retVal);
810
811 EXPECT_NE(0u, retImageRowPitch);
812 EXPECT_NE(0u, retImageSlicePitch);
813
814 image.reset(ImageHelper<Image1dDefaults>::create(context));
815 pCmdQ->enqueueMapImage(image.get(), true, CL_MAP_READ, origin, region,
816 &retImageRowPitch, &retImageSlicePitch,
817 1, &blockingClEvent, nullptr, retVal);
818 EXPECT_EQ(CL_SUCCESS, retVal);
819
820 EXPECT_NE(0u, retImageRowPitch);
821 EXPECT_EQ(0u, retImageSlicePitch);
822 }
823
TEST_F(EnqueueMapImageTest,givenZeroCopyImageWhenMappedOnCpuThenReturnImageRowAndSlicePitch)824 TEST_F(EnqueueMapImageTest, givenZeroCopyImageWhenMappedOnCpuThenReturnImageRowAndSlicePitch) {
825 const size_t origin[3] = {0, 0, 0};
826 const size_t region[3] = {1, 1, 1};
827 size_t retImageRowPitch = 0;
828 size_t retImageSlicePitch = 0;
829
830 std::unique_ptr<Image> image(ImageHelper<Image1dArrayDefaults>::create(context));
831 EXPECT_TRUE(image->mappingOnCpuAllowed());
832 EXPECT_TRUE(image->isMemObjZeroCopy());
833
834 pCmdQ->enqueueMapImage(image.get(), true, CL_MAP_READ, origin, region,
835 &retImageRowPitch, &retImageSlicePitch,
836 0, nullptr, nullptr, retVal);
837 EXPECT_EQ(CL_SUCCESS, retVal);
838
839 EXPECT_EQ(image->getImageDesc().image_row_pitch, retImageRowPitch);
840 EXPECT_EQ(image->getImageDesc().image_slice_pitch, retImageSlicePitch);
841 }
842
TEST_F(EnqueueMapImageTest,givenNonZeroCopyImageWhenMappedOnCpuThenReturnHostRowAndSlicePitch)843 TEST_F(EnqueueMapImageTest, givenNonZeroCopyImageWhenMappedOnCpuThenReturnHostRowAndSlicePitch) {
844 const size_t origin[3] = {0, 0, 0};
845 const size_t region[3] = {1, 1, 1};
846 size_t retImageRowPitch = 0;
847 size_t retImageSlicePitch = 0;
848
849 std::unique_ptr<Image> image(ImageHelper<ImageUseHostPtr<Image1dArrayDefaults>>::create(context));
850 EXPECT_TRUE(image->mappingOnCpuAllowed());
851 EXPECT_FALSE(image->isMemObjZeroCopy());
852
853 pCmdQ->enqueueMapImage(image.get(), true, CL_MAP_READ, origin, region,
854 &retImageRowPitch, &retImageSlicePitch,
855 0, nullptr, nullptr, retVal);
856 EXPECT_EQ(CL_SUCCESS, retVal);
857
858 EXPECT_EQ(image->getHostPtrRowPitch(), retImageRowPitch);
859 EXPECT_EQ(image->getHostPtrSlicePitch(), retImageSlicePitch);
860 }
861
TEST_F(EnqueueMapImageTest,givenZeroCopyImageWhenMappedOnGpuThenReturnHostRowAndSlicePitch)862 TEST_F(EnqueueMapImageTest, givenZeroCopyImageWhenMappedOnGpuThenReturnHostRowAndSlicePitch) {
863 const size_t origin[3] = {0, 0, 0};
864 const size_t region[3] = {1, 1, 1};
865 size_t retImageRowPitch = 0;
866 size_t retImageSlicePitch = 0;
867
868 std::unique_ptr<Image> image(ImageHelper<Image1dArrayDefaults>::create(context));
869 image->setSharingHandler(new SharingHandler());
870 EXPECT_FALSE(image->mappingOnCpuAllowed());
871 EXPECT_TRUE(image->isMemObjZeroCopy());
872
873 pCmdQ->enqueueMapImage(image.get(), true, CL_MAP_READ, origin, region,
874 &retImageRowPitch, &retImageSlicePitch,
875 0, nullptr, nullptr, retVal);
876 EXPECT_EQ(CL_SUCCESS, retVal);
877
878 EXPECT_EQ(image->getHostPtrRowPitch(), retImageRowPitch);
879 EXPECT_EQ(image->getHostPtrSlicePitch(), retImageSlicePitch);
880 }
881
TEST_F(EnqueueMapImageTest,givenNonZeroCopyImageWhenMappedOnGpuThenReturnHostRowAndSlicePitch)882 TEST_F(EnqueueMapImageTest, givenNonZeroCopyImageWhenMappedOnGpuThenReturnHostRowAndSlicePitch) {
883 const size_t origin[3] = {0, 0, 0};
884 const size_t region[3] = {1, 1, 1};
885 size_t retImageRowPitch = 0;
886 size_t retImageSlicePitch = 0;
887
888 std::unique_ptr<Image> image(ImageHelper<ImageUseHostPtr<Image1dArrayDefaults>>::create(context));
889 image->setSharingHandler(new SharingHandler());
890 EXPECT_FALSE(image->mappingOnCpuAllowed());
891 EXPECT_FALSE(image->isMemObjZeroCopy());
892
893 pCmdQ->enqueueMapImage(image.get(), true, CL_MAP_READ, origin, region,
894 &retImageRowPitch, &retImageSlicePitch,
895 0, nullptr, nullptr, retVal);
896 EXPECT_EQ(CL_SUCCESS, retVal);
897
898 EXPECT_EQ(image->getHostPtrRowPitch(), retImageRowPitch);
899 EXPECT_EQ(image->getHostPtrSlicePitch(), retImageSlicePitch);
900 }
901
TEST_F(EnqueueMapImageTest,givenMipMapImageWhenMappedThenReturnHostRowAndSlicePitch)902 TEST_F(EnqueueMapImageTest, givenMipMapImageWhenMappedThenReturnHostRowAndSlicePitch) {
903 const size_t origin[4] = {0, 0, 0, 1};
904 const size_t region[3] = {1, 1, 1};
905 size_t retImageRowPitch = 0;
906 size_t retImageSlicePitch = 0;
907
908 cl_image_desc imageDesc = {};
909 imageDesc.image_type = CL_MEM_OBJECT_IMAGE3D;
910 imageDesc.num_mip_levels = 10;
911 imageDesc.image_width = 4;
912 imageDesc.image_height = 4;
913 imageDesc.image_depth = 4;
914
915 std::unique_ptr<Image> image(ImageHelper<Image3dDefaults>::create(context, &imageDesc));
916 image->setSharingHandler(new SharingHandler());
917 EXPECT_FALSE(image->mappingOnCpuAllowed());
918
919 pCmdQ->enqueueMapImage(image.get(), true, CL_MAP_READ, origin, region,
920 &retImageRowPitch, &retImageSlicePitch,
921 0, nullptr, nullptr, retVal);
922 EXPECT_EQ(CL_SUCCESS, retVal);
923
924 EXPECT_EQ(image->getHostPtrRowPitch(), retImageRowPitch);
925 EXPECT_EQ(image->getHostPtrSlicePitch(), retImageSlicePitch);
926 }
927
TEST_F(EnqueueMapImageTest,givenImage1DArrayWhenEnqueueMapImageIsCalledThenReturnRowAndSlicePitchAreEqual)928 TEST_F(EnqueueMapImageTest, givenImage1DArrayWhenEnqueueMapImageIsCalledThenReturnRowAndSlicePitchAreEqual) {
929 class MockImage : public Image {
930 public:
931 MockImage(Context *context, cl_mem_flags flags, GraphicsAllocation *allocation, const ClSurfaceFormatInfo &surfaceFormat,
932 const cl_image_format &imageFormat, const cl_image_desc &imageDesc)
933 : Image(context, ClMemoryPropertiesHelper::createMemoryProperties(flags, 0, 0, &context->getDevice(0)->getDevice()), flags, 0,
934 0, nullptr, nullptr,
935 imageFormat, imageDesc,
936 true,
937 GraphicsAllocationHelper::toMultiGraphicsAllocation(allocation),
938 false, 0, 0,
939 surfaceFormat, nullptr) {
940 }
941
942 void setImageArg(void *memory, bool isMediaBlockImage, uint32_t mipLevel, uint32_t rootDeviceIndex, bool useGlobalAtomics) override {}
943 void setMediaImageArg(void *memory, uint32_t rootDeviceIndex) override {}
944 void setMediaSurfaceRotation(void *memory) override {}
945 void setSurfaceMemoryObjectControlState(void *memory, uint32_t value) override {}
946 void transformImage2dArrayTo3d(void *memory) override {}
947 void transformImage3dTo2dArray(void *memory) override {}
948 };
949
950 const size_t origin[3] = {0, 0, 0};
951 const size_t region[3] = {1, 1, 1};
952 size_t retImageRowPitch = 0;
953 size_t retImageSlicePitch = 0;
954
955 cl_mem_flags flags = CL_MEM_READ_ONLY;
956
957 cl_image_desc imageDesc = {};
958 imageDesc.image_type = CL_MEM_OBJECT_IMAGE1D_ARRAY;
959 imageDesc.image_width = 329;
960 imageDesc.image_array_size = 48;
961 imageDesc.image_row_pitch = 2688;
962 imageDesc.image_slice_pitch = 10752;
963 imageDesc.num_mip_levels = 0;
964
965 size_t imgSize = imageDesc.image_slice_pitch * imageDesc.image_array_size;
966
967 cl_image_format imageFormat = {};
968 imageFormat.image_channel_order = CL_RGBA;
969 imageFormat.image_channel_data_type = CL_UNSIGNED_INT16;
970
971 const ClSurfaceFormatInfo *surfaceFormat = Image::getSurfaceFormatFromTable(flags, &imageFormat, context->getDevice(0)->getHardwareInfo().capabilityTable.supportsOcl21Features);
972 auto allocation = context->getMemoryManager()->allocateGraphicsMemoryWithProperties(MockAllocationProperties{context->getDevice(0)->getRootDeviceIndex(), imgSize});
973 ASSERT_NE(allocation, nullptr);
974
975 MockImage image(context, flags, allocation, *surfaceFormat, imageFormat, imageDesc);
976
977 EXPECT_TRUE(image.mappingOnCpuAllowed());
978 EXPECT_TRUE(image.isMemObjZeroCopy());
979
980 pCmdQ->enqueueMapImage(&image, true, CL_MAP_READ, origin, region,
981 &retImageRowPitch, &retImageSlicePitch,
982 0, nullptr, nullptr, retVal);
983 EXPECT_EQ(CL_SUCCESS, retVal);
984
985 EXPECT_EQ(retImageRowPitch, retImageSlicePitch);
986 }
987
988 struct EnqueueMapImageTypeTest : public CommandEnqueueFixture,
989 public ::testing::Test {
990
991 typedef CommandQueueHwFixture CommandQueueFixture;
992 using CommandQueueHwFixture::pCmdQ;
993
EnqueueMapImageTypeTestEnqueueMapImageTypeTest994 EnqueueMapImageTypeTest(void) {
995 }
996
SetUpEnqueueMapImageTypeTest997 void SetUp() override {
998 CommandEnqueueFixture::SetUp();
999 image = ImageHelper<ImageUseHostPtr<Image2dDefaults>>::create(&context);
1000 }
1001
TearDownEnqueueMapImageTypeTest1002 void TearDown() override {
1003 delete image;
1004 CommandEnqueueFixture::TearDown();
1005 }
1006
1007 protected:
1008 template <typename FamilyType>
enqueueMapImageEnqueueMapImageTypeTest1009 void enqueueMapImage(cl_bool blocking = CL_TRUE) {
1010 typedef ImageUseHostPtr<Image2dDefaults> Traits;
1011
1012 size_t imageRowPitch;
1013 size_t imageSlicePitch;
1014 size_t origin[3] = {0, 0, 0};
1015 size_t region[3] = {Traits::imageDesc.image_width, Traits::imageDesc.image_height, Traits::imageDesc.image_depth};
1016
1017 cl_int retVal = 0;
1018 auto mappedPtr = pCmdQ->enqueueMapImage(
1019 image,
1020 blocking,
1021 Traits::flags,
1022 origin,
1023 region,
1024 &imageRowPitch,
1025 &imageSlicePitch,
1026 0,
1027 nullptr,
1028 nullptr,
1029 retVal);
1030 EXPECT_NE(nullptr, mappedPtr);
1031 EXPECT_EQ(CL_SUCCESS, retVal);
1032
1033 parseCommands<FamilyType>(*pCmdQ);
1034 }
1035 MockContext context;
1036 Image *image = nullptr;
1037 };
1038
HWCMDTEST_F(IGFX_GEN8_CORE,EnqueueMapImageTypeTest,GiveRequirementForPipeControlWorkaroundWhenMappingImageThenAdditionalPipeControlIsProgrammed)1039 HWCMDTEST_F(IGFX_GEN8_CORE, EnqueueMapImageTypeTest, GiveRequirementForPipeControlWorkaroundWhenMappingImageThenAdditionalPipeControlIsProgrammed) {
1040 typedef typename FamilyType::PIPE_CONTROL PIPE_CONTROL;
1041
1042 // Set taskCount to 1 to call finish on map operation
1043 pCmdQ->taskCount = 1;
1044 bool blocking = true;
1045 enqueueMapImage<FamilyType>(blocking);
1046
1047 auto itorWalker = find<typename FamilyType::GPGPU_WALKER *>(cmdList.begin(), cmdList.end());
1048
1049 auto itorCmd = find<PIPE_CONTROL *>(itorWalker, cmdList.end());
1050 auto *cmd = (PIPE_CONTROL *)*itorCmd;
1051 EXPECT_NE(cmdList.end(), itorCmd);
1052
1053 if (UnitTestHelper<FamilyType>::isPipeControlWArequired(pDevice->getHardwareInfo())) {
1054 // SKL: two PIPE_CONTROLs following GPGPU_WALKER: first has DcFlush and second has Write HwTag
1055 EXPECT_FALSE(cmd->getDcFlushEnable());
1056 // Move to next PPC
1057 auto itorCmdP = ++((GenCmdList::iterator)itorCmd);
1058 EXPECT_NE(cmdList.end(), itorCmdP);
1059 auto itorCmd2 = find<PIPE_CONTROL *>(itorCmdP, cmdList.end());
1060 cmd = (PIPE_CONTROL *)*itorCmd2;
1061 EXPECT_TRUE(cmd->getDcFlushEnable());
1062 } else {
1063 // single PIPE_CONTROL following GPGPU_WALKER has DcFlush and Write HwTag
1064 EXPECT_TRUE(cmd->getDcFlushEnable());
1065 }
1066 }
1067