From dbd33bc9cfd2ace62445a812a6aabb901c2f7e74 Mon Sep 17 00:00:00 2001 From: Nikhil Joshi Date: Tue, 4 Oct 2022 21:30:03 +0530 Subject: [PATCH] External sharing new updates (#1482) * Fix enqueue_flags test to use correct barrier type. Currently, enqueue_flags test uses CLK_LOCAL_MEM_FENCE. Use CLK_GLOBAL_MEM_FENCE instead as all threads across work-groups need to wait here. * Add check for support for Read-Wrie images Read-Write images have required OpenCL 2.x. Read-Write image tests are already being skipped for 1.x devices. With OpenCL 3.0, read-write images being optional, the tests should be run or skipped depending on the implementation support. Add a check to decide if Read-Write images are supported or required to be supported depending on OpenCL version and decide if the tests should be run on skipped. Fixes issue #894 * Fix formatting in case of Read-Write image checks. Fix formatting in case of Read-write image checks. Also, combine two ifs into one in case of kerne_read_write tests * Fix some more formatting for RW-image checks Remove unnecessary spaces at various places. Also, fix lengthy lines. * Fix malloc-size calculation in test imagedim unsigned char size is silently assumed to be 1 in imagedim test of test_basic. Pass sizeof(type) in malloc size calculation. Also, change loop variable from signed to unsigned. Add checks for null pointer for malloced memory. * Initial CTS for external sharing extensions Initial set of tests for below extensions with Vulkan as producer 1. cl_khr_external_memory 2. cl_khr_external_memory_win32 3. cl_khr_external_memory_opaque_fd 4. cl_khr_external_semaphore 5. cl_khr_external_semaphore_win32 6. cl_khr_external_semaphore_opaque_fd * Updates to external sharing CTS Updates to external sharing CTS 1. Fix some build issues to remove unnecessary, non-existent files 2. Add new tests for platform and device queries. 3. Some added checks for VK Support. * Update CTS build script for Vulkan Headers Update CTS build to clone Vulkan Headers repo and pass it to CTS build in preparation for external memory and semaphore tests * Fix Vulkan header path Fix Vulkan header include path. * Add Vulkan loader dependency Vulkan loader is required to build test_vulkan of OpenCL-CTS. Clone and build Vulkan loader as prerequisite to OpenCL-CTS. * Fix Vulkan loader path in test_vulkan Remove arch/os suffix in Vulkan loader path to match vulkan loader repo build. * Fix warnings around getHandle API. Return type of getHandle is defined differently based on win or linux builds. Use appropriate guards when using API at other places. While at it remove duplicate definition of ARRAY_SIZE. * Use ARRAY_SIZE in harness. Use already defined ARRAY_SIZE macro from test_harness. * Fix build issues for test_vulkan Fix build issues for test_vulkan 1. Add cl_ext.h in common files 2. Replace cl_mem_properties_khr with cl_mem_properties 3. Replace cl_external_mem_handle_type_khr with cl_external_memory_handle_type_khr 4. Type-cast malloc as required. * Fix code formatting. Fix code formatting to get CTS CI builds clean. * Fix formatting fixes part-2 Another set of formatting fixes. * Fix code formatting part-3 Some more code formatting fixes. * Fix code formatting issues part-4 More code formatting fixes. * Formatting fixes part-5 Some more formatting fixes * Fix formatting part-6 More formatting fixes continued. * Code formatting fixes part-7 Code formatting fixes for image * Code formatting fixes part-8 Fixes for platform and device query tests. * Code formatting fixes part-9 More formatting fixes for vulkan_wrapper * Code formatting fixes part-10 More fixes to wrapper header * Code formatting fixes part-11 Formatting fixes for api_list * Code formatting fixes part-12 Formatting fixes for api_list_map. * Code formatting changes part-13 Code formatting changes for utility. * Code formatting fixes part-15 Formatting fixes for wrapper. * Misc Code formatting fixes Some more misc code formatting fixes. * Fix build breaks due to code formatting Fix build issues arised with recent code formatting issues. * Fix presubmit script after merge Fix presubmit script after merge conflicts. * Fix Vulkan loader build in presubmit script. Use cmake ninja and appropriate toolchain for Vulkan loader dependency to fix linking issue on arm/aarch64. * Use static array sizes Use static array sizes to fix windows builds. * Some left-out formatting fixes. Fix remaining formatting issues. * Fix harness header path Fix harness header path While at it, remove Misc and test pragma. * Add/Fix license information Add Khronos License info for test_vulkan. Replace Apple license with Khronos as applicable. * Fix headers for Mac OSX builds. Use appropriate headers for Mac OSX builds * Fix Mac OSX builds. Use appropriate headers for Mac OSX builds. Also, fix some build issues due to type-casting. * Fix new code formatting issues Fix new code formatting issues with recent MacOS fixes. * Add back missing case statement Add back missing case statement that was accidentally removed. * Disable USE_GAS for Vulkan Loader build. Disable USE_GAS for Vulkan Loader build to fix aarch64 build. * Fixes to OpenCL external sharing tests Fix clReleaseSemaphore() API. Fix copyright year. Some other minor fixes. * Improvements to OpenCL external sharing CTS Use SPIR-V shaders instead of NV extension path from GLSL to Vulkan shaders. Fixes for lower end GPUs to use limited memory. Update copy-right year at some more places. * Fix new code formatting issues. Fix code formatting issues with recent changes for external sharing tests. * More formatting fixes. More formatting fixes for recent updates to external sharing tests. * Final code formatting fixes. Minor formatting fixes to get format checks clean. --- test_conformance/vulkan/main.cpp | 4 +- test_conformance/vulkan/shaders/buffer.comp | 28 ++++ test_conformance/vulkan/shaders/buffer.spv | Bin 0 -> 2168 bytes test_conformance/vulkan/shaders/image2D.comp | 31 ++++ .../vulkan/shaders/image2D_r16i.spv | Bin 0 -> 3264 bytes .../vulkan/shaders/image2D_r16ui.spv | Bin 0 -> 3264 bytes .../vulkan/shaders/image2D_r32f.spv | Bin 0 -> 3268 bytes .../vulkan/shaders/image2D_r32i.spv | Bin 0 -> 3256 bytes .../vulkan/shaders/image2D_r32ui.spv | Bin 0 -> 3256 bytes .../vulkan/shaders/image2D_r8i.spv | Bin 0 -> 3264 bytes .../vulkan/shaders/image2D_r8ui.spv | Bin 0 -> 3264 bytes .../vulkan/shaders/image2D_rg16i.spv | Bin 0 -> 3264 bytes .../vulkan/shaders/image2D_rg16ui.spv | Bin 0 -> 3264 bytes .../vulkan/shaders/image2D_rg32f.spv | Bin 0 -> 3276 bytes .../vulkan/shaders/image2D_rg32i.spv | Bin 0 -> 3264 bytes .../vulkan/shaders/image2D_rg32ui.spv | Bin 0 -> 3264 bytes .../vulkan/shaders/image2D_rg8i.spv | Bin 0 -> 3264 bytes .../vulkan/shaders/image2D_rg8ui.spv | Bin 0 -> 3264 bytes .../vulkan/shaders/image2D_rgba16i.spv | Bin 0 -> 3256 bytes .../vulkan/shaders/image2D_rgba16ui.spv | Bin 0 -> 3256 bytes .../vulkan/shaders/image2D_rgba32f.spv | Bin 0 -> 3268 bytes .../vulkan/shaders/image2D_rgba32i.spv | Bin 0 -> 3256 bytes .../vulkan/shaders/image2D_rgba32ui.spv | Bin 0 -> 3256 bytes .../vulkan/shaders/image2D_rgba8i.spv | Bin 0 -> 3256 bytes .../vulkan/shaders/image2D_rgba8ui.spv | Bin 0 -> 3256 bytes .../vulkan/test_vulkan_api_consistency.cpp | 14 +- .../vulkan/test_vulkan_interop_buffer.cpp | 36 +---- .../vulkan/test_vulkan_interop_image.cpp | 142 ++++++------------ .../opencl_vulkan_wrapper.cpp | 51 ++++++- .../opencl_vulkan_wrapper.hpp | 6 +- .../vulkan_interop_common/vulkan_list_map.hpp | 7 +- .../vulkan_interop_common/vulkan_utility.cpp | 105 +++++++------ .../vulkan_interop_common/vulkan_utility.hpp | 1 + .../vulkan_interop_common/vulkan_wrapper.cpp | 14 +- .../vulkan_interop_common/vulkan_wrapper.hpp | 3 +- 35 files changed, 230 insertions(+), 212 deletions(-) create mode 100644 test_conformance/vulkan/shaders/buffer.comp create mode 100644 test_conformance/vulkan/shaders/buffer.spv create mode 100644 test_conformance/vulkan/shaders/image2D.comp create mode 100644 test_conformance/vulkan/shaders/image2D_r16i.spv create mode 100644 test_conformance/vulkan/shaders/image2D_r16ui.spv create mode 100644 test_conformance/vulkan/shaders/image2D_r32f.spv create mode 100644 test_conformance/vulkan/shaders/image2D_r32i.spv create mode 100644 test_conformance/vulkan/shaders/image2D_r32ui.spv create mode 100644 test_conformance/vulkan/shaders/image2D_r8i.spv create mode 100644 test_conformance/vulkan/shaders/image2D_r8ui.spv create mode 100644 test_conformance/vulkan/shaders/image2D_rg16i.spv create mode 100644 test_conformance/vulkan/shaders/image2D_rg16ui.spv create mode 100644 test_conformance/vulkan/shaders/image2D_rg32f.spv create mode 100644 test_conformance/vulkan/shaders/image2D_rg32i.spv create mode 100644 test_conformance/vulkan/shaders/image2D_rg32ui.spv create mode 100644 test_conformance/vulkan/shaders/image2D_rg8i.spv create mode 100644 test_conformance/vulkan/shaders/image2D_rg8ui.spv create mode 100644 test_conformance/vulkan/shaders/image2D_rgba16i.spv create mode 100644 test_conformance/vulkan/shaders/image2D_rgba16ui.spv create mode 100644 test_conformance/vulkan/shaders/image2D_rgba32f.spv create mode 100644 test_conformance/vulkan/shaders/image2D_rgba32i.spv create mode 100644 test_conformance/vulkan/shaders/image2D_rgba32ui.spv create mode 100644 test_conformance/vulkan/shaders/image2D_rgba8i.spv create mode 100644 test_conformance/vulkan/shaders/image2D_rgba8ui.spv diff --git a/test_conformance/vulkan/main.cpp b/test_conformance/vulkan/main.cpp index 6cbde5cc..2eeb0c36 100644 --- a/test_conformance/vulkan/main.cpp +++ b/test_conformance/vulkan/main.cpp @@ -134,7 +134,6 @@ cl_device_id *devices; const size_t bufsize = BUFFERSIZE; char buf[BUFFERSIZE]; cl_uchar uuid[CL_UUID_SIZE_KHR]; -VulkanDevice vkDevice; unsigned int numCQ; bool multiImport; bool multiCtx; @@ -220,9 +219,12 @@ int main(int argc, const char *argv[]) if (!checkVkSupport()) { log_info("Vulkan supported GPU not found \n"); + log_info("TEST SKIPPED \n"); return 0; } + VulkanDevice vkDevice; + cl_device_type requestedDeviceType = CL_DEVICE_TYPE_GPU; char *force_cpu = getenv("CL_DEVICE_TYPE"); if (force_cpu != NULL) diff --git a/test_conformance/vulkan/shaders/buffer.comp b/test_conformance/vulkan/shaders/buffer.comp new file mode 100644 index 00000000..d8756f92 --- /dev/null +++ b/test_conformance/vulkan/shaders/buffer.comp @@ -0,0 +1,28 @@ +#version 450 +#extension GL_ARB_separate_shader_objects : enable +#extension GL_EXT_shader_explicit_arithmetic_types_int8 : enable +#extension GL_EXT_shader_explicit_arithmetic_types_int32 : enable + +#define MAX_BUFFERS 5 + +layout(binding = 0) buffer Params +{ + uint32_t numBuffers; + uint32_t bufferSize; + uint32_t interBufferOffset; +}; +layout(binding = 1) buffer Buffer +{ + uint8_t ptr[]; +} bufferPtrList[MAX_BUFFERS]; +layout(local_size_x = 512) in; +void main() { + for (uint32_t bufIdx = 0; bufIdx < numBuffers; bufIdx++) { + uint32_t ptrIdx = gl_GlobalInvocationID.x; + uint32_t limit = bufferSize; + while (ptrIdx < limit) { + bufferPtrList[bufIdx].ptr[ptrIdx]++; + ptrIdx += (gl_NumWorkGroups.x * gl_WorkGroupSize.x); + } + } +} \ No newline at end of file diff --git a/test_conformance/vulkan/shaders/buffer.spv b/test_conformance/vulkan/shaders/buffer.spv new file mode 100644 index 0000000000000000000000000000000000000000..685523ba5fbed55d32c179504d1d26834767d78e GIT binary patch literal 2168 zcmb7^Sx*yD6vwYp$|m4~J76j93!*Gm5EK+gVk1P1`Zk%iQyuM=>6Cy^KKOzB41Or1 zCMN!W?Y+sodYhAT&j0M^-f7>_*#TpYm|-(&JkvWT%uyl6z^_>zoG|hkFnhai{g*FZ z`-M`J`0b<_2OmPBMohLF;}!RcbM2%ux41M9K4$u0UE8?g`z2>Y^&o11Pl+Chjw>Eq zWsJ&p{pq1kKdI)2l#CfNUeVupz3I2ZRuBhC=(j%xl`!_3rB7iwY2%|givHI7H_j9u zv}#dV-3D=#e5{8_RQ8j@R@nBVMzU~Ux&Je!pb*6jQ^(uc@rql1xz*A#4r`%biyEO{ ziDGpV?Q1g`Gnn#=8s%E2l6L5qM|vOVgUVm(RJSV!M*GTWH^%I0+`8t=>W9v0bn2U( zYBh{gI6G6Uz`f{;>SvV=e|D&O!?+i*Q?0hcL|DJ}ds3W3tqbCs9Q&~X^Eg0Vk`F((*1^fVi>yg*N0%$XN)oKF87bl&2K zXwvF);#m=O$j@vX*9ZB*Tz+Ot=SL?$@#N>7lmDF6pIiOB_$l2Gz0W8c`H)X5hWa@X zCnsGWYnQ-pS^E%1ooTx(Y>tTGn<5{(ShYI4a5YXz&WY%goahrG4)LeO@g?VkB>F!q z{G!F^7eDq!4shzAvr{na(^hu9QzvgRX1ybsd*e)L3+x|kS~7b^cQemR=D%c8areZr zgHeZXnts9Q1D(47!@g`~a#Q!AICbfh{OBuIcXM!e&QAZ2MI?0nY+3tlZHf7xi|_%5 zmenowd8-_D_DcveP=kA=2e{Mu7SsBS)z=k`FZcdiIfli_1^=mB$HYfP9PDG_*x_rk z)Bm`L92@F@`#mG#O>pnd=d8uQ%ZEGXoe~HBM|seB&(tG7GWodYOCs)`9$mhiWN`Xt zHgs2iRxGWHY$?0ahEg2th zW<$R!A~(5^UF@P{cIWgZ$>iS9ZSl4qh|eDf^i_+kh%+ZPFt^7w z$-KSv9@nLVe^UuCj literal 0 HcmV?d00001 diff --git a/test_conformance/vulkan/shaders/image2D.comp b/test_conformance/vulkan/shaders/image2D.comp new file mode 100644 index 00000000..42fa2f73 --- /dev/null +++ b/test_conformance/vulkan/shaders/image2D.comp @@ -0,0 +1,31 @@ +#version 450 +#extension GL_ARB_separate_shader_objects : enable +#extension GL_EXT_shader_explicit_arithmetic_types_int32 : enable + +#define MAX_2D_IMAGES 5 +#define MAX_2D_IMAGE_MIP_LEVELS 11 +#define MAX_2D_IMAGE_DESCRIPTORS MAX_2D_IMAGES * MAX_2D_IMAGE_MIP_LEVELS + +layout(binding = 0) buffer Params +{ + uint32_t numImage2DDescriptors; +}; +layout(binding = 1, rgba32f ) uniform image2D image2DList[ MAX_2D_IMAGE_DESCRIPTORS ]; +layout(local_size_x = 32, local_size_y = 32) in; +void main() { + uvec3 numThreads = gl_NumWorkGroups * gl_WorkGroupSize; + for (uint32_t image2DIdx = 0; image2DIdx < numImage2DDescriptors; image2DIdx++) { + ivec2 imageDim = imageSize(image2DList[image2DIdx]); + uint32_t heightBy2 = imageDim.y / 2; + for (uint32_t row = gl_GlobalInvocationID.y; row < heightBy2; row += numThreads.y) { + for (uint32_t col = gl_GlobalInvocationID.x; col < imageDim.x; col += numThreads.x) { + ivec2 coordsA = ivec2(col, row); + ivec2 coordsB = ivec2(col, imageDim.y - row - 1); + vec4 dataA = imageLoad(image2DList[image2DIdx], coordsA); + vec4 dataB = imageLoad(image2DList[image2DIdx], coordsB); + imageStore(image2DList[image2DIdx], coordsA, dataB); + imageStore(image2DList[image2DIdx], coordsB, dataA); + } + } + } +} \ No newline at end of file diff --git a/test_conformance/vulkan/shaders/image2D_r16i.spv b/test_conformance/vulkan/shaders/image2D_r16i.spv new file mode 100644 index 0000000000000000000000000000000000000000..00c5c2833f43ad1d329f438b740b24f020c9e4fe GIT binary patch literal 3264 zcmZ9N4Obgg5QbMsD74ZRp~d=915~QEqS6o1N?T$ujg18ah`CFI-|}nPJfL#17aKp{YIwe zrH?s7DPOJMt~&kAA11CBhe@*? zB%$vn``w`DhMnXRy%%&s?WwM;RO`#G-|f1+WWN=-t*{fgjWAZ`a5o@sI2RXo{MJq* zA5c_8{%l2^c0~C*JMHym9C!`&SJWMF=B<`{XQzEHinptAw9}ObU(tC=I&6De!F;9G z*n=xL&xl*qaodU-QbsUKd29X9O3?G;u$x3NiHf?j=hBFN&RN-M&Pv#Zrp^W9<}3BE zm!uruvl)b2&7{1~y^v=?I*uO79!|fL`mVO3O|Mn!>_)zqgi)tf$zqqK{ivnbg0q&| zQQYV)rFi2%Tu!mB(`k5#w}3P{D2{p+DJ}Y5S>ji0I$u}WuTAH(;lO`4)aW6{K6Z{#??0ISYT;jD>yGpvk#7(uSm1t=^e~>WZ{3cZ>D>|%rHGr*UoZ9HZ|@3`R?>Y{M(t$**KFg z&0gL#{)RO9;4R;wZ1!K(jFOjF7Bldtn4&bbm<4_d*bnBL@zxD5LQ7oE80t za-Ei*l(6vghVi4R#r^WG$ic4zGgXo>2h5tqOvx@uepAe}^o)dk@P8{0_>_ctvzmn3k!IrcLbmnHDODh7NtAjQC!1dR=9aH28=hX7MKHbGl4d`eUf-9#En(rWOXEjd zt`B6B!|wZz>{*G`<9=5p@OHPW8E5DDP?$XMoCjQ!P@j6(R`0HC>XD1w>k>Hge)6^aBYo@YD!0&T6!T^EKMS)PT1d9og_!Ba+?kwO^wvObk3Vwx!{$#@Dj*HDY0E zz*~);YKhz)= zKOAr6wluSgW_HG;zm<$g#wFIrql`EFoiOuvkj={Xvf+s1uJFO--z)C~d{dg5Xsh{y zY--xPGJgjWc&quN>;;L{d@LL8KtfG?aMW0orpAp7Pi7cx`F@hk-CDk%Wy2jx$cGP( Rd=t{-J1zZ>{v$6){snCV9kBoa literal 0 HcmV?d00001 diff --git a/test_conformance/vulkan/shaders/image2D_r16ui.spv b/test_conformance/vulkan/shaders/image2D_r16ui.spv new file mode 100644 index 0000000000000000000000000000000000000000..87514d9fab1809a8a08645f9fc8a489463837db2 GIT binary patch literal 3264 zcmZ9NiE>m`5QeXigiQow6GaS*pn|Aj5k;27#32#^0}8IgWagS&GMh6uKok`b^>us( zAIegztn&M2Zd^mq&h0ZvPBZdE@xWQ?I+@sbsBc;MXL+^O%krgG znrHQDzS>Ru)vTSLu=f#8$ezm8tChK#YQ5X7_Vd+dT5V?SwA#peWX|rU#;q^KW$k)% zrBMzTC8B(`QBJ#o{Oy(2jip{%YtY{)ci@>9o7F2Tt(%?Ra;4W<>Eh8hnmmopTD8S= zVtTf5Uv4CM7Va9yZ4os_MzM>$v%}Eov|sOK-MrH?(I|JecMvhmc^G?;b2@8DTjywS z6Vr29KOb;@&!se5T*{|by%+PGLH9cMu-6a16Ma{jorPL+wtcr#ujN^%Jv&{*&Y*iazuDX7sr~@@&U;?JMqS*y~90 z9Gv}oclIFm%yFjYvBwc>noGP7aarpHw6W&ffpq}w*$<(O)o0v`XnT_1!^|E*yT`oc zW&B|z`hE%fO=K@QUPil5&YZ8HUGnxW_B*=tzdAGrd&JJLJ=j;A9dnsgBn6P3v!<%_krE9%H-zG&5>mV_nX`pD?3nYdH)47Pw#RH}kCr_P)G_!1{fU zW$arlclIY))tGEaLn{PUGJxc;o%vV(;)A+B2%X;|u7Eh)e$@TE9B(U<%tD zf#1cphq%vaY-8lTPw{!g+~(4bUf;vESHF4Wt6m|73{-E)boC?BJy#!*9tDq^C8$g@}5ULi&)=!+EMR1w)Mt- z?!x~uw%j$oowITiagV&`pD(!B{|UI9|5LDBobNMi&llg`9zI9p<9uHfT;%)`Y)_Bq zVGLbGjFAuSHnzNZqxTy2Z6tbMz_uTGd%uO2i#lInm*37J*ta8ZjBjTNkr#U(8R8u* z&|4jS2Qfz88fn2rjV8EUqXo8xeAH-T%SVk4_ORDsjV{<2d21}A<)X&d*yS2Mur=hP zMju-~YUJ4S1@;^(hK9gP40 literal 0 HcmV?d00001 diff --git a/test_conformance/vulkan/shaders/image2D_r32f.spv b/test_conformance/vulkan/shaders/image2D_r32f.spv new file mode 100644 index 0000000000000000000000000000000000000000..e82c9c190df248347e32b8354713ffb95d8d23f0 GIT binary patch literal 3268 zcmZ9NiINjl6oy+!hE4WO6frE~0*=Zii{Q*SOc;%TfC6qIlQfftY$iRxC@Lbz>-Y>l zl%-Z#<@a^^T3z0$Q|J8uIrl8L@6D`Sv42$*Jr}Kw)<@5x@!A}%fQcgY5%jufBE!|% z)!O7JZA~6K{*I9wqLr%5w=r53m9a}mw~_SZ*CD5oP3)K7j8t+l|4RN=pnnIHqBT*q zR-e8yQy;~H#;}pb^-;Ugiih?7!kxI8j*MAv%=sJFf-8P7=p@Y~tv7~A+U~|_(yXV; zgLqU=dg)PnFEJT=s+TTRYv=3D!Js}$mpgI2ll05L=3_W_Ewr z8ck>T=6`r5!?(e$M%tK08{f(C47v*TjFcCU3q*vO_k&$kdXw{Z8SlLN5Ldu^$1uar zTl9-@v8B+nTM$0MUwP@>! z&*a#ea)-d?TaQ!qF zOfO?s5Nn!Cya#bv>lL)I=G%s~5AEy+(8lUB?p3rs$?u?t*D^kec5#1S#}-?oxZLa6 zg&E$!eh=AAzBkeClQaJzv`gMO#C}tk{>ibK?GYaZ+mAhlxlUkPGkgp0fIS(1Cbyl_ z`C4duE+P8)Imq(KhkQ@5-G74nYF=Yq&h#ft8Eq}^g})2-Bli3G?gRGh4i>O}-(&EL zeV0W%SHSM~{RO|+_ZM)Cefxjd_ox3X{ivTuyqizxTD&ReynpAI$o}QD?cMwK7c*`{ zJFEH+a_vCxL|poP&-&HY@_c<`=J1!o`&B`_lis@!vlqL9{6@?^^nS#B@_&;@ydANg z`L)B|x3EKg@5S4QcSYV9@0I_oy|G*awkN&-1;E$cMB2H0Q!RSHR|xcOLOai1n?f9qL`fww}4n{TU(` z{GVfccENugTkbO7&bx90agV(7U(LC&{|j(2|CeC7FyB|$&KJJDJ$#MGhxxwAxsbCC zwx`GRFoC{>7$YCJ2KFt)yrK66Z28c86We~|?foWNF4S4X-h$ZcB)X0GcI1um?Iei2 z*z>r9cm@mXX`$~T#>iXacFu(wU2w5R4{QzjP@|76A8HJ+?KRZ+7TXwkYjn_Zp~euq zSYrgXhJ2`zV#|jbOW4jKc8+_9d7EHk^~;4C_pysL9)PVOA8IUP%ZD1@VLL~t@jbTZ zV+~{V%lTGnXz#At@39&E5Ltt4MM57x&o<2PE4JKI V#C-bX%(o3~z8&a)_>X)N`4=`19>M?s literal 0 HcmV?d00001 diff --git a/test_conformance/vulkan/shaders/image2D_r32i.spv b/test_conformance/vulkan/shaders/image2D_r32i.spv new file mode 100644 index 0000000000000000000000000000000000000000..7ea8d26f9c273accbaee061c9b89743cfeb9aa3d GIT binary patch literal 3256 zcmZ9NiE>m`5QeXigiQow6GaS*pn|Aj5k;27#32#^0}8IgWagS&GMh6uKok`b^>us( zAIegztn&M2Zd^mq&h0ZTNzbL&R9()-BRJG z=)~2tOkB((0b-^|CR#pL2qXQox%N7v!3=&V#d(ph-){PM>61y83(UN z%PHcGXzPj37TB6{N5SUXjEtA?Q`q{C;dP&V8KdYkuWv@5`y^G6UBj00e_n&4)&1Plxxn_G{tqt~_#FKh_ax#xKBjB&)`D~9&T}IBm(#X)XY4P|xC8B3 z)qjv{7kW41((m`xueO%=>$@_CzYfmSIN}^SYZ0>_dmQm8i2LOKCXaY0Vm9yEjyNmw#vDPvhR830525{sCFlA47GnFGK%YS5d~@!1E>0ryzY-%pU2x9W z17bX-*@-h)e$*TE9AS zeSmF_xbG|2hmok~{a!`n<8H4NT%6}auzBP?k9ZcbzV)=D-gRv2nakWa5V`Pwgzeph z|6^>qYkWIryOPz9_iJ`6bw%9?`=X zx{4ShAKYzhdGki^HSF6+^uB;?Kl1i|3oRFQzQQiQokg&3N8T9U&JrRo_C7MiJ6NE% zI{FS`jJ!3{f{PkWaJfbcYz_IS(Z-gK8XfFmufrN$urc!1SVqf5jjyrGHF{uc$VZJn zwtUpcvF8iyIaUz!);XuK`sJd=UF>p=dthtGM~(Z~@=;?I+jEGm@eShrSi@NTa=w+z zXlGaL>}*4Si)=u)Bhkl$f)D&1*!g={%*yxJa>jXA`sB**)psJkh_5m<0 literal 0 HcmV?d00001 diff --git a/test_conformance/vulkan/shaders/image2D_r32ui.spv b/test_conformance/vulkan/shaders/image2D_r32ui.spv new file mode 100644 index 0000000000000000000000000000000000000000..dbcdbc5f59437f9b98223a30b7e547e1f83e30e9 GIT binary patch literal 3256 zcmZ9NiE>m`5QeXigiQow6Gcn_K?PC6B8n`Di9;j;1{7R}$;>smWHx7RSQHfz^>us( zAIegztn&M2Zd^mq&h0Z*tMQCr-X@OMUbVA$JFVw^V>TNzd-Hm@(tF)zR?qTkt(WD?tu)W- z)%<=p?N_sQe%#(iI2n5?SFTj%XRGyYx7yF|H`8h}Yp2yl)+2LvCpB(;DK2Z*o2!j- zz$g*rvyF1v4didHwyrPt(prQ5M!5seywt2-UTxjz^j0dp&T1ErzR~1qbk?dZrIRyr zjeBw<$+K`*Ic|%nF*1r>HFS`c}}Byox9lU2j7XlE6vVgtvT1e)2Y|;tka&GDPm{Q^-h!6 zkz`?DcY2Nf)BxZ74^Iy;*O@eGd2I@9{7QkR(G{?Fq&$OMBqGjy0PL#Jo1D{SyyxAI zxB^cM;Q)KyvR|y9u^RZDLF<{v{aXsW2fdvMbOtA=W#uNuGnV zfA7v7#GX0M^gQ-BVoh_2_aQE8y?{2>d^@oAqdogUw6Xe(dl79<@_U%s!)W)Ix4ev> zK%(!Lu-`=XlH+Bx`{c~|3fd)a?_$5BOaH4wbFfG34BLZ!#aT{bTQlz8?`}`Vzg^h& z>6z+iZQr#12HJe`k?%3K`%f{W<~7#k4EzZ*inf-s;BSHZ#eOs2dSLJCPzmezJ%(TG zTP)+*5_Z3DF8pHOTwr@o|A$u!e3pKUdlGRTAJesXYr#2l=Qxr5%W2!YGxir}+=2G2 z>OaV}3%wh0>GuumS6j;y(Gm$s^v0SkL_0 zvG*u;-v-<fBc4O7Z$0g(cMaQm<}&wnL@xXvVS9Js z{}@~DD&Njoxq-Mx-t#XMT4i@RH zj=qf;BX5ng;G#woT&~dqTSGo-w6W!*MhAP?>##-_Y>d1$R?u=$<7@14jULz<@=>FY zEgv;N^o%KwDECHGjaiX3VSe z_Yjefnm=NnMxy2;Y`KSsHTB6^<1E@5=L@{M!0O2N6Sj97`F_ThdyJS*pPc!2pv|`n L{SW_Vz@!~S literal 0 HcmV?d00001 diff --git a/test_conformance/vulkan/shaders/image2D_r8i.spv b/test_conformance/vulkan/shaders/image2D_r8i.spv new file mode 100644 index 0000000000000000000000000000000000000000..1a64147563f2356166183920112ca58aa9f3aeb2 GIT binary patch literal 3264 zcmZ9NiE>m`5QeXigiQow6GaS*pn|Aj5k;27#32#^0}8IgWagS&GMh6uKok`b^>us( zAIegztn&M2Zd^mq&h0ZvPBZdE@xWQ?I+@sbsBc;MXL+^O%krgG znrHQDzS>Ru)vTSLu=f#8$ezm8tChK#YQ5X7_Vd+dT5V?SwA#peWX|rU#;q^KW$k)% zrBMzTC8B(`QBJ#o{Oy(2jip{%YtY{)ci@>9o7F2Tt(%?Ra;4W<>Eh8hnmmopTD8S= zVtTf5Uv4CM7Va9yZ4os_MzM>$v%}Eov|sOK-MrH?(I|JecMvhmc^G?;b2@8DTjywS z6Vr29KOb;@&!se5T*{|by%+PGLH9cMu-6a16Ma{jorPL+wtcr#ujN^%Jv&{*&Y*iazuDX7sr~@@&U;?JMqS*y~90 z9Gv}oclIFm%yFjYvBwc>noGP7aarpHw6W&ffpq}w*$<(O)o0v`XnT_1!^|E*>o;$C z8807wzl8lJvX>k$qunQG&R5Vbd3zW89bNig9h!qZVrSSMBER|+V$HaJzq>sd|8`+} zHqTT?yT>=JzkxQNeB8}rZ1`sM^1l)zK3#Cm z*aKobr?V_?PQRzvoX*-=#QBSy#=leW#{0d+-r+g4XHji0;`4~P&7~c^zK3nEe)Gsz!IrmIXHje~m(cE4+w1%2%ZN+=99q9R za(#epj=1kD*oTp*=lxzqL$G<|J&$-6vA*@RquzCF>zT{kHxRk-e}wJb zh5uu0xodnoXXPg19(m6{UvRPi6L2~Ir(n4_-)GpKFTTA!e2&P+`MxN)$oVDMo*vP| z7`loWBOlytYtU`VvM{s(t?W`O>ntJ3v3PfsL{rjj~X59VXwm)U9d6o)>uZ%MUAhq%QbpnYsg29 zKDK<+$g$@O>^W8t^VT`1vHInr#$D`kjeB5g$VZL)*z!?h729)&t?>=w{aC|T{c^sQ z%V=j;?d)tre~WBDwjm`5QeXigiQow6GaS*pn|Aj5k;27#32#^0}8IgWagS&GMh6uKok`b^>us( zAIegztn&M2Zd^mq&h0ZvPBZdE@xWQ?I+@sbsBc;MXL+^O%krgG znrHQDzS>Ru)vTSLu=f#8$ezm8tChK#YQ5X7_Vd+dT5V?SwA#peWX|rU#;q^KW$k)% zrBMzTC8B(`QBJ#o{Oy(2jip{%YtY{)ci@>9o7F2Tt(%?Ra;4W<>Eh8hnmmopTD8S= zVtTf5Uv4CM7Va9yZ4os_MzM>$v%}Eov|sOK-MrH?(I|JecMvhmc^G?;b2@8DTjywS z6Vr29KOb;@&!se5T*{|by%+PGLH9cMu-6a16Ma{jorPL+wtcr#ujN^%Jv&{*&Y*iazuDX7sr~@@&U;?JMqS*y~90 z9Gv}oclIFm%yFjYvBwc>noGP7aarpHw6W&ffpq}w*$<(O)o0v`XnT_1!^|E*yT`oc zWxT%V`z7o*k-g-2x$w!E^A)s9-rmK2N0W(b~Rg{SCDFZ1`ji0;`4~P&7~c^zK3nEe)Gsz!IrmIXHje~m(cE4+w1%2%ZN+=99q9R za(#epj=1kD*oTp*=lxzqL$G<|J&$-6vA*@RquzCF>zT{kHxRk-e}wJb zh5uu0xodnoXXPg19(m6{UvRPi6L2~Ir(n4_-)GpKFTTA!e2&P+`MxN)$oVDMo*vP| z7`loWBOlytYtU`VvM{s(t?W`O>ntJ3v3PfsL{rjj~X59VXwm)U9d6o)>uZ%MUAhq%QbpnYsg29 zKDK<+$g$@O>^W8t^VT`1vHInr#$D`kjeB5g$VZL)*z!?h729)&t?>=w{aC|T{c^sQ z%V=j;?d)tre~WBDwjm`5QeXigiQow6GaS*pn|Aj5k;27#32#^0}8IgWagS&GMh6uKok`b^>us( zAIegztn&M2Zd^mq&h0ZvPBZdE@xWQ?I+@sbsBc;MXL+^O%krgG znrHQDzS>Ru)vTSLu=f#8$ezm8tChK#YQ5X7_Vd+dT5V?SwA#peWX|rU#;q^KW$k)% zrBMzTC8B(`QBJ#o{Oy(2jip{%YtY{)ci@>9o7F2Tt(%?Ra;4W<>Eh8hnmmopTD8S= zVtTf5Uv4CM7Va9yZ4os_MzM>$v%}Eov|sOK-MrH?(I|JecMvhmc^G?;b2@8DTjywS z6Vr29KOb;@&!se5T*{|by%+PGLH9cMu-6a16Ma{jorPL+wtcr#ujN^%Jv&{*&Y*iazuDX7sr~@@&U;?JMqS*y~90 z9Gv}oclIFm%yFjYvBwc>noGP7aarpHw6W&ffpq}w*$<(O)o0v`XnT_1!^|E*>o;$C z8E-Gq_e_W6E$`QNWe$HGoT+idIds+{Ww|bDK*$dVLSuUj62guYxUaug;>_UM`{CueR6s(U%dI{yDUM zb>#X0+Z=J5*w!&F-AVP+t~8vjoxe6w~^?50o#7$?fn*7F6w-RU4A=@VBe0sF}|H8L|*KDWQcdL zKyP*Q9mE)UYorAiHJadZjTYD%@=>FWEgv;H*u!3jHM(G92?L_TW%hm`5QeXigiQow6GaS*pn|Aj5k;27#32#^0}8IgWagS&GMh6uKok`b^>us( zAIegztn&M2Zd^mq&h0ZvPBZdE@xWQ?I+@sbsBc;MXL+^O%krgG znrHQDzS>Ru)vTSLu=f#8$ezm8tChK#YQ5X7_Vd+dT5V?SwA#peWX|rU#;q^KW$k)% zrBMzTC8B(`QBJ#o{Oy(2jip{%YtY{)ci@>9o7F2Tt(%?Ra;4W<>Eh8hnmmopTD8S= zVtTf5Uv4CM7Va9yZ4os_MzM>$v%}Eov|sOK-MrH?(I|JecMvhmc^G?;b2@8DTjywS z6Vr29KOb;@&!se5T*{|by%+PGLH9cMu-6a16Ma{jorPL+wtcr#ujN^%Jv&{*&Y*iazuDX7sr~@@&U;?JMqS*y~90 z9Gv}oclIFm%yFjYvBwc>noGP7aarpHw6W&ffpq}w*$<(O)o0v`XnT_1!^|E*yT`oc zWxO%b_eE#6vi&fIxUWdCy7_U?@R#Tj>? zJ*)Z;a_vIzMqK)R!}`_M@=kqM=J40SnHoo&LuV~w_G6DDzY%i)eGqY<{NLme??kL; ze(l(M3_J2W7snB2Mc$Yr=+_YWCGa7%|FGmdpWi}ke-r2vh@5ZE{m#@$ME+M|#HS0+ z8GAsC=X93k&FS|Po6}i4i#UIg)A)A^-gv*a*gHIj_Ka%p_yYPO;?h5f)~}8`n8G$k z;CHd@A?|Y;+ZcK8Q+ysVx4E>V*Y~jP)o&j8D%kS&>MV-wFMJ?wQ@qYE}h-Wtnjxv23qcDY6mYz_IS z(Z`mL8aeiSfj!3xV%|FEG*-V{)VPaXu5k}+4f&{XA6q_ZtYUi(u{FLyydP^At6$Ey zavAOHs-2x}=x>n?$aW<9cu??xzXLme4~tp(9$U^h?@FIs`Mvs1#23-lR7cGpu&o*M z>ij)KdJI|6b~B1MjF>g?M5pe*8B5!;$}KBX00*jZ(Iwm_`#r)G?TR67$#}E8>dOL zo-PgIQ9bFUN9?`CgzTwax>T*5uQvyS`Y2uM#Pv?ni|egqNao~TY~0F1T+(ZH7F)%D zG7-hu%1pb8{Jq8Qwe~P>wCJzQ894Jor+#sGYDec}}8-{rlJ}v+qRT)lPrD(V6Sr>o*%|((ld9Cc2((3&f8_Y^X@@h0q-8e z3_EYpFV@es%KWF$)-#X$*XMW#dNViBJGh%_*3;ff%mliExVD3NEE#7^1-ue1r-)ah zttUR6V{6JC0Gn?uQYqkd*!mCRb)S71qv$iQZ$_W{L!Qmpu3h;#je9LcufW;AXJ-#$ zXAU#Hid{jhX)f_D#AU75(8ii?3)UXAv+qM2tIxRC(e@<2jUL{}_$b=J{e2T#Y>nay zdBY5E;d~$2Nxrwy?vpeB0klirImCWbm;TAIne7oD0o#u~hPjSoTQhtM?|?lSe>%6F z)A?FxdoCgR`8mk)$%lN;u-$*0`)XceUEYU3VajN0c`y83uphDC&vzfNXSctA_4^)! zU+lXq;@JXrzwa;j#lF9QW9-}i!@fWLXXr=$BI4bAO4s5IIp_U5%S84sr)}@vx4)Qi z3))%Le~@b%dOPCM?|as-wwCAX8#9N$6yC22;+^!~g_zyg734Q!_MrD7?vwwUJmRg0 z^~|pw_P&E1@_R4dMZ7EW#(1y%XC>bT??WHTIp_0Ri0#jNbQF>E{kh*ebqtYzPK?-l zE9V`1NQ`ryK;+Ho_Y|AcyY>O%{R=sbKa=yu`@O}U;aRjZsy*X#=qbdde;Tb{9nRn* z>~W59FEe0!24%vwr3ao*Rka;^XtK6&LJt$@TZl38foovjLd+X_pU0LDy*IJ#N8a9VqUA!J1?)|Ty-uRrh;K*U7~f8U z$csIXJBVj6&z=_gE@F(lHE!ozsL=%%YxKa@kPkKb*z%#q0NY+ejqk9Hk+((%Ef;DG zv5PfEU~9;S8Y#AXsIiFc9Af9VhnTksHden}sBs^=SmOcM8uFpW61IG(@jbS4gc?6! zdp_1MR==EYrH1zIs{I}t(GQVT$Yvz;@ng;h{1bQ(B|8Q{)hj_Cy;*uv{WAn literal 0 HcmV?d00001 diff --git a/test_conformance/vulkan/shaders/image2D_rg32i.spv b/test_conformance/vulkan/shaders/image2D_rg32i.spv new file mode 100644 index 0000000000000000000000000000000000000000..b7d302f4a70b50c0cae567bade858cec703577d9 GIT binary patch literal 3264 zcmZ9NiE>m`5QeXigiQow6GaS*pn|Aj5k;27#32#^0}8IgWagS&GMh6uKok`b^>us( zAIegztn&M2Zd^mq&h0ZvPBZdE@xWQ?I+@sbsBc;MXL+^O%krgG znrHQDzS>Ru)vTSLu=f#8$ezm8tChK#YQ5X7_Vd+dT5V?SwA#peWX|rU#;q^KW$k)% zrBMzTC8B(`QBJ#o{Oy(2jip{%YtY{)ci@>9o7F2Tt(%?Ra;4W<>Eh8hnmmopTD8S= zVtTf5Uv4CM7Va9yZ4os_MzM>$v%}Eov|sOK-MrH?(I|JecMvhmc^G?;b2@8DTjywS z6Vr29KOb;@&!se5T*{|by%+PGLH9cMu-6a16Ma{jorPL+wtcr#ujN^%Jv&{*&Y*iazuDX7sr~@@&U;?JMqS*y~90 z9Gv}oclIFm%yFjYvBwc>noGP7aarpHw6W&ffpq}w*$<(O)o0v`XnT_1!^|E*>o;$C z8SfoO-!EamiR>lE%V_t>ne!F2OWxkaen*%7SBK_ckJuTu2kXXJPGMU!?%(fjPsYDp z*q+TZ)zR+tP3v!<%_krE9%H-zG&5>mV_nX`pD?3nYdH)47Pw#RH}kCr_P!36uzufT z_{F}(GM*`6_xt9;FZRs^w)gaZc)7sm=*PGx5$EwSU5mFCoHKWx6WPC`nQwY*=yxjFoGaHhr)=g?V;nElw}$Zy0PKp#ZhC;vBj#5)n| znO{5h9>b3O&c$)WS&=v92>LZdehGXC?LRCz&*!%g+usEG1S02~bH8(O5|RIv81d&mB#2ESDZez=vH+rvO-$tVM1#J6~xA$9Uxv29McKPirf_*#k#`t!Y5P7lpks;o} z0=?DIcMxObt&tX7)M$dsHCkY6$VZJfwtUp+U=Mp8*64zbk+;S&S}tmQja{zM16xBr zYV@(?qehNBUtrI%f|$3?IgQmX7d7r;muuVuTSGo-+{c!W8mrizLu`$25bwtt#_E^z ztz1SsyJ}}=8~R&h1F{{7J{}Z&;P1fB-@{^7zQ>j`&b!hlSAMU)6Y)i~HPun`2W)G` zygGjm5&5Y3Bla02YCgi2dx%(5pPV($p{;SDz`F~qj(k61d$*DAXKcC0i23x%nQsT$ Ne7n&9@E`dM@-I{S9j*WX literal 0 HcmV?d00001 diff --git a/test_conformance/vulkan/shaders/image2D_rg32ui.spv b/test_conformance/vulkan/shaders/image2D_rg32ui.spv new file mode 100644 index 0000000000000000000000000000000000000000..6cf2f1b8ff79b915b32ea0c3763689988fd4eaa3 GIT binary patch literal 3264 zcmZ9NiE>m`5QeXigiQow6GaS*pn|Aj5k;27#32#^0}8IgWagS&GMh6uKok`b^>us( zAIegztn&M2Zd^mq&h0ZvPBZdE@xWQ?I+@sbsBc;MXL+^O%krgG znrHQDzS>Ru)vTSLu=f#8$ezm8tChK#YQ5X7_Vd+dT5V?SwA#peWX|rU#;q^KW$k)% zrBMzTC8B(`QBJ#o{Oy(2jip{%YtY{)ci@>9o7F2Tt(%?Ra;4W<>Eh8hnmmopTD8S= zVtTf5Uv4CM7Va9yZ4os_MzM>$v%}Eov|sOK-MrH?(I|JecMvhmc^G?;b2@8DTjywS z6Vr29KOb;@&!se5T*{|by%+PGLH9cMu-6a16Ma{jorPL+wtcr#ujN^%Jv&{*&Y*iazuDX7sr~@@&U;?JMqS*y~90 z9Gv}oclIFm%yFjYvBwc>noGP7aarpHw6W&ffpq}w*$<(O)o0v`XnT_1!^|E*yT`oc zWxU@s`hE%fO=K@QUPil5&YZ8HUGnxW_B*=tzdAGrd&JJLJ=j;A9dnsgBn6P3v!<%_krE9%H-zG&5>mV_nX`pD?3nYdH)47Pw#RH}kCr_P!36uzufT z_{F}(GM*`6_xt9;FZRs^w)gaZc)7sm=*PGx5$EwSU5mFCoHKWx6WPC`nQwY*c`l{x%%aHhr)=g?V;nElw}$Zy0PKszVylmDAM;+=@~ z%Lk6}lC=i)fxtjHU41pOK!zXU#n_8*p<=kr^L?Qa5o0+I90x!;*OiOBy-jQDiH zIb#oq@tn@GygB`zVskocXA$QwavJ|m!5i=Q7JG;1(4JB49bZ6SL|pnO(fZYK2UFPQ z2>dR#J;Z%ZV;dvyeTvT`<~Emh^!gsQz52}~Ujd5s0wmIUyuV5cWqMrAA6_Jm-y;g8>o)5w1k@q~}S;YF*(~f%Av8`t=bKgMZ!v7Jr zcNhMTvE{Du?VOdHhEf+Pu#xB?Bfvq7Q zHTu}{Q6tBmFR$0ojg39}fyX@ONP6?_n`3-($-e=UwTOE5BFYiTEPgn(CX)B`4=;29f|+| literal 0 HcmV?d00001 diff --git a/test_conformance/vulkan/shaders/image2D_rg8i.spv b/test_conformance/vulkan/shaders/image2D_rg8i.spv new file mode 100644 index 0000000000000000000000000000000000000000..a71b9bf08a654f4a68efb190656be3d91cefa7ce GIT binary patch literal 3264 zcmZ9NiE>m`5QeXigiQow6GaS*pn|Aj5k;27#32#^0}8IgWagS&GMh6uKok`b^>us( zAIegztn&M2Zd^mq&h0ZvPBZdE@xWQ?I+@sbsBc;MXL+^O%krgG znrHQDzS>Ru)vTSLu=f#8$ezm8tChK#YQ5X7_Vd+dT5V?SwA#peWX|rU#;q^KW$k)% zrBMzTC8B(`QBJ#o{Oy(2jip{%YtY{)ci@>9o7F2Tt(%?Ra;4W<>Eh8hnmmopTD8S= zVtTf5Uv4CM7Va9yZ4os_MzM>$v%}Eov|sOK-MrH?(I|JecMvhmc^G?;b2@8DTjywS z6Vr29KOb;@&!se5T*{|by%+PGLH9cMu-6a16Ma{jorPL+wtcr#ujN^%Jv&{*&Y*iazuDX7sr~@@&U;?JMqS*y~90 z9Gv}oclIFm%yFjYvBwc>noGP7aarpHw6W&ffpq}w*$<(O)o0v`XnT_1!^|E*>o;$C z8E+5K_e_W6E$`QNWe$HGoT+idIds+{Ww|bDK*$dVLSuUj62guYxUaug;>_UM`{CueR6s(U%dI{yDUM zb>#X0+Z=J5*w!&F-AVP+t~8vjoxe6w~^?50o#7$?fn*7F6w-RU4A=@VBe0sF}|H8L|*KDWQcdL zKyP*Q9mE)UYorAiHJadZjTYD%@=>FWEgv;H*u!3jHM(G92?L_TW%hm`5QeXigiQow6GaS*pn|Aj5k;27#32#^0}8IgWagS&GMh6uKok`b^>us( zAIegztn&M2Zd^mq&h0ZvPBZdE@xWQ?I+@sbsBc;MXL+^O%krgG znrHQDzS>Ru)vTSLu=f#8$ezm8tChK#YQ5X7_Vd+dT5V?SwA#peWX|rU#;q^KW$k)% zrBMzTC8B(`QBJ#o{Oy(2jip{%YtY{)ci@>9o7F2Tt(%?Ra;4W<>Eh8hnmmopTD8S= zVtTf5Uv4CM7Va9yZ4os_MzM>$v%}Eov|sOK-MrH?(I|JecMvhmc^G?;b2@8DTjywS z6Vr29KOb;@&!se5T*{|by%+PGLH9cMu-6a16Ma{jorPL+wtcr#ujN^%Jv&{*&Y*iazuDX7sr~@@&U;?JMqS*y~90 z9Gv}oclIFm%yFjYvBwc>noGP7aarpHw6W&ffpq}w*$<(O)o0v`XnT_1!^|E*yT`oc zWxTP`_ev9JEgc(I!%USTZ!2M#snQuL?cW}6b_4^*f zFZL~#@k|N3-!~V2v2QN0y{G@f%LP71KgKMj%;B$tGc}Glht68W?8hEQek0}p`XJ&y`M=2{-icVx z{Mxbi7UqCc5&5{=YXukQ`4DU#dCw!BMXYZYewx>t* zFov!o#>fYE8(ZGI(R&U1HWIxrVB3$pz28F1MV+s(%Wr29?Awty#<#PC$cw#?4Dk*Y z=&g>vgBT-kjkMsRMiX4F(E?jTK5DeF<)cOid)Vu+Mi*?1yfv25a#7=J>~f7B*c$Rt zqmL~gHFE6v0(*`X#JqLRX{>&^sBss&T;m?t8uC%&KDK<+SjF}nVrzVZct6%KR==EY z2-jzPN@_Y52h%ch8sg9aIU|Tch z)%kmf$Vbf|vCkk;^AWb(L&Tc;m`5QeXigiQow6GaS*pn|Aj5k;27#32#^0}8IgWagS&GMh6uKok`b^>us( zAIegztn&M2Zd^mq&h0ZTNzbL&R9()-BRJG z=)~2tOkB((0b-^|CR#pL2qXQox%N7v!3=&V#d(ph-){PM>61y83(UN z%PHcGXzPj37TB6{N5SUXjEtA?Q`q{C;dP&V8KdYkuWv@5`y^G6UBj00e_n&4)&1Plxxn_G{tqt~_#FKh_ax#xKBjB&)`D~9&T}IBm(#X)XY4P|xC8B3 z)qjx7@3|Xs>GuumS6j>b^e^i%w_Hyh+OzT!uIaM z|1q}QHNKs*auacnyyu@UxY+**xSan}uw0z)Gi=Wn-`*ZRN95ytUld&A{1R+WkLY0x zT}6zM5AHU$ym_Pd8uo1@dSAe{A9;Jfg_estUtyQu&LY^iBX5jvX988a=QzzvbA{c=&`E_S)bJ+L+8qsDz~`KYmq?K#BO_y+NQtYNHvIp4}< zw6m*rcDA9vMK&PYk?7+=!3X{h?EF0}X61WqIpe%5eRAdZ>N^o%L|aoGHGjaiX3VSe z_Yjefnm=NnL89g(Y`KSsHTB6^;~d%=7Ye+)!0O2N6Sj97`F_ThdyJS*pPc!2pv|`n L{SW_<&mjK-*xwyz literal 0 HcmV?d00001 diff --git a/test_conformance/vulkan/shaders/image2D_rgba16ui.spv b/test_conformance/vulkan/shaders/image2D_rgba16ui.spv new file mode 100644 index 0000000000000000000000000000000000000000..84c3d3db7a6147478654fed0abb9f462eb5831c4 GIT binary patch literal 3256 zcmZ9NiE>m`5QeXigiQow6GaS*pn|Aj5k;27#32#^0}8IgWagS&GMh6uKok`b^>us( zAIegztn&M2Zd^mq&h0ZTNzbL&R9()-BRJG z=)~2tOkB((0b-^|CR#pL2qXQox%N7v!3=&V#d(ph-){PM>61y83(UN z%PHcGXzPj37TB6{N5SUXjEtA?Q`q{C;dP&V8KdYkuWv@5`yjdUSnO(z@IRqXlpqO{ua1j>^JkR2ll=Wm#}`{WBA3s z#WJ2LVfXvy!Y}sC1-AF}e|WjT=jg|{ClTlIFoVL9?V}EhR9ca(0 z{)1e*(7O?re&4WuwY9uc-<3K1b@2O-BhI0-7BTy=$C2NNIe^+7Z`JId7h_fPZ%n|f!i2M@x5ZZrOa-Pp`A-2B>^a(`HH|Kt5>LeonD>35J1?P-C zAjWe#%kt**dy38Jter)izsPC)I|XmN-&^b*oRn`?^Q%T?)F;2#d$sin@8UBh-VS&TTeUcUB|Ybxy*e7kqiGv*xp_E zKgO23#VxTw(tmus}Z){u`HZEX3d(ZL?}I;_zJ8zXOxWwcz>_!_%hqX)KzeAMV; z%SVkId%nP)V+Ap9opTziUoL9g#V*&l2eyWM)VPl=A2n97J%`vD-yq(PHH_6S=Ucgq zc6Qay&NlS7$OdFP5`8=<_`u(Roxg|0tbC6xXPkGXPp-_O``j}i0flQZ8AwE1?S K|KUIK8RTED^c_L~ literal 0 HcmV?d00001 diff --git a/test_conformance/vulkan/shaders/image2D_rgba32f.spv b/test_conformance/vulkan/shaders/image2D_rgba32f.spv new file mode 100644 index 0000000000000000000000000000000000000000..35136c581b75c00ea5c883aea3a67fdbbc0204ac GIT binary patch literal 3268 zcmZ9NiINjl6oy+!hE4WO6frE~0*=Zii{Q*SOc;%TfC6qIlQfftY$iRxC@Lbz>-Y>l zl%-Z#<@a^^T3z0$Q|J8uIrl8L@6D`Sv42$*Jr}Kw)<@5x@!A}%fQcgY5%jufBE!|% z)!O7JZA~6K{*I9wqLr%5w=r53m9a}mw~_SZ*CD5oP3)K7j8t+l|4RN=pnnIHqBT*q zR-e8yQy;~H#;}pb^-;Ugiih?7!kxI8j*MAv%=sJFf-8P7=p@Y~tv7~A+U~|_(yXV; zgLqU=dg)PnFEJT=s+TTRYv=3D!Js}$mpgI2ll05L=3_W_Ewr z8ck>T=6`r5!?(e$M%tK08{f(C47v*TjFcCU3q*vO_k&$kdXw{Z8SlLN5Ldu^$1uar zTl9-@v8B+nTM$0MUwP@>! z&*a#ea)-d?TaQ!qF zOfO?s5Nn!Cya#bv>lL)I=G%s~5AEy+(8lUB?p3rs$?u?t*D^kec5#1S#}-?oxZE4+ zn*R;tJ!Cie-bA}k&iseaE_vq=`%PW?C&y;CM|>1)KlT{rI)QD?@GZOp);Io4Zrg|R zwb1rlLiF=|ajX-o0;sG2=G0 zv#S3f*ADbf#HHW&tY2*{&(}9*4u2`UUlqhV>Aedvd$B9XZ^Z0F??>Dx|2KKW+Y#%T zUpwr53p?cZUc8NXSLBWHUir^Tz70NremCcw&u<~NKkv~oM9%l;e(%(AME)5uV(+b- zckCfC&Uq4%H>clIY)UOz_H5SRW> z(E8OO*F|h|g!8_HeGm!tJm1TRd^p=rb1uwt1#BL9=MjH|Sl@cuq24uY>zT{kpCNL= z|2ej27yQ?;yel^l_sBc{)tn3azW^8We+iZg^L>TweBs;M!`Fy>nD3jM3pwjx zdwNU{6X;urG4g?HVBbQ_8+u>BmJhu*vF%6R-fyDiLY+nIEr`8NqT7gXN8T9UPJ+mb zJ&!wxXRyGY7Wyt?jJ!2&=Uk}K1s7}dz}AotHTu}{p~e8)UPFycn*1#-<6-R<&5*J^vM<9tM5eY?}|0m zq2?oOYkIFjji0d}74T#1QyDkDpTA(s8D~v>a@IJ9w#HPBcjj0f@;|}$Y{Lw{V#_^6 V%%@MzeB03G+kyUv|Hvnie*q<59>4$q literal 0 HcmV?d00001 diff --git a/test_conformance/vulkan/shaders/image2D_rgba32i.spv b/test_conformance/vulkan/shaders/image2D_rgba32i.spv new file mode 100644 index 0000000000000000000000000000000000000000..4d1ae58107ad1e187b2a0c8a0048026b900523ce GIT binary patch literal 3256 zcmZ9NiE>m`5QeXiu*xQzC}LOy6+{h-D6%9b4v`2LP;eb4GuPyj*_^onqNs?duj4cL zP?lO{mESjWTW@(%UETlReY*Rc+h?+R)xot%@<#m$mE7*k%FiAK4zy@QBh&coP)oKsm#+B!#r8=soZ z`uTu+9d0qr78di#74OA7XVAUQJ?wRZ??m5~W@ob0y+(gxfN%YaCkL48Od7SkHi0&Nslb!y3fMbRo<=Sa5obOKc2($2&gnAV^BzE4 zf%gyL0DIoDU#y?88u*<->zT*>TME1fy`2ek2KQ6VdfGdQS%V%!T)V+Mk^yJT7Sqd-g+UWAz#L0@|MB_b{_Z(E80=Uc=8S|_F&yO%PDMY#{K)`%_P1MOMW zzmsbhdN<Hy~W<)Ikaa~d&d{h7ZI2K3ABE7+`%NaIRd|f zZ4YsuQ`pAHd!OR-h`G(B9lgGbZLfaw$XCIZw^wITY%iD4?pNFEd+5uEOaCldzdCZg zk8O^)?05CgL7>&p%givHxRmIsYeMxj5ga*q$%Gy*+$}$jA9UFSyA01=yY*(Zd>a z6){FWxZBwB=8fKK*te1BeIDC>3t-=lyfMC=MMPfgePoDtFi&rF z^c}<)d26Hv7d4vTa*Y<)8uC%2jV&KFI@rTrhc&ujW8|%|gqDjMUtyPP^uX4Tj~ab! z`KXa&&lT8nEFm`5QeXigiQow6GaS*pn|Aj5k;27#32#^0}8IgWagS&GMh6uKok`b^>us( zAIegztn&M2Zd^mq&h0ZTNzbL&R9()-BRJG z=)~2tOkB((0b-^|CR#pL2qXQox%N7v!3=&V#d(ph-){PM>61y83(UN z%PHcGXzPj37TB6{N5SUXjEtA?Q`q{C;dP&V8KdYkuWv@5`yjdUSnO(z@IRqXlpqO{ua1j>^JkR2ll=Wm#}`{WBA3s z#WJ2LVfXvy!Y}sC1-AF}e|WjT=jg|{ClTlIFoVL9?V}EhR9ca(0 z{)1e*(7O?re&4WuwY9uc-<3K1b#SJ}5$DiZh?xD@zQ9W z_8!BI{LaO3#95Iy<_P*VM1Bc;2<<;CInU>}5Zm7b`UE29n{&T2brO;Pl^F5qf^)_m z5aT(WWqEV@J;mm9*3Kf%U*t6Yoq{*s?=AKY&!Ih|+B?30zKFQ=Ponj!;|`{<%@O!r zY9tG`qh!^ z18j4|eP6*oj6^-}_bMVEcYCei;yfRM%_Hx5#IuO?t*0IJu47xzT;{%k$c6tSZ0|1o zA7jg1Tu|Vc$lg_XTYGk+=6-Xt}8K6?XaUEP{PI^2YdfmJoTd_mLsq!2-S2 z(RUDI&16xBrYTU<`j~c7ko` literal 0 HcmV?d00001 diff --git a/test_conformance/vulkan/shaders/image2D_rgba8i.spv b/test_conformance/vulkan/shaders/image2D_rgba8i.spv new file mode 100644 index 0000000000000000000000000000000000000000..edf8c58ce7cac895e99bc4c63cd0960c96e7f944 GIT binary patch literal 3256 zcmZ9NiE>m`5QeXigiQow6GaS*pn|Aj5k;27#32#^0}8IgWagS&GMh6uKok`b^>us( zAIegztn&M2Zd^mq&h0ZTNzbL&R9()-BRJG z=)~2tOkB((0b-^|CR#pL2qXQox%N7v!3=&V#d(ph-){PM>61y83(UN z%PHcGXzPj37TB6{N5SUXjEtA?Q`q{C;dP&V8KdYkuWv@5`y^G6UBj00e_n&4)&1Plxxn_G{tqt~_#FKh_ax#xKBjB&)`D~9&T}IBm(#X)XY4P|xC8B3 z)qjv{7us*?((fDAueO%=>$@_CzYfmSIN}^SYZ0>_dmQm8i2LOKCXaY0Vm9yEjyNmw#vDPvhR830525{sCFlA47GnFGK%YS5d~@!1E>0ryzY-%pU2x9W z17bX-*@-h)e$*TE9AS zeSmF_xbG|2hmok~{a!`n<8H4NT%6}auzBP?k9ZcbzV)=D-gRv2nakWa5V`Pwgzeph z|6^>qYkWIryOPz9_iJ`6bw%9?`=X zx{4ShAKYzhdGki^HSF6+^uB;?Kl1i|3oRFQzQQiQokg&3N8T9U&JrRo_C7MiJ6NE% zI{FS`jJ!3{f{PkWaJfbcYz_IS(Z-gK8XfFmufrN$urc!1SVqf5jjyrGHF{uc$VZJn zwtUpcvF8iyIaUz!);XuK`sJd=UF>p=dthtGM~(Z~@=;?I+jEGm@eShrSi@NTa=w+z zXlGaL>}*4Si)=u)Bhkl$f)D&1*!g={%*yxJa>jXA`sB**)psJkh_m`5QeXiu*xQzC}LOy6+{h-D6%9b4v`2LP;eb4GuPyj*_^onqNs?duj4cL zP?lO{mESjWTW@(%UETlReY*Rc+h?+R)xot%@<#m$mE7*k%FiAK4zy@QBh&coP)oKsm#+B!#r8=soZ z`uTu+9d0qr78di#74OA7XVAUQJ?wRZ??m5~W@ob0y+(gxfN%YaCkL48Od7SkHi0&Nslb!y3fMbRo<=Sa5obOKc2($2&gnAV^BzE4 zf%gyL0DIoDU#y?88u*<->zT*>TME1fy`2ek2KQ6VdfGdQS%V%!T)V+Mk^yJT7Sqd-g+UWAz#L0@|MB_b{_Z(C#sBc@5+L z>-$BVZyvBiuFT=DgEKXTIET(!#O%i&Lw+UZ0Qw-}KKZ}MBi@Nv&-~i4 z_ZW8McP@@2&WgM-N6@b#@{8a@X#Zi!c|N~|*#5@RClEQ`oco=rlZgB;#E4H9oHO=- z7|-b}%bU~hDK@9Gb{28|BB$|h7rgO)Z?SiH4(%D$-th(WMZ~3l0Q_q?bTTn+sh@i`_=aP9{Muk(m#vVuZ~>r zW1A!H`wI49B9ou^5GWQKcF8m*2dw1dg z2wUzN-_BXNiMU7J^UoDr?Ee^C&i@HmF3$HUw&#m)Zx5d#@^QY;3odeg0k)?{^sok9 zMU0UT?l!i(d879l_H873pU1Wzd3(QwmWw)HVwc~}0@$}BZ;Wqe5s?>r9~t5u%+p&P zeFrf{-WqAaMU5u7T%!fHhJ4g$W6MX44)(CuVT~@>7qmJ+L+8qedTF zK5FFHa|QMs%ZPdFoYPqSa#7pYl1zsI&_%&YVF z5Rs3XKVY9hqUIxPxrc}~^~qV|9NHQe3cS0(>d5yaws#x(e!`Y}jF?ZKocVU3&9@8v KH~*2(ApZckOdUf2 literal 0 HcmV?d00001 diff --git a/test_conformance/vulkan/test_vulkan_api_consistency.cpp b/test_conformance/vulkan/test_vulkan_api_consistency.cpp index 2987418f..f22ac319 100644 --- a/test_conformance/vulkan/test_vulkan_api_consistency.cpp +++ b/test_conformance/vulkan/test_vulkan_api_consistency.cpp @@ -238,7 +238,7 @@ int test_consistency_external_image(cl_device_id deviceID, cl_context _context, const VulkanMemoryTypeList& memoryTypeList = vkImage2D->getMemoryTypeList(); uint64_t totalImageMemSize = vkImage2D->getSize(); - log_info("Memory type index: %d\n", (uint32_t)memoryTypeList[0]); + log_info("Memory type index: %lu\n", (uint32_t)memoryTypeList[0]); log_info("Memory type property: %d\n", memoryTypeList[0].getMemoryTypeProperty()); log_info("Image size : %d\n", totalImageMemSize); @@ -552,17 +552,17 @@ int test_consistency_external_semaphore(cl_device_id deviceID, // Pass invalid object to release call - errNum = clReleaseSemaphoreObjectKHRptr(NULL); + errNum = clReleaseSemaphoreKHRptr(NULL); test_failure_error(errNum, CL_INVALID_VALUE, - "clReleaseSemaphoreObjectKHRptr fails with " + "clReleaseSemaphoreKHRptr fails with " "CL_INVALID_VALUE when NULL semaphore object is passed"); // Release both semaphore objects - errNum = clReleaseSemaphoreObjectKHRptr(clVk2Clsemaphore); - test_error(errNum, "clReleaseSemaphoreObjectKHRptr failed"); + errNum = clReleaseSemaphoreKHRptr(clVk2Clsemaphore); + test_error(errNum, "clReleaseSemaphoreKHRptr failed"); - errNum = clReleaseSemaphoreObjectKHRptr(clCl2Vksemaphore); - test_error(errNum, "clReleaseSemaphoreObjectKHRptr failed"); + errNum = clReleaseSemaphoreKHRptr(clCl2Vksemaphore); + test_error(errNum, "clReleaseSemaphoreKHRptr failed"); return TEST_PASS; } diff --git a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp index 7daf96de..9b0bc9de 100644 --- a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp +++ b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp @@ -39,35 +39,6 @@ struct Params }; } -static const char *vkBufferShader = - "#version 450\n" - "#extension GL_ARB_separate_shader_objects : enable\n" - "#extension GL_NV_gpu_shader5 : enable\n" - "layout(binding = 0) buffer Params\n" - "{\n" - " uint32_t numBuffers;\n" - " uint32_t bufferSize;\n" - " uint32_t interBufferOffset;\n" - "};\n" - "layout(binding = 1) buffer Buffer\n" - "{\n" - " uint8_t ptr[];\n" - "} bufferPtrList[" STRING( - MAX_BUFFERS) "];\n" - "layout(local_size_x = 512) in;\n" - "void main() {\n" - " for (uint32_t bufIdx = 0; bufIdx < numBuffers;" - " bufIdx++) {\n" - " uint32_t ptrIdx = gl_GlobalInvocationID.x;\n" - " uint32_t limit = bufferSize;\n" - " while (ptrIdx < limit) {\n" - " bufferPtrList[bufIdx].ptr[ptrIdx]++;\n" - " ptrIdx += (gl_NumWorkGroups.x * " - "gl_WorkGroupSize.x);\n" - " }\n" - " }\n" - "}\n"; - const char *kernel_text_numbuffer_1 = " \ __kernel void clUpdateBuffer(int bufferSize, __global unsigned char *a) { \n\ int gid = get_global_id(0); \n\ @@ -149,6 +120,8 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, VulkanQueue &vkQueue = vkDevice.getQueue(); + std::vector vkBufferShader = readFile("buffer.spv"); + VulkanShaderModule vkBufferShaderModule(vkDevice, vkBufferShader); VulkanDescriptorSetLayoutBindingList vkDescriptorSetLayoutBindingList( MAX_BUFFERS + 1, VULKAN_DESCRIPTOR_TYPE_STORAGE_BUFFER); @@ -446,6 +419,7 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, VulkanQueue &vkQueue = vkDevice.getQueue(); + std::vector vkBufferShader = readFile("buffer.spv"); VulkanShaderModule vkBufferShaderModule(vkDevice, vkBufferShader); VulkanDescriptorSetLayoutBindingList vkDescriptorSetLayoutBindingList( MAX_BUFFERS + 1, VULKAN_DESCRIPTOR_TYPE_STORAGE_BUFFER); @@ -716,6 +690,8 @@ int run_test_with_multi_import_same_ctx( VulkanQueue &vkQueue = vkDevice.getQueue(); + std::vector vkBufferShader = readFile("buffer.spv"); + VulkanShaderModule vkBufferShaderModule(vkDevice, vkBufferShader); VulkanDescriptorSetLayoutBindingList vkDescriptorSetLayoutBindingList( MAX_BUFFERS + 1, VULKAN_DESCRIPTOR_TYPE_STORAGE_BUFFER); @@ -1050,6 +1026,8 @@ int run_test_with_multi_import_diff_ctx( VulkanQueue &vkQueue = vkDevice.getQueue(); + std::vector vkBufferShader = readFile("buffer.spv"); + VulkanShaderModule vkBufferShaderModule(vkDevice, vkBufferShader); VulkanDescriptorSetLayoutBindingList vkDescriptorSetLayoutBindingList( MAX_BUFFERS + 1, VULKAN_DESCRIPTOR_TYPE_STORAGE_BUFFER); diff --git a/test_conformance/vulkan/test_vulkan_interop_image.cpp b/test_conformance/vulkan/test_vulkan_interop_image.cpp index f1d0af1f..7577de09 100644 --- a/test_conformance/vulkan/test_vulkan_interop_image.cpp +++ b/test_conformance/vulkan/test_vulkan_interop_image.cpp @@ -25,8 +25,6 @@ #define MAX_2D_IMAGE_ELEMENT_SIZE 16 #define MAX_2D_IMAGE_MIP_LEVELS 11 #define MAX_2D_IMAGE_DESCRIPTORS MAX_2D_IMAGES *MAX_2D_IMAGE_MIP_LEVELS -#define GLSL_FORMAT_STRING "" -#define GLSL_TYPE_PREFIX_STRING "" #define NUM_THREADS_PER_GROUP_X 32 #define NUM_THREADS_PER_GROUP_Y 32 #define NUM_BLOCKS(size, blockSize) \ @@ -54,61 +52,8 @@ struct Params } static cl_uchar uuid[CL_UUID_SIZE_KHR]; static cl_device_id deviceId = NULL; - -static const char *vkImage2DShader = - "#version 450\n" - "#extension GL_ARB_separate_shader_objects : enable\n" - "#extension GL_NV_gpu_shader5 : enable\n" - "layout(binding = 0) buffer Params\n" - "{\n" - " uint32_t numImage2DDescriptors;\n" - "};\n" - "layout(binding = 1, " GLSL_FORMAT_STRING - ") uniform " GLSL_TYPE_PREFIX_STRING "image2D image2DList[" STRING( - MAX_2D_IMAGE_DESCRIPTORS) "];\n" - "layout(local_size_x = 32, local_size_y = " - "32) in;\n" - "void main() {\n" - " uvec3 numThreads = gl_NumWorkGroups * " - "gl_WorkGroupSize;\n" - " for (uint32_t image2DIdx = 0; " - "image2DIdx < numImage2DDescriptors; " - "image2DIdx++)" - " {\n" - " ivec2 imageDim = " - "imageSize(image2DList[image2DIdx]);\n" - " uint32_t heightBy2 = imageDim.y / " - "2;\n" - " for (uint32_t row = " - "gl_GlobalInvocationID.y; row < heightBy2; " - "row += numThreads.y)" - " {\n" - " for (uint32_t col = " - "gl_GlobalInvocationID.x; col < imageDim.x; " - "col += numThreads.x)" - " {\n" - " ivec2 coordsA = ivec2(col, " - "row);\n" - " ivec2 coordsB = ivec2(col, " - "imageDim.y - row - 1);\n" - " " GLSL_TYPE_PREFIX_STRING - "vec4 dataA = " - "imageLoad(image2DList[image2DIdx], " - "coordsA);\n" - " " GLSL_TYPE_PREFIX_STRING - "vec4 dataB = " - "imageLoad(image2DList[image2DIdx], " - "coordsB);\n" - " " - "imageStore(image2DList[image2DIdx], " - "coordsA, dataB);\n" - " " - "imageStore(image2DList[image2DIdx], " - "coordsB, dataA);\n" - " }\n" - " }\n" - " }\n" - "}\n"; +size_t max_width = MAX_2D_IMAGE_WIDTH; +size_t max_height = MAX_2D_IMAGE_HEIGHT; const char *kernel_text_numImage_1 = " \ __constant sampler_t smpImg = CLK_NORMALIZED_COORDS_FALSE|CLK_ADDRESS_NONE|CLK_FILTER_NEAREST;\n\ @@ -268,8 +213,8 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, VULKAN_MEMORY_TYPE_PROPERTY_HOST_VISIBLE_COHERENT)); vkParamsDeviceMemory.bindBuffer(vkParamsBuffer); - uint64_t maxImage2DSize = MAX_2D_IMAGE_WIDTH * MAX_2D_IMAGE_HEIGHT - * MAX_2D_IMAGE_ELEMENT_SIZE * 2; + uint64_t maxImage2DSize = + max_width * max_height * MAX_2D_IMAGE_ELEMENT_SIZE * 2; VulkanBuffer vkSrcBuffer(vkDevice, maxImage2DSize); VulkanDeviceMemory vkSrcBufferDeviceMemory( vkDevice, vkSrcBuffer.getSize(), @@ -310,6 +255,12 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, clCl2VkExternalSemaphore = new clExternalSemaphore( vkCl2VkSemaphore, context, vkExternalSemaphoreHandleType, deviceId); + std::vector vkNonDedicatedImage2DListDeviceMemory1; + std::vector vkNonDedicatedImage2DListDeviceMemory2; + std::vector nonDedicatedExternalMemory1; + std::vector nonDedicatedExternalMemory2; + std::vector vkImage2DShader; + for (size_t fIdx = 0; fIdx < vkFormatList.size(); fIdx++) { VulkanFormat vkFormat = vkFormatList[fIdx]; @@ -317,15 +268,13 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, uint32_t elementSize = getVulkanFormatElementSize(vkFormat); ASSERT_LEQ(elementSize, (uint32_t)MAX_2D_IMAGE_ELEMENT_SIZE); log_info("elementSize= %d\n", elementSize); - std::map patternToSubstituteMap; - patternToSubstituteMap[GLSL_FORMAT_STRING] = - getVulkanFormatGLSLFormat(vkFormat); - patternToSubstituteMap[GLSL_TYPE_PREFIX_STRING] = - getVulkanFormatGLSLTypePrefix(vkFormat); - VulkanShaderModule vkImage2DShaderModule( - vkDevice, - prepareVulkanShader(vkImage2DShader, patternToSubstituteMap)); + std::string fileName = "image2D_" + + std::string(getVulkanFormatGLSLFormat(vkFormat)) + ".spv"; + log_info("Load %s file", fileName.c_str()); + vkImage2DShader = readFile(fileName); + VulkanShaderModule vkImage2DShaderModule(vkDevice, vkImage2DShader); + VulkanComputePipeline vkComputePipeline(vkDevice, vkPipelineLayout, vkImage2DShaderModule); @@ -333,13 +282,13 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, { uint32_t width = widthList[wIdx]; log_info("Width: %d\n", width); - ASSERT_LEQ(width, (uint32_t)MAX_2D_IMAGE_WIDTH); + if (width > max_width) continue; region[0] = width; for (size_t hIdx = 0; hIdx < ARRAY_SIZE(heightList); hIdx++) { uint32_t height = heightList[hIdx]; log_info("Height: %d", height); - ASSERT_LEQ(height, (uint32_t)MAX_2D_IMAGE_HEIGHT); + if (height > max_height) continue; region[1] = height; uint32_t numMipLevels = 1; @@ -418,14 +367,6 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, const VulkanMemoryTypeList &memoryTypeList = vkDummyImage2D.getMemoryTypeList(); - std::vector - vkNonDedicatedImage2DListDeviceMemory1; - std::vector - vkNonDedicatedImage2DListDeviceMemory2; - std::vector - nonDedicatedExternalMemory1; - std::vector - nonDedicatedExternalMemory2; for (size_t mtIdx = 0; mtIdx < memoryTypeList.size(); mtIdx++) { @@ -834,6 +775,8 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, } } } + + vkImage2DShader.clear(); } CLEANUP: if (clVk2CLExternalSemaphore) delete clVk2CLExternalSemaphore; @@ -866,8 +809,8 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, VULKAN_MEMORY_TYPE_PROPERTY_HOST_VISIBLE_COHERENT)); vkParamsDeviceMemory.bindBuffer(vkParamsBuffer); - uint64_t maxImage2DSize = MAX_2D_IMAGE_WIDTH * MAX_2D_IMAGE_HEIGHT - * MAX_2D_IMAGE_ELEMENT_SIZE * 2; + uint64_t maxImage2DSize = + max_width * max_height * MAX_2D_IMAGE_ELEMENT_SIZE * 2; VulkanBuffer vkSrcBuffer(vkDevice, maxImage2DSize); VulkanDeviceMemory vkSrcBufferDeviceMemory( vkDevice, vkSrcBuffer.getSize(), @@ -908,6 +851,12 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, clCl2VkExternalSemaphore = new clExternalSemaphore( vkCl2VkSemaphore, context, vkExternalSemaphoreHandleType, deviceId); + std::vector vkNonDedicatedImage2DListDeviceMemory1; + std::vector vkNonDedicatedImage2DListDeviceMemory2; + std::vector nonDedicatedExternalMemory1; + std::vector nonDedicatedExternalMemory2; + std::vector vkImage2DShader; + for (size_t fIdx = 0; fIdx < vkFormatList.size(); fIdx++) { VulkanFormat vkFormat = vkFormatList[fIdx]; @@ -915,15 +864,13 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, uint32_t elementSize = getVulkanFormatElementSize(vkFormat); ASSERT_LEQ(elementSize, (uint32_t)MAX_2D_IMAGE_ELEMENT_SIZE); log_info("elementSize= %d\n", elementSize); - std::map patternToSubstituteMap; - patternToSubstituteMap[GLSL_FORMAT_STRING] = - getVulkanFormatGLSLFormat(vkFormat); - patternToSubstituteMap[GLSL_TYPE_PREFIX_STRING] = - getVulkanFormatGLSLTypePrefix(vkFormat); - VulkanShaderModule vkImage2DShaderModule( - vkDevice, - prepareVulkanShader(vkImage2DShader, patternToSubstituteMap)); + std::string fileName = "image2D_" + + std::string(getVulkanFormatGLSLFormat(vkFormat)) + ".spv"; + log_info("Load %s file", fileName.c_str()); + vkImage2DShader = readFile(fileName); + VulkanShaderModule vkImage2DShaderModule(vkDevice, vkImage2DShader); + VulkanComputePipeline vkComputePipeline(vkDevice, vkPipelineLayout, vkImage2DShaderModule); @@ -931,13 +878,13 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, { uint32_t width = widthList[wIdx]; log_info("Width: %d\n", width); - ASSERT_LEQ(width, (uint32_t)MAX_2D_IMAGE_WIDTH); + if (width > max_width) continue; region[0] = width; for (size_t hIdx = 0; hIdx < ARRAY_SIZE(heightList); hIdx++) { uint32_t height = heightList[hIdx]; log_info("Height: %d\n", height); - ASSERT_LEQ(height, (uint32_t)MAX_2D_IMAGE_HEIGHT); + if (height > max_height) continue; region[1] = height; uint32_t numMipLevels = 1; @@ -1016,14 +963,6 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, const VulkanMemoryTypeList &memoryTypeList = vkDummyImage2D.getMemoryTypeList(); - std::vector - vkNonDedicatedImage2DListDeviceMemory1; - std::vector - vkNonDedicatedImage2DListDeviceMemory2; - std::vector - nonDedicatedExternalMemory1; - std::vector - nonDedicatedExternalMemory2; for (size_t mtIdx = 0; mtIdx < memoryTypeList.size(); mtIdx++) { @@ -1368,6 +1307,7 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, } } } + vkImage2DShader.clear(); } CLEANUP: if (clVk2CLExternalSemaphore) delete clVk2CLExternalSemaphore; @@ -1494,6 +1434,14 @@ int test_image_common(cl_device_id device_, cl_context context_, goto CLEANUP; } deviceId = devices[device_no]; + err = setMaxImageDimensions(deviceId, max_width, max_height); + if (CL_SUCCESS != err) + { + print_error(err, "error setting max image dimensions"); + goto CLEANUP; + } + log_info("Set max_width to %lu and max_height to %lu\n", max_width, + max_height); context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU, NULL, NULL, &err); if (CL_SUCCESS != err) diff --git a/test_conformance/vulkan/vulkan_interop_common/opencl_vulkan_wrapper.cpp b/test_conformance/vulkan/vulkan_interop_common/opencl_vulkan_wrapper.cpp index 136818f6..9d9a6601 100644 --- a/test_conformance/vulkan/vulkan_interop_common/opencl_vulkan_wrapper.cpp +++ b/test_conformance/vulkan/vulkan_interop_common/opencl_vulkan_wrapper.cpp @@ -23,6 +23,7 @@ #include #define ASSERT(x) assert((x)) +#define GB(x) ((unsigned long long)(x) << 30) pfnclCreateSemaphoreWithPropertiesKHR clCreateSemaphoreWithPropertiesKHRptr; pfnclEnqueueWaitSemaphoresKHR clEnqueueWaitSemaphoresKHRptr; @@ -31,7 +32,7 @@ pfnclEnqueueAcquireExternalMemObjectsKHR clEnqueueAcquireExternalMemObjectsKHRptr; pfnclEnqueueReleaseExternalMemObjectsKHR clEnqueueReleaseExternalMemObjectsKHRptr; -pfnclReleaseSemaphoreObjectKHR clReleaseSemaphoreObjectKHRptr; +pfnclReleaseSemaphoreKHR clReleaseSemaphoreKHRptr; void init_cl_vk_ext(cl_platform_id opencl_platform) { @@ -51,13 +52,13 @@ void init_cl_vk_ext(cl_platform_id opencl_platform) throw std::runtime_error("Failed to get the function pointer of " "clEnqueueSignalSemaphoresKHRptr!"); } - clReleaseSemaphoreObjectKHRptr = (pfnclReleaseSemaphoreObjectKHR) - clGetExtensionFunctionAddressForPlatform(opencl_platform, - "clReleaseSemaphoreObjectKHR"); - if (NULL == clReleaseSemaphoreObjectKHRptr) + clReleaseSemaphoreKHRptr = + (pfnclReleaseSemaphoreKHR)clGetExtensionFunctionAddressForPlatform( + opencl_platform, "clReleaseSemaphoreKHR"); + if (NULL == clReleaseSemaphoreKHRptr) { throw std::runtime_error("Failed to get the function pointer of " - "clReleaseSemaphoreObjectKHRptr!"); + "clReleaseSemaphoreKHRptr!"); } clCreateSemaphoreWithPropertiesKHRptr = (pfnclCreateSemaphoreWithPropertiesKHR) @@ -70,6 +71,40 @@ void init_cl_vk_ext(cl_platform_id opencl_platform) } } +cl_int setMaxImageDimensions(cl_device_id deviceID, size_t &max_width, + size_t &max_height) +{ + cl_int result = CL_SUCCESS; + cl_ulong val; + size_t paramSize; + + result = clGetDeviceInfo(deviceID, CL_DEVICE_GLOBAL_MEM_SIZE, + sizeof(cl_ulong), &val, ¶mSize); + + if (result != CL_SUCCESS) + { + return result; + } + + if (val < GB(4)) + { + max_width = 256; + max_height = 256; + } + else if (val < GB(8)) + { + max_width = 512; + max_height = 256; + } + else + { + max_width = 1024; + max_height = 512; + } + + return result; +} + cl_int getCLFormatFromVkFormat(VkFormat vkFormat, cl_image_format *clImageFormat) { @@ -798,10 +833,10 @@ clExternalSemaphore::clExternalSemaphore( clExternalSemaphore::~clExternalSemaphore() { - cl_int err = clReleaseSemaphoreObjectKHRptr(m_externalSemaphore); + cl_int err = clReleaseSemaphoreKHRptr(m_externalSemaphore); if (err != CL_SUCCESS) { - throw std::runtime_error("clReleaseSemaphoreObjectKHR failed!"); + throw std::runtime_error("clReleaseSemaphoreKHR failed!"); } } diff --git a/test_conformance/vulkan/vulkan_interop_common/opencl_vulkan_wrapper.hpp b/test_conformance/vulkan/vulkan_interop_common/opencl_vulkan_wrapper.hpp index c1d2a766..d9f8dccb 100644 --- a/test_conformance/vulkan/vulkan_interop_common/opencl_vulkan_wrapper.hpp +++ b/test_conformance/vulkan/vulkan_interop_common/opencl_vulkan_wrapper.hpp @@ -49,7 +49,7 @@ typedef cl_int (*pfnclEnqueueReleaseExternalMemObjectsKHR)( cl_command_queue command_queue, cl_uint num_mem_objects, const cl_mem *mem_objects, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event); -typedef cl_int (*pfnclReleaseSemaphoreObjectKHR)(cl_semaphore_khr sema_object); +typedef cl_int (*pfnclReleaseSemaphoreKHR)(cl_semaphore_khr sema_object); extern pfnclCreateSemaphoreWithPropertiesKHR clCreateSemaphoreWithPropertiesKHRptr; @@ -59,7 +59,7 @@ extern pfnclEnqueueAcquireExternalMemObjectsKHR clEnqueueAcquireExternalMemObjectsKHRptr; extern pfnclEnqueueReleaseExternalMemObjectsKHR clEnqueueReleaseExternalMemObjectsKHRptr; -extern pfnclReleaseSemaphoreObjectKHR clReleaseSemaphoreObjectKHRptr; +extern pfnclReleaseSemaphoreKHR clReleaseSemaphoreKHRptr; cl_int getCLImageInfoFromVkImageInfo(const VkImageCreateInfo *, size_t, cl_image_format *, cl_image_desc *); @@ -69,6 +69,8 @@ cl_int check_external_memory_handle_type( cl_int check_external_semaphore_handle_type( cl_device_id deviceID, cl_external_semaphore_handle_type_khr requiredHandleType); +cl_int setMaxImageDimensions(cl_device_id deviceID, size_t &width, + size_t &height); class clExternalMemory { protected: diff --git a/test_conformance/vulkan/vulkan_interop_common/vulkan_list_map.hpp b/test_conformance/vulkan/vulkan_interop_common/vulkan_list_map.hpp index 831403e1..10a7b221 100644 --- a/test_conformance/vulkan/vulkan_interop_common/vulkan_list_map.hpp +++ b/test_conformance/vulkan/vulkan_interop_common/vulkan_list_map.hpp @@ -335,11 +335,8 @@ const VulkanWrapper & template VulkanWrapper &VulkanList::operator[](size_t idx) { - if (idx < m_wrapperList.size()) - { - // CHECK_LT(idx, m_wrapperList.size()); - return m_wrapperList[idx].get(); - } + // CHECK_LT(idx, m_wrapperList.size()); + return m_wrapperList[idx].get(); } template diff --git a/test_conformance/vulkan/vulkan_interop_common/vulkan_utility.cpp b/test_conformance/vulkan/vulkan_interop_common/vulkan_utility.cpp index 81e12621..4e6118b1 100644 --- a/test_conformance/vulkan/vulkan_interop_common/vulkan_utility.cpp +++ b/test_conformance/vulkan/vulkan_interop_common/vulkan_utility.cpp @@ -18,6 +18,7 @@ #include "vulkan_wrapper.hpp" #include #include +#include #include #include #include @@ -541,59 +542,6 @@ const char *getVulkanFormatGLSLFormat(VulkanFormat format) return (const char *)size_t(0); } -const char *getVulkanFormatGLSLTypePrefix(VulkanFormat format) -{ - switch (format) - { - case VULKAN_FORMAT_R8_UINT: - case VULKAN_FORMAT_R8G8_UINT: - case VULKAN_FORMAT_R8G8B8A8_UINT: - case VULKAN_FORMAT_R16_UINT: - case VULKAN_FORMAT_R16G16_UINT: - case VULKAN_FORMAT_R16G16B16A16_UINT: - case VULKAN_FORMAT_R32_UINT: - case VULKAN_FORMAT_R32G32_UINT: - case VULKAN_FORMAT_R32G32B32A32_UINT: return "u"; - - case VULKAN_FORMAT_R8_SINT: - case VULKAN_FORMAT_R8G8_SINT: - case VULKAN_FORMAT_R8G8B8A8_SINT: - case VULKAN_FORMAT_R16_SINT: - case VULKAN_FORMAT_R16G16_SINT: - case VULKAN_FORMAT_R16G16B16A16_SINT: - case VULKAN_FORMAT_R32_SINT: - case VULKAN_FORMAT_R32G32_SINT: - case VULKAN_FORMAT_R32G32B32A32_SINT: return "i"; - - case VULKAN_FORMAT_R32_SFLOAT: - case VULKAN_FORMAT_R32G32_SFLOAT: - case VULKAN_FORMAT_R32G32B32A32_SFLOAT: return ""; - - default: ASSERT(0); std::cout << "Unknown format"; - } - - return ""; -} - -std::string prepareVulkanShader( - std::string shaderCode, - const std::map &patternToSubstituteMap) -{ - for (std::map::const_iterator psIt = - patternToSubstituteMap.begin(); - psIt != patternToSubstituteMap.end(); ++psIt) - { - std::string::size_type pos = 0u; - while ((pos = shaderCode.find(psIt->first, pos)) != std::string::npos) - { - shaderCode.replace(pos, psIt->first.length(), psIt->second); - pos += psIt->second.length(); - } - } - - return shaderCode; -} - std::ostream &operator<<(std::ostream &os, VulkanMemoryTypeProperty memoryTypeProperty) { @@ -691,3 +639,54 @@ std::ostream &operator<<(std::ostream &os, VulkanFormat format) return os; } + +static char *findFilePath(const std::string filename) +{ + const char *searchPath[] = { + "./", // Same dir + "./shaders/", // In shaders folder in same dir + "../test_conformance/vulkan/shaders/" // In src folder + }; + for (unsigned int i = 0; i < sizeof(searchPath) / sizeof(char *); ++i) + { + std::string path(searchPath[i]); + + path.append(filename); + FILE *fp; + fp = fopen(path.c_str(), "rb"); + + if (fp != NULL) + { + fclose(fp); + // File found + char *file_path = (char *)(malloc(path.length() + 1)); + strncpy(file_path, path.c_str(), path.length() + 1); + return file_path; + } + if (fp) + { + fclose(fp); + } + } + // File not found + return 0; +} + +std::vector readFile(const std::string &filename) +{ + char *file_path = findFilePath(filename); + + std::ifstream file(file_path, std::ios::ate | std::ios::binary); + + if (!file.is_open()) + { + throw std::runtime_error("failed to open shader spv file!\n"); + } + size_t fileSize = (size_t)file.tellg(); + std::vector buffer(fileSize); + file.seekg(0); + file.read(buffer.data(), fileSize); + file.close(); + printf("filesize is %d", fileSize); + return buffer; +} diff --git a/test_conformance/vulkan/vulkan_interop_common/vulkan_utility.hpp b/test_conformance/vulkan/vulkan_interop_common/vulkan_utility.hpp index 7022fd5a..04f5a594 100644 --- a/test_conformance/vulkan/vulkan_interop_common/vulkan_utility.hpp +++ b/test_conformance/vulkan/vulkan_interop_common/vulkan_utility.hpp @@ -66,4 +66,5 @@ operator<<(std::ostream& os, VulkanExternalSemaphoreHandleType externalSemaphoreHandleType); std::ostream& operator<<(std::ostream& os, VulkanFormat format); +std::vector readFile(const std::string& filename); #endif // _vulkan_utility_hpp_ diff --git a/test_conformance/vulkan/vulkan_interop_common/vulkan_wrapper.cpp b/test_conformance/vulkan/vulkan_interop_common/vulkan_wrapper.cpp index c044e009..e5d3a271 100644 --- a/test_conformance/vulkan/vulkan_interop_common/vulkan_wrapper.cpp +++ b/test_conformance/vulkan/vulkan_interop_common/vulkan_wrapper.cpp @@ -201,7 +201,8 @@ VulkanInstance::VulkanInstance(): m_vkInstance(VK_NULL_HANDLE) if (physicalDeviceCount == uint32_t(0)) { - throw std::runtime_error("failed to find GPUs with Vulkan support!"); + std::cout << "failed to find GPUs with Vulkan support!\n"; + return; } std::vector vkPhysicalDeviceList(physicalDeviceCount, @@ -846,23 +847,18 @@ VulkanShaderModule::VulkanShaderModule(const VulkanShaderModule &shaderModule) {} VulkanShaderModule::VulkanShaderModule(const VulkanDevice &device, - const std::string &code) + const std::vector &code) : m_device(device) { - std::string paddedCode = code; - while (paddedCode.size() % 4) - { - paddedCode += " "; - } VkShaderModuleCreateInfo vkShaderModuleCreateInfo = {}; vkShaderModuleCreateInfo.sType = VK_STRUCTURE_TYPE_SHADER_MODULE_CREATE_INFO; vkShaderModuleCreateInfo.pNext = NULL; vkShaderModuleCreateInfo.flags = 0; - vkShaderModuleCreateInfo.codeSize = paddedCode.size(); + vkShaderModuleCreateInfo.codeSize = code.size(); vkShaderModuleCreateInfo.pCode = - (const uint32_t *)(void *)paddedCode.c_str(); + reinterpret_cast(code.data()); vkCreateShaderModule(m_device, &vkShaderModuleCreateInfo, NULL, &m_vkShaderModule); diff --git a/test_conformance/vulkan/vulkan_interop_common/vulkan_wrapper.hpp b/test_conformance/vulkan/vulkan_interop_common/vulkan_wrapper.hpp index 1f68a92b..37925ee4 100644 --- a/test_conformance/vulkan/vulkan_interop_common/vulkan_wrapper.hpp +++ b/test_conformance/vulkan/vulkan_interop_common/vulkan_wrapper.hpp @@ -240,7 +240,8 @@ protected: VulkanShaderModule(const VulkanShaderModule &shaderModule); public: - VulkanShaderModule(const VulkanDevice &device, const std::string &code); + VulkanShaderModule(const VulkanDevice &device, + const std::vector &code); virtual ~VulkanShaderModule(); operator VkShaderModule() const; };