1 //
2 // Copyright (c) 2017 The Khronos Group Inc.
3 //
4 // Licensed under the Apache License, Version 2.0 (the "License");
5 // you may not use this file except in compliance with the License.
6 // You may obtain a copy of the License at
7 //
8 // http://www.apache.org/licenses/LICENSE-2.0
9 //
10 // Unless required by applicable law or agreed to in writing, software
11 // distributed under the License is distributed on an "AS IS" BASIS,
12 // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13 // See the License for the specific language governing permissions and
14 // limitations under the License.
15 //
16 #define _CRT_SECURE_NO_WARNINGS
17 #include "harness.h"
18 #include <vector>
19
20 Texture2DSize texture2DSizes[] =
21 {
22 {
23 4, // Width
24 4, // Height
25 1, // MipLevels
26 1, // ArraySize
27 1, // SubResourceCount
28 { // SubResources
29 {0, 0}, // MipLevel, ArraySlice
30 {0, 0}, // MipLevel, ArraySlice
31 {0, 0}, // MipLevel, ArraySlice
32 {0, 0}, // MipLevel, ArraySlice
33 },
34 0, // MiscFlags
35 },
36 {
37 15, // Width
38 37, // Height
39 2, // MipLevels
40 1, // ArraySize
41 2, // SubResourceCount
42 { // SubResources
43 {0, 0}, // MipLevel, ArraySlice
44 {1, 0}, // MipLevel, ArraySlice
45 {0, 0}, // MipLevel, ArraySlice
46 {0, 0}, // MipLevel, ArraySlice
47 },
48 0, // MiscFlags
49 },
50 {
51 65, // Width
52 17, // Height
53 1, // MipLevels
54 1, // ArraySize
55 1, // SubResourceCount
56 { // SubResources
57 {0, 0}, // MipLevel, ArraySlice
58 {0, 0}, // MipLevel, ArraySlice
59 {0, 0}, // MipLevel, ArraySlice
60 {0, 0}, // MipLevel, ArraySlice
61 },
62 D3D11_RESOURCE_MISC_SHARED, // MiscFlags
63 },
64
65 {
66 127, // Width
67 125, // Height
68 4, // MipLevels
69 1, // ArraySize
70 4, // SubResourceCount
71 { // SubResources
72 {3, 0}, // MipLevel, ArraySlice
73 {2, 0}, // MipLevel, ArraySlice
74 {1, 0}, // MipLevel, ArraySlice
75 {0, 0}, // MipLevel, ArraySlice
76 },
77 0, // MiscFlags
78 },
79 {
80 128, // Width
81 128, // Height
82 4, // MipLevels
83 6, // ArraySize
84 4, // SubResourceCount
85 { // SubResources
86 {0, 1}, // MipLevel, ArraySlice
87 {1, 0}, // MipLevel, ArraySlice
88 {0, 2}, // MipLevel, ArraySlice
89 {3, 5}, // MipLevel, ArraySlice
90 },
91 0, // MiscFlags
92 },
93 {
94 256, // Width
95 256, // Height
96 0, // MipLevels
97 256, // ArraySize
98 4, // SubResourceCount
99 { // SubResources
100 {0, 0}, // MipLevel, ArraySlice
101 {1, 255}, // MipLevel, ArraySlice
102 {2, 127}, // MipLevel, ArraySlice
103 {3, 128}, // MipLevel, ArraySlice
104 },
105 0, // MiscFlags
106 },
107 {
108 258, // Width
109 511, // Height
110 1, // MipLevels
111 1, // ArraySize
112 1, // SubResourceCount
113 { // SubResources
114 {0, 0}, // MipLevel, ArraySlice
115 {0, 0}, // MipLevel, ArraySlice
116 {0, 0}, // MipLevel, ArraySlice
117 {0, 0}, // MipLevel, ArraySlice
118 },
119 0, // MiscFlags
120 },
121 {
122 767, // Width
123 1025, // Height
124 4, // MipLevels
125 1, // ArraySize
126 1, // SubResourceCount
127 { // SubResources
128 {0, 0}, // MipLevel, ArraySlice
129 {0, 0}, // MipLevel, ArraySlice
130 {0, 0}, // MipLevel, ArraySlice
131 {0, 0}, // MipLevel, ArraySlice
132 },
133 0, // MiscFlags
134 },
135 {
136 2048, // Width
137 2048, // Height
138 1, // MipLevels
139 1, // ArraySize
140 1, // SubResourceCount
141 { // SubResources
142 {0, 0}, // MipLevel, ArraySlice
143 {0, 0}, // MipLevel, ArraySlice
144 {0, 0}, // MipLevel, ArraySlice
145 {0, 0}, // MipLevel, ArraySlice
146 },
147 0, // MiscFlags
148 },
149 };
150 UINT texture2DSizeCount = sizeof(texture2DSizes)/sizeof(texture2DSizes[0]);
151
152 const char *
153 texture2DPatterns[2][2] =
154 {
155 {"aAbBcCdDeEfFgGhHiIjJ", "AaBbCcDdEeFfGgHhIiJj"},
156 {"zZyYxXwWvVuUtTsSrRqQ", "ZzYyXxWwVvUuTtSsRrQq"},
157 };
158
SubTestTexture2D(cl_context context,cl_command_queue command_queue,cl_kernel kernel,ID3D11Device * pDevice,ID3D11DeviceContext * pDC,const TextureFormat * format,const Texture2DSize * size)159 void SubTestTexture2D(
160 cl_context context,
161 cl_command_queue command_queue,
162 cl_kernel kernel,
163 ID3D11Device* pDevice,
164 ID3D11DeviceContext* pDC,
165 const TextureFormat* format,
166 const Texture2DSize* size)
167 {
168 ID3D11Texture2D* pTexture = NULL;
169 HRESULT hr = S_OK;
170 cl_image_format clFormat;
171 cl_int result = CL_SUCCESS;
172
173 HarnessD3D11_TestBegin("2D Texture: Format=%s, Width=%d, Height=%d, MipLevels=%d, ArraySize=%d",
174 format->name_format,
175 size->Width,
176 size->Height,
177 size->MipLevels,
178 size->ArraySize);
179
180 struct
181 {
182 cl_mem mem;
183 UINT subResource;
184 UINT width;
185 UINT height;
186 }
187 subResourceInfo[4];
188
189 cl_event events[4] = {NULL, NULL, NULL, NULL};
190
191 // create the D3D11 resources
192 {
193 D3D11_TEXTURE2D_DESC desc;
194 memset(&desc, 0, sizeof(desc) );
195 desc.Width = size->Width;
196 desc.Height = size->Height;
197 desc.MipLevels = size->MipLevels;
198 desc.ArraySize = size->ArraySize;
199 desc.Format = format->format;
200 desc.SampleDesc.Count = 1;
201 desc.SampleDesc.Quality = 0;
202 desc.Usage = D3D11_USAGE_DEFAULT;
203 desc.BindFlags = D3D11_BIND_SHADER_RESOURCE | D3D11_BIND_RENDER_TARGET;
204 desc.CPUAccessFlags = 0;
205 desc.MiscFlags = 0;
206
207 hr = pDevice->CreateTexture2D(&desc, NULL, &pTexture);
208 TestRequire(SUCCEEDED(hr), "ID3D11Device::CreateTexture2D failed (non-OpenCL D3D error, but test is invalid).");
209 }
210
211 // initialize some useful variables
212 for (UINT i = 0; i < size->SubResourceCount; ++i)
213 {
214 // compute the expected values for the subresource
215 subResourceInfo[i].subResource = D3D11CalcSubresource(
216 size->subResources[i].MipLevel,
217 size->subResources[i].ArraySlice,
218 size->MipLevels);
219 subResourceInfo[i].width = size->Width;
220 subResourceInfo[i].height = size->Height;
221 for (UINT j = 0; j < size->subResources[i].MipLevel; ++j) {
222 subResourceInfo[i].width /= 2;
223 subResourceInfo[i].height /= 2;
224 }
225 subResourceInfo[i].mem = NULL;
226 }
227
228 // copy a pattern into the corners of the image, coordinates
229 // (0,0), (w,0-1), (0,h-1), (w-1,h-1)
230 for (UINT i = 0; i < size->SubResourceCount; ++i)
231 for (UINT x = 0; x < 2; ++x)
232 for (UINT y = 0; y < 2; ++y)
233 {
234 // create the staging buffer
235 ID3D11Texture2D* pStagingBuffer = NULL;
236 {
237 D3D11_TEXTURE2D_DESC desc = {0};
238 desc.Width = 1;
239 desc.Height = 1;
240 desc.MipLevels = 1;
241 desc.ArraySize = 1;
242 desc.Format = format->format;
243 desc.SampleDesc.Count = 1;
244 desc.SampleDesc.Quality = 0;
245 desc.Usage = D3D11_USAGE_STAGING;
246 desc.BindFlags = 0;
247 desc.CPUAccessFlags = D3D11_CPU_ACCESS_READ | D3D11_CPU_ACCESS_WRITE;
248 desc.MiscFlags = 0;
249 hr = pDevice->CreateTexture2D(&desc, NULL, &pStagingBuffer);
250 TestRequire(SUCCEEDED(hr), "ID3D11Device::CreateTexture2D failed (non-OpenCL D3D error, but test is invalid).");
251 }
252
253 // write the data to the staging buffer
254 {
255 D3D11_MAPPED_SUBRESOURCE mappedTexture;
256 hr = pDC->Map(
257 pStagingBuffer,
258 0,
259 D3D11_MAP_READ_WRITE,
260 0,
261 &mappedTexture);
262 memcpy(mappedTexture.pData, texture2DPatterns[x][y], format->bytesPerPixel);
263 pDC->Unmap(pStagingBuffer, 0);
264 }
265
266 // copy the data to to the texture
267 {
268 D3D11_BOX box = {0};
269 box.front = 0; box.back = 1;
270 box.top = 0; box.bottom = 1;
271 box.left = 0; box.right = 1;
272 pDC->CopySubresourceRegion(
273 pTexture,
274 subResourceInfo[i].subResource,
275 x ? subResourceInfo[i].width - 1 : 0,
276 y ? subResourceInfo[i].height - 1 : 0,
277 0,
278 pStagingBuffer,
279 0,
280 &box);
281 }
282
283 pStagingBuffer->Release();
284 }
285
286 // create the cl_mem objects for the resources and verify its sanity
287 for (UINT i = 0; i < size->SubResourceCount; ++i)
288 {
289 // create a cl_mem for the resource
290 subResourceInfo[i].mem = clCreateFromD3D11Texture2DKHR(
291 context,
292 0,
293 pTexture,
294 subResourceInfo[i].subResource,
295 &result);
296 if (CL_IMAGE_FORMAT_NOT_SUPPORTED == result)
297 {
298 goto Cleanup;
299 }
300 TestRequire(result == CL_SUCCESS, "clCreateFromD3D11Texture2DKHR failed");
301
302 // query resource pointer and verify
303 ID3D11Resource* clResource = NULL;
304 result = clGetMemObjectInfo(
305 subResourceInfo[i].mem,
306 CL_MEM_D3D11_RESOURCE_KHR,
307 sizeof(clResource),
308 &clResource,
309 NULL);
310 TestRequire(result == CL_SUCCESS, "clGetMemObjectInfo for CL_MEM_D3D11_RESOURCE_KHR failed.");
311 TestRequire(clResource == pTexture, "clGetMemObjectInfo for CL_MEM_D3D11_RESOURCE_KHR returned incorrect value.");
312
313 // query subresource and verify
314 UINT clSubResource;
315 result = clGetImageInfo(
316 subResourceInfo[i].mem,
317 CL_IMAGE_D3D11_SUBRESOURCE_KHR,
318 sizeof(clSubResource),
319 &clSubResource,
320 NULL);
321 TestRequire(result == CL_SUCCESS, "clGetImageInfo for CL_IMAGE_D3D11_SUBRESOURCE_KHR failed");
322 TestRequire(clSubResource == subResourceInfo[i].subResource, "clGetImageInfo for CL_IMAGE_D3D11_SUBRESOURCE_KHR returned incorrect value.");
323
324 // query format and verify
325 result = clGetImageInfo(
326 subResourceInfo[i].mem,
327 CL_IMAGE_FORMAT,
328 sizeof(clFormat),
329 &clFormat,
330 NULL);
331 TestRequire(result == CL_SUCCESS, "clGetImageInfo for CL_IMAGE_FORMAT failed");
332 TestRequire(clFormat.image_channel_order == format->channel_order, "clGetImageInfo for CL_IMAGE_FORMAT returned incorrect channel order.");
333 TestRequire(clFormat.image_channel_data_type == format->channel_type, "clGetImageInfo for CL_IMAGE_FORMAT returned incorrect channel data type.");
334
335 // query width
336 size_t width;
337 result = clGetImageInfo(
338 subResourceInfo[i].mem,
339 CL_IMAGE_WIDTH,
340 sizeof(width),
341 &width,
342 NULL);
343 TestRequire(result == CL_SUCCESS, "clGetImageInfo for CL_IMAGE_WIDTH failed");
344 TestRequire(width == subResourceInfo[i].width, "clGetImageInfo for CL_IMAGE_HEIGHT returned incorrect value.");
345
346 // query height
347 size_t height;
348 result = clGetImageInfo(
349 subResourceInfo[i].mem,
350 CL_IMAGE_HEIGHT,
351 sizeof(height),
352 &height,
353 NULL);
354 TestRequire(result == CL_SUCCESS, "clGetImageInfo for CL_IMAGE_HEIGHT failed");
355 TestRequire(height == subResourceInfo[i].height, "clGetImageInfo for CL_IMAGE_HEIGHT returned incorrect value.");
356
357 }
358
359 // acquire the resources for OpenCL
360 for (UINT i = 0; i < 2; ++i)
361 {
362 cl_uint memCount = 0;
363 cl_mem memToAcquire[MAX_REGISTERED_SUBRESOURCES];
364
365 // cut the registered sub-resources into two sets and send the acquire calls for them separately
366 if (i == 0)
367 {
368 for(UINT j = 0; j < size->SubResourceCount/2; ++j)
369 {
370 memToAcquire[memCount++] = subResourceInfo[j].mem;
371 }
372 }
373 else
374 {
375 for(UINT j = size->SubResourceCount/2; j < size->SubResourceCount; ++j)
376 {
377 memToAcquire[memCount++] = subResourceInfo[j].mem;
378 }
379 }
380 if (!memCount) continue;
381
382 // do the acquire
383 result = clEnqueueAcquireD3D11ObjectsKHR(
384 command_queue,
385 memCount,
386 memToAcquire,
387 0,
388 NULL,
389 &events[0+i]);
390 TestRequire(result == CL_SUCCESS, "clEnqueueAcquireD3D11ObjectsKHR failed.");
391 TestRequire(events[0+i], "clEnqueueAcquireD3D11ObjectsKHR did not return an event.");
392
393 // make sure the event type is correct
394 cl_uint eventType = 0;
395 result = clGetEventInfo(
396 events[0+i],
397 CL_EVENT_COMMAND_TYPE,
398 sizeof(eventType),
399 &eventType,
400 NULL);
401 TestRequire(result == CL_SUCCESS, "clGetEventInfo for event created by clEnqueueAcquireD3D11ObjectsKHR failed.");
402 TestRequire(eventType == CL_COMMAND_ACQUIRE_D3D11_OBJECTS_KHR, "clGetEventInfo for CL_EVENT_COMMAND_TYPE was not CL_COMMAND_ACQUIRE_D3D11_OBJECTS_KHR.");
403 }
404
405 // download the data using OpenCL & compare with the expected results
406 for (UINT i = 0; i < size->SubResourceCount; ++i)
407 {
408 size_t origin[3] = {0,0,0};
409 size_t region[3] = {subResourceInfo[i].width, subResourceInfo[i].height, 1};
410 cl_mem tempImage;
411 cl_image_desc image_desc = { 0 };
412 image_desc.image_depth = 1;
413 image_desc.image_height = subResourceInfo[i].height;
414 image_desc.image_width = subResourceInfo[i].width;
415 image_desc.image_type = CL_MEM_OBJECT_IMAGE2D;
416
417 tempImage = clCreateImage(context, 0, &clFormat, &image_desc, NULL, &result);
418 TestRequire(result == CL_SUCCESS, "clCreateImage failed");
419
420 result = clEnqueueCopyImage(command_queue, subResourceInfo[i].mem, tempImage,
421 origin, origin, region, 0, NULL, NULL);
422 TestRequire(result == CL_SUCCESS, "clEnqueueCopyImage failed");
423
424 // copy (0,0) to (1,1) and (w-1,h-1) to (w-2,h-2) using a kernel
425 {
426 result = clSetKernelArg(
427 kernel,
428 0,
429 sizeof(cl_mem),
430 (void *)&tempImage);
431 result = clSetKernelArg(
432 kernel,
433 1,
434 sizeof(cl_mem),
435 (void *)&subResourceInfo[i].mem);
436
437 TestRequire(CL_SUCCESS == result, "clSetKernelArg failed");
438
439 size_t localWorkSize[] = {1};
440 size_t globalWorkSize[] = {1};
441 result = clEnqueueNDRangeKernel(
442 command_queue,
443 kernel,
444 1,
445 NULL,
446 globalWorkSize,
447 localWorkSize,
448 0,
449 NULL,
450 NULL);
451 TestRequire(CL_SUCCESS == result, "clEnqueueNDRangeKernel failed");
452 }
453 // copy (w-1,0) to (w-2,1) and (0,h) to (1,h-2) using a memcpy
454 for (UINT x = 0; x < 2; ++x)
455 for (UINT y = 0; y < 2; ++y)
456 {
457 if (x == y)
458 {
459 continue;
460 }
461
462 size_t src[3] =
463 {
464 x ? subResourceInfo[i].width - 1 : 0,
465 y ? subResourceInfo[i].height - 1 : 0,
466 0,
467 };
468 size_t dst[3] =
469 {
470 x ? subResourceInfo[i].width - 2 : 1,
471 y ? subResourceInfo[i].height - 2 : 1,
472 0,
473 };
474 size_t region[3] =
475 {
476 1,
477 1,
478 1,
479 };
480 result = clEnqueueCopyImage(
481 command_queue,
482 subResourceInfo[i].mem,
483 subResourceInfo[i].mem,
484 src,
485 dst,
486 region,
487 0,
488 NULL,
489 NULL);
490 TestRequire(result == CL_SUCCESS, "clEnqueueCopyImage failed.");
491 }
492 clReleaseMemObject(tempImage);
493 }
494
495 // release the resource from OpenCL
496 for (UINT i = 0; i < 2; ++i)
497 {
498 cl_uint memCount = 0;
499 cl_mem memToAcquire[MAX_REGISTERED_SUBRESOURCES];
500
501 // cut the registered sub-resources into two sets and send the release calls for them separately
502 if (i == 0)
503 {
504 for(UINT j = size->SubResourceCount/4; j < size->SubResourceCount; ++j)
505 {
506 memToAcquire[memCount++] = subResourceInfo[j].mem;
507 }
508 }
509 else
510 {
511 for(UINT j = 0; j < size->SubResourceCount/4; ++j)
512 {
513 memToAcquire[memCount++] = subResourceInfo[j].mem;
514 }
515 }
516 if (!memCount) continue;
517
518 // do the release
519 result = clEnqueueReleaseD3D11ObjectsKHR(
520 command_queue,
521 memCount,
522 memToAcquire,
523 0,
524 NULL,
525 &events[2+i]);
526 TestRequire(result == CL_SUCCESS, "clEnqueueReleaseD3D11ObjectsKHR failed.");
527 TestRequire(events[2+i], "clEnqueueReleaseD3D11ObjectsKHR did not return an event.");
528
529 // make sure the event type is correct
530 cl_uint eventType = 0;
531 result = clGetEventInfo(
532 events[2+i],
533 CL_EVENT_COMMAND_TYPE,
534 sizeof(eventType),
535 &eventType,
536 NULL);
537 TestRequire(result == CL_SUCCESS, "clGetEventInfo for event created by clEnqueueReleaseD3D11ObjectsKHR failed.");
538 TestRequire(eventType == CL_COMMAND_RELEASE_D3D11_OBJECTS_KHR, "clGetEventInfo for CL_EVENT_COMMAND_TYPE was not CL_COMMAND_RELEASE_D3D11_OBJECTS_KHR.");
539 }
540
541 for (UINT i = 0; i < size->SubResourceCount; ++i)
542 for (UINT x = 0; x < 2; ++x)
543 for (UINT y = 0; y < 2; ++y)
544 {
545 // create the staging buffer
546 ID3D11Texture2D* pStagingBuffer = NULL;
547 {
548 D3D11_TEXTURE2D_DESC desc = {0};
549 desc.Width = 1;
550 desc.Height = 1;
551 desc.MipLevels = 1;
552 desc.ArraySize = 1;
553 desc.Format = format->format;
554 desc.SampleDesc.Count = 1;
555 desc.SampleDesc.Quality = 0;
556 desc.Usage = D3D11_USAGE_STAGING;
557 desc.BindFlags = 0;
558 desc.CPUAccessFlags = D3D11_CPU_ACCESS_READ | D3D11_CPU_ACCESS_WRITE;
559 desc.MiscFlags = 0;
560 hr = pDevice->CreateTexture2D(&desc, NULL, &pStagingBuffer);
561 TestRequire(SUCCEEDED(hr), "Failed to create staging buffer.");
562 }
563
564 // wipe out the staging buffer to make sure we don't get stale values
565 {
566 D3D11_MAPPED_SUBRESOURCE mappedTexture;
567 hr = pDC->Map(
568 pStagingBuffer,
569 0,
570 D3D11_MAP_READ_WRITE,
571 0,
572 &mappedTexture);
573 TestRequire(SUCCEEDED(hr), "Failed to map staging buffer");
574 memset(mappedTexture.pData, 0, format->bytesPerPixel);
575 pDC->Unmap(pStagingBuffer, 0);
576 }
577
578 // copy the pixel to the staging buffer
579 {
580 D3D11_BOX box = {0};
581 box.left = x ? subResourceInfo[i].width - 2 : 1; box.right = box.left + 1;
582 box.top = y ? subResourceInfo[i].height - 2 : 1; box.bottom = box.top + 1;
583 box.front = 0; box.back = 1;
584 pDC->CopySubresourceRegion(
585 pStagingBuffer,
586 0,
587 0,
588 0,
589 0,
590 pTexture,
591 subResourceInfo[i].subResource,
592 &box);
593 }
594
595 // make sure we read back what was written next door
596 {
597 D3D11_MAPPED_SUBRESOURCE mappedTexture;
598 hr = pDC->Map(
599 pStagingBuffer,
600 0,
601 D3D11_MAP_READ_WRITE,
602 0,
603 &mappedTexture);
604 TestRequire(SUCCEEDED(hr), "Failed to map staging buffer");
605 TestRequire(
606 !memcmp(mappedTexture.pData, texture2DPatterns[x][y], format->bytesPerPixel),
607 "Failed to map staging buffer");
608 pDC->Unmap(pStagingBuffer, 0);
609 }
610
611 pStagingBuffer->Release();
612 }
613
614
615 Cleanup:
616
617 if (pTexture)
618 {
619 pTexture->Release();
620 }
621 for (UINT i = 0; i < size->SubResourceCount; ++i)
622 {
623 clReleaseMemObject(subResourceInfo[i].mem);
624 }
625 for (UINT i = 0; i < 4; ++i)
626 {
627 if (events[i])
628 {
629 result = clReleaseEvent(events[i]);
630 TestRequire(result == CL_SUCCESS, "clReleaseEvent for event failed.");
631 }
632 }
633
634
635 HarnessD3D11_TestEnd();
636 }
637
is_format_supported(cl_channel_order channel_order,cl_channel_type channel_type,const std::vector<cl_image_format> & supported_image_formats)638 bool is_format_supported(
639 cl_channel_order channel_order,
640 cl_channel_type channel_type,
641 const std::vector<cl_image_format> &supported_image_formats)
642 {
643 for (std::vector<cl_image_format>::const_iterator it = supported_image_formats.begin(); it != supported_image_formats.end(); ++it)
644 if (it->image_channel_data_type == channel_type && it->image_channel_order == channel_order)
645 return true;
646
647 return false;
648 }
649
TestDeviceTexture2D(cl_device_id device,cl_context context,cl_command_queue command_queue,ID3D11Device * pDevice,ID3D11DeviceContext * pDC)650 void TestDeviceTexture2D(
651 cl_device_id device,
652 cl_context context,
653 cl_command_queue command_queue,
654 ID3D11Device* pDevice,
655 ID3D11DeviceContext* pDC)
656 {
657 cl_int result = CL_SUCCESS;
658 cl_kernel kernels[3] = {NULL, NULL, NULL};
659
660 const char *sourceRaw =
661 " \
662 __kernel void texture2D\n\
663 ( \n\
664 __read_only image2d_t texIn, \n\
665 __write_only image2d_t texOut \n\
666 ) \n\
667 { \n\
668 const sampler_t smp = CLK_FILTER_NEAREST; \n\
669 CLK_NORMALIZED_COORDS_FALSE |\n\
670 CLK_ADDRESS_CLAMP_TO_EDGE; \n\
671 %s value; \n\
672 int2 coordIn; \n\
673 int2 coordOut; \n\
674 int w = get_image_width(texIn); \n\
675 int h = get_image_height(texIn); \n\
676 \n\
677 coordIn = (int2)(0, 0); \n\
678 coordOut = (int2)(1, 1); \n\
679 value = read_image%s(texIn, smp, coordIn); \n\
680 write_image%s(texOut, coordOut, value); \n\
681 \n\
682 coordIn = (int2)(w-1, h-1); \n\
683 coordOut = (int2)(w-2, h-2); \n\
684 value = read_image%s(texIn, smp, coordIn); \n\
685 write_image%s(texOut, coordOut, value); \n\
686 } \n\
687 ";
688
689 cl_uint supported_formats_count;
690 std::vector<cl_image_format> supported_image_formats;
691 result = clGetSupportedImageFormats(context, CL_MEM_READ_ONLY, CL_MEM_OBJECT_IMAGE2D, 0, NULL, &supported_formats_count);
692 TestRequire(CL_SUCCESS == result, "clGetSupportedImageFormats failed.");
693
694 supported_image_formats.resize(supported_formats_count);
695 result = clGetSupportedImageFormats(context, CL_MEM_READ_ONLY, CL_MEM_OBJECT_IMAGE2D, supported_formats_count, &supported_image_formats[0], NULL);
696 TestRequire(CL_SUCCESS == result, "clGetSupportedImageFormats failed.");
697
698 char source[2048];
699 sprintf(source, sourceRaw, "float4", "f", "f", "f", "f");
700 result = HarnessD3D11_CreateKernelFromSource(&kernels[0], device, context, source, "texture2D");
701 TestRequire(CL_SUCCESS == result, "HarnessD3D11_CreateKernelFromSource failed.");
702
703 sprintf(source, sourceRaw, "uint4", "ui", "ui", "ui", "ui");
704 result = HarnessD3D11_CreateKernelFromSource(&kernels[1], device, context, source, "texture2D");
705 TestRequire(CL_SUCCESS == result, "HarnessD3D11_CreateKernelFromSource failed.");
706
707 sprintf(source, sourceRaw, "int4", "i", "i", "i", "i");
708 result = HarnessD3D11_CreateKernelFromSource(&kernels[2], device, context, source, "texture2D");
709 TestRequire(CL_SUCCESS == result, "HarnessD3D11_CreateKernelFromSource failed.");
710
711 for (UINT format = 0, size = 0; format < formatCount; ++size, ++format)
712 {
713 if (!is_format_supported(formats[format].channel_order, formats[format].channel_type, supported_image_formats))
714 {
715 HarnessD3D11_TestBegin("2D_texture: Format=%s, Width=%d, Height=%d, MipLevels=%d, ArraySize=%d\n",
716 formats[format].name_format,
717 texture2DSizes[size % texture2DSizeCount].Width,
718 texture2DSizes[size % texture2DSizeCount].Height,
719 texture2DSizes[size % texture2DSizeCount].MipLevels,
720 texture2DSizes[size % texture2DSizeCount].ArraySize);
721 log_info("\tFormat not supported, skipping test!\n");
722 HarnessD3D11_TestEnd();
723
724 continue;
725 }
726
727 SubTestTexture2D(
728 context,
729 command_queue,
730 kernels[formats[format].generic],
731 pDevice,
732 pDC,
733 &formats[format],
734 &texture2DSizes[size % texture2DSizeCount]);
735 }
736
737 Cleanup:
738
739
740 for (UINT i = 0; i < 3; ++i)
741 {
742 if (kernels[i])
743 {
744 clReleaseKernel(kernels[i]);
745 }
746 }
747 }
748
749
750