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