From ab106e97f96fbf66008c46dbc07cd8d0b39c4acf Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Fri, 4 Sep 2020 03:46:22 -0700 Subject: [PATCH] specialization constant test fixes (#927) * remove SPV_KHR_no_integer_wrap_decoration * address review comments * remove the assemble_spirv script There is another PR in flight that adds a much more complete script that we should merge instead. --- .../op_spec_constant_false_simple.spvasm32 | 2 -- .../op_spec_constant_false_simple.spvasm64 | 2 -- .../op_spec_constant_true_simple.spvasm32 | 2 -- .../op_spec_constant_true_simple.spvasm64 | 2 -- .../op_spec_constant_uchar_simple.spvasm32 | 2 -- .../op_spec_constant_uchar_simple.spvasm64 | 2 -- .../op_spec_constant_ushort_simple.spvasm32 | 2 -- .../op_spec_constant_ushort_simple.spvasm64 | 2 -- .../op_spec_constant_false_simple.spv32 | Bin 836 -> 784 bytes .../op_spec_constant_false_simple.spv64 | Bin 836 -> 784 bytes .../op_spec_constant_true_simple.spv32 | Bin 836 -> 784 bytes .../op_spec_constant_true_simple.spv64 | Bin 836 -> 784 bytes .../op_spec_constant_uchar_simple.spv32 | Bin 628 -> 576 bytes .../op_spec_constant_uchar_simple.spv64 | Bin 628 -> 576 bytes .../op_spec_constant_ushort_simple.spv32 | Bin 616 -> 564 bytes .../op_spec_constant_ushort_simple.spv64 | Bin 616 -> 564 bytes .../spirv_new/test_op_spec_constant.cpp | 4 +--- 17 files changed, 1 insertion(+), 19 deletions(-) diff --git a/test_conformance/spirv_new/spirv_asm/op_spec_constant_false_simple.spvasm32 b/test_conformance/spirv_new/spirv_asm/op_spec_constant_false_simple.spvasm32 index 36be62fd..8492474c 100644 --- a/test_conformance/spirv_new/spirv_asm/op_spec_constant_false_simple.spvasm32 +++ b/test_conformance/spirv_new/spirv_asm/op_spec_constant_false_simple.spvasm32 @@ -7,7 +7,6 @@ OpCapability Linkage OpCapability Kernel OpCapability Int8 - OpExtension "SPV_KHR_no_integer_wrap_decoration" %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical32 OpenCL OpEntryPoint Kernel %11 "spec_const_kernel" @@ -21,7 +20,6 @@ OpName %if_end "if.end" OpDecorate %test_value FuncParamAttr Zext OpDecorate %test_value LinkageAttributes "test_value" Export - OpDecorate %add NoSignedWrap OpDecorate %false SpecId 101 %uchar = OpTypeInt 8 0 %uint = OpTypeInt 32 0 diff --git a/test_conformance/spirv_new/spirv_asm/op_spec_constant_false_simple.spvasm64 b/test_conformance/spirv_new/spirv_asm/op_spec_constant_false_simple.spvasm64 index c9f53443..c76bccc8 100644 --- a/test_conformance/spirv_new/spirv_asm/op_spec_constant_false_simple.spvasm64 +++ b/test_conformance/spirv_new/spirv_asm/op_spec_constant_false_simple.spvasm64 @@ -7,7 +7,6 @@ OpCapability Linkage OpCapability Kernel OpCapability Int8 - OpExtension "SPV_KHR_no_integer_wrap_decoration" %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %11 "spec_const_kernel" @@ -21,7 +20,6 @@ OpName %if_end "if.end" OpDecorate %test_value FuncParamAttr Zext OpDecorate %test_value LinkageAttributes "test_value" Export - OpDecorate %add NoSignedWrap OpDecorate %false SpecId 101 %uchar = OpTypeInt 8 0 %uint = OpTypeInt 32 0 diff --git a/test_conformance/spirv_new/spirv_asm/op_spec_constant_true_simple.spvasm32 b/test_conformance/spirv_new/spirv_asm/op_spec_constant_true_simple.spvasm32 index 5e182e1a..834b85df 100644 --- a/test_conformance/spirv_new/spirv_asm/op_spec_constant_true_simple.spvasm32 +++ b/test_conformance/spirv_new/spirv_asm/op_spec_constant_true_simple.spvasm32 @@ -7,7 +7,6 @@ OpCapability Linkage OpCapability Kernel OpCapability Int8 - OpExtension "SPV_KHR_no_integer_wrap_decoration" %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical32 OpenCL OpEntryPoint Kernel %11 "spec_const_kernel" @@ -21,7 +20,6 @@ OpName %if_end "if.end" OpDecorate %test_value FuncParamAttr Zext OpDecorate %test_value LinkageAttributes "test_value" Export - OpDecorate %add NoSignedWrap OpDecorate %true SpecId 101 %uchar = OpTypeInt 8 0 %uint = OpTypeInt 32 0 diff --git a/test_conformance/spirv_new/spirv_asm/op_spec_constant_true_simple.spvasm64 b/test_conformance/spirv_new/spirv_asm/op_spec_constant_true_simple.spvasm64 index 720121f6..83ce4d62 100644 --- a/test_conformance/spirv_new/spirv_asm/op_spec_constant_true_simple.spvasm64 +++ b/test_conformance/spirv_new/spirv_asm/op_spec_constant_true_simple.spvasm64 @@ -7,7 +7,6 @@ OpCapability Linkage OpCapability Kernel OpCapability Int8 - OpExtension "SPV_KHR_no_integer_wrap_decoration" %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %11 "spec_const_kernel" @@ -21,7 +20,6 @@ OpName %if_end "if.end" OpDecorate %test_value FuncParamAttr Zext OpDecorate %test_value LinkageAttributes "test_value" Export - OpDecorate %add NoSignedWrap OpDecorate %true SpecId 101 %uchar = OpTypeInt 8 0 %uint = OpTypeInt 32 0 diff --git a/test_conformance/spirv_new/spirv_asm/op_spec_constant_uchar_simple.spvasm32 b/test_conformance/spirv_new/spirv_asm/op_spec_constant_uchar_simple.spvasm32 index af778e37..0e832b42 100644 --- a/test_conformance/spirv_new/spirv_asm/op_spec_constant_uchar_simple.spvasm32 +++ b/test_conformance/spirv_new/spirv_asm/op_spec_constant_uchar_simple.spvasm32 @@ -6,7 +6,6 @@ OpCapability Addresses OpCapability Kernel OpCapability Int8 - OpExtension "SPV_KHR_no_integer_wrap_decoration" %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical32 OpenCL OpEntryPoint Kernel %9 "spec_const_kernel" @@ -18,7 +17,6 @@ OpName %add "add" OpDecorate %spec_const_kernel_spec_constant_value Constant OpDecorate %spec_const_kernel_spec_constant_value Alignment 1 - OpDecorate %add NoSignedWrap OpDecorate %uchar_0 SpecId 101 %uchar = OpTypeInt 8 0 %uint = OpTypeInt 32 0 diff --git a/test_conformance/spirv_new/spirv_asm/op_spec_constant_uchar_simple.spvasm64 b/test_conformance/spirv_new/spirv_asm/op_spec_constant_uchar_simple.spvasm64 index a05d7b7a..89150fc6 100644 --- a/test_conformance/spirv_new/spirv_asm/op_spec_constant_uchar_simple.spvasm64 +++ b/test_conformance/spirv_new/spirv_asm/op_spec_constant_uchar_simple.spvasm64 @@ -6,7 +6,6 @@ OpCapability Addresses OpCapability Kernel OpCapability Int8 - OpExtension "SPV_KHR_no_integer_wrap_decoration" %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %9 "spec_const_kernel" @@ -18,7 +17,6 @@ OpName %add "add" OpDecorate %spec_const_kernel_spec_constant_value Constant OpDecorate %spec_const_kernel_spec_constant_value Alignment 1 - OpDecorate %add NoSignedWrap OpDecorate %uchar_0 SpecId 101 %uchar = OpTypeInt 8 0 %uint = OpTypeInt 32 0 diff --git a/test_conformance/spirv_new/spirv_asm/op_spec_constant_ushort_simple.spvasm32 b/test_conformance/spirv_new/spirv_asm/op_spec_constant_ushort_simple.spvasm32 index 75d2e24f..d5012070 100644 --- a/test_conformance/spirv_new/spirv_asm/op_spec_constant_ushort_simple.spvasm32 +++ b/test_conformance/spirv_new/spirv_asm/op_spec_constant_ushort_simple.spvasm32 @@ -6,7 +6,6 @@ OpCapability Addresses OpCapability Kernel OpCapability Int16 - OpExtension "SPV_KHR_no_integer_wrap_decoration" %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical32 OpenCL OpEntryPoint Kernel %9 "spec_const_kernel" @@ -17,7 +16,6 @@ OpName %entry "entry" OpDecorate %spec_const_kernel_spec_constant_value Constant OpDecorate %spec_const_kernel_spec_constant_value Alignment 2 - OpDecorate %add NoSignedWrap OpDecorate %ushort_0 SpecId 101 %ushort = OpTypeInt 16 0 %uint = OpTypeInt 32 0 diff --git a/test_conformance/spirv_new/spirv_asm/op_spec_constant_ushort_simple.spvasm64 b/test_conformance/spirv_new/spirv_asm/op_spec_constant_ushort_simple.spvasm64 index 4f3582eb..27f5fb8f 100644 --- a/test_conformance/spirv_new/spirv_asm/op_spec_constant_ushort_simple.spvasm64 +++ b/test_conformance/spirv_new/spirv_asm/op_spec_constant_ushort_simple.spvasm64 @@ -6,7 +6,6 @@ OpCapability Addresses OpCapability Kernel OpCapability Int16 - OpExtension "SPV_KHR_no_integer_wrap_decoration" %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %9 "spec_const_kernel" @@ -17,7 +16,6 @@ OpName %entry "entry" OpDecorate %spec_const_kernel_spec_constant_value Constant OpDecorate %spec_const_kernel_spec_constant_value Alignment 2 - OpDecorate %add NoSignedWrap OpDecorate %ushort_0 SpecId 101 %ushort = OpTypeInt 16 0 %uint = OpTypeInt 32 0 diff --git a/test_conformance/spirv_new/spirv_bin/op_spec_constant_false_simple.spv32 b/test_conformance/spirv_new/spirv_bin/op_spec_constant_false_simple.spv32 index 045abe523e6a2c2a77830470d1300d4786eb8b2e..adc01353be692106471cd9c1d55dfded4304a7d9 100644 GIT binary patch literal 784 zcmZXS$w~u36h&Vq;}l~|oLwXciUSR992UB9;YRQinh?tYrPFpgLUw+E{~`XM3X1o1 zmm*l9a^G<4z9A_U*Q;g)GppK~zBvsmTG?a8@wasU*3WZgo0P#e)soGNIVD7&gFh{7{TcaakcVIQeoGsj zmEF-U8ba)*VspZ_xJ@BDpk_yYOGu9Rp5DNxzkRW=nS+DYn{;^lC4_UR4&d>g-;smf zfppMgx4dDOM^Zv>UKh<5W$9tT$C!s1Uwz$0pCk667x?Umy`zbJ>@}8T|NNKO$%FXp lWLcNolQi*He2iVeq5i6mpNNA7KC__nF6*Zr_^Bu}e*?giN9h0n literal 836 zcmZXS%T5$Q6owBoj3c)}1jO3{LIM%UByLm|vT=b4CK^rHRyspDOd!=wPc?)cPvCnP zUrJu?Ltna!GArDXALL~gM|m)tu};nWx$xn`r~Wt|nf0t)`�ew}%EV`nd1|H^{>{ zx)Wape=1+s=7ctD94;$|E?ySC{~xwr4UV&NYe(32{h1X%CRDvs&k?;?&$H9YTkT;c zA7Xha*-;1dkXNy&?obU*i@ArJRJUfM@BMh}L(MbQ*gj~keN|#-8@il*b&vf@_kR5* z*475is)Z(RMmjH`_H0>Mvy#-}V9!aeV(4tAc&{<mjm zQNBgVRbiJT*#|Y(WG_n+!~a@2;OXy;Q0UCWf$O(;$o(e?=AAkKM?P=Lf$yz&;A79^ zu-9EN!S_TL&Yz0X!%Ffo4>SHsb)TjfzQ4-AW=FgMoapZp$1~}-|0{NKCOkWNu1oAn hoY=epK6V9$`Y&|x?+F73HnYISVR!jf@ diff --git a/test_conformance/spirv_new/spirv_bin/op_spec_constant_false_simple.spv64 b/test_conformance/spirv_new/spirv_bin/op_spec_constant_false_simple.spv64 index 701527d85ea3ccea6e30a2d194cfa3f4f92cafe8..2046ab64ba834a20300443facbd3447220ff4f05 100644 GIT binary patch literal 784 zcmZXS$w~uJ5Jf9P8>bkjI4=?e#eoJl4h!A5$VTuJnh>8x3~ifq1b2Rc{~`XM3X1o1 zKP6y=mzr+f8j@0Rvtm{-vx;r#8=14BmAwwU9`kxzddh;V;xcjZ`9(NLugqqwq$mo@={Ru zAe%3BVts7)hR(7Gv${C#KS>9gfEn(GX)^pnZIq^0Tccl1X4=ee>?%HWzB$?8H*NyCrePfIs&Bzrzz&D}qLMm`+m!PmUMsEf|Z zZt50u(%1`%%}Y0gU6AGm^lZvslqN^~KsDf*?@%ag_CVlz7Z0`H(qN7?0XXXU9TE7B z!~-9$e~JOHp?I0mD~E>Hq)$ literal 836 zcmZXS+e#ck5Qa`DAC_UqBkK&zA;Co13 zPy-SFZ>AL?HS|_j)n9)ddRxOQ9kU@b>)1Q}BC|GZBT2WDewOr_^tisC-+#F;pMSWJ zyMBoN(3kGK%nEnldwH3~ejdywtX(sID}4C$wL6FhW>eM@yVYmfMh#xrXI5Mdsd~GfBYLr(Wha%l+QUpf z#PUMAtq$lRuVPU>q8gkQ^8`1dZq53?yYa|}nrE!Beb8L{s>IG#b&!2^kG-z*w0;w7 zX@h3fY?C*kJ1?O2Y*AUyrKyErPf9OgNH$Y^sxcq!Y1wcP2S1bS1#MSlY*~Amljc4z z-@NpaunW@cgPO~-7o~~ef3FX)d4v2c~bzd@GppLFzBx@RTG?a8+9tJ<{KW_h67i?6%MD00C+ZjknGdhyPlxTWl47CjeK zy^{5pIH^AN?o(%3gkD{k#3Nxv9nixdO``OdmGd<4A6{AAn%!On$&Cv$or&Ccq2}tS z61&*c<#g0N{A7 zve&hXrVzWO*sO3(+?J3XP;*^=TS$)hzTUv6zXP$bnS+DYn{;^lC4_UN4&d>g-;smf zp>)t=x4dDO$5KLXP8ZD=W$9tw$C!s1Uwz#LpCk667x?U)J)sGnc#TEbKmR3m@*qAt lS<)r5J=osIztH)Np;gx4PnPG@P82h zpqNNJUw27JtWM3Td+(`pA2sdP>#o^|nRRVNzpQC%*=WK}!Y2ve3ZLscKK$ywd_Hm^ zcY_f9sW07CnHBED_wzD~gFKjxS*KP$D17*|-y6mgvk7b0zLs<*?Wv*5d06*rw@ z-FIIGe=c9w#>F;k94;$QUA!)Q?>}y@>Yruh#+JD4`ZFtjj3|4j?jw4sUS{W&x7tHb zKE(1`*ii-4kXNy&9#{=egL%XoRkdb=Z@qZtL#^|yv3*cqd&$?#U~o`!Jd?j266aP$)3}6mB!XIml+}F zS@~v#tK!ZHnFl%7WX}tU;osIB_|*4KEY9@ALhF}wxO*>zv!e>&ai2G3q4!=o=rL#R zFxL-KLhq$EnzxEl!>i<@AA0tBbqRNShV70j$+>-y%*TG7fL1CL`KcZ5|7vW%ELbE$E{8rYJoudLQ#J`|rQRP9DW) mCo8(dUZjb=>V51A4)xc(|6Ck2u$cv&Prax~J@8vzX8r}jGe_wF literal 836 zcmZXSOG_L<6os#jlQBMH)R>qo6d}PR46gE6=q3vV14;z9MKe~ML8)$gsuSGh7x+IU ze=rIL&(~dwU^O+T?!BkZeblr@*E(h+X4bJc`en`8sEs9TCw!9dop3_m!TzD!-TC1{ z?gkQ?Vr8#i>ZhlTICULJc)gF5C zA(rRDwkn{8yoyD2&uVxY%mdz-sx=$@?8d$iwa(MV_CbAZD-$z&tIgV0_1Np$5A!#% zmL_Oc%{BKXb>;}lZ&hSrA^pJ~i{dsa3Y#Nj_n_JXFXG`6a_JQs4F zmv3IUBJP5ad608e_M(s&{texMPko=n;!IC0wEjtlyZ?o7zNi9t+~-YM=zW$Bdd!(S z%=N33(0iec=B1+4u#|lCLyvz`-j^wc?^-vonLl$v6WmN1uXMiqUon$2@tMi8HnBg_ h#O4j~F)KLaf31ywOB^(?=>?rDoybW(@J3#Gz6CjHR!jf@ diff --git a/test_conformance/spirv_new/spirv_bin/op_spec_constant_uchar_simple.spv32 b/test_conformance/spirv_new/spirv_bin/op_spec_constant_uchar_simple.spv32 index 94ce6666a3b474b9c94c0de4161e8be15e7f8e52..ac0536a7abcebc4483e39db268636c42e99c4426 100644 GIT binary patch delta 15 Wcmeyua)4!m&c+ZX#?48LLW}?|hXkqs delta 63 zcmX@W@`YuBjv5yO7ejDBSiHAKP<&o~d}dxrYIe0fn~L3~PTa(+=_NoIcDMlBXb S4rT^+1_p*w!OdEXVvGQe5E6y} diff --git a/test_conformance/spirv_new/spirv_bin/op_spec_constant_uchar_simple.spv64 b/test_conformance/spirv_new/spirv_bin/op_spec_constant_uchar_simple.spv64 index 6dc981679c8c0154d65576f3d3bcfeda65cb2ece..5f2e95ddec579e9e84fd2393c6985f0048f23ea0 100644 GIT binary patch delta 15 Wcmeyua)4!m&c+ZX#?48LLW}?|hXkqs delta 63 zcmX@W@`YuBjv5yO7ejDBSiHAKP<&o~d}dxrYIe0fn~L3~PTa(+=_NoIcDMlBXb S4rT^+1_p*w!OdEXVvGQe5E6y} diff --git a/test_conformance/spirv_new/spirv_bin/op_spec_constant_ushort_simple.spv32 b/test_conformance/spirv_new/spirv_bin/op_spec_constant_ushort_simple.spv32 index c81e309a1f31b113221ac0f7176794711086a4cf..84926b4cbd3911b150ecb4146d2d7e4cf3c73f82 100644 GIT binary patch delta 271 zcmXw!O%4H35QNLq{uyLqKlWA@Nc^lE!4;f?BM?_{C?^mbiLb^?CH4Asbyp?*if7yw z_tX0l>PTw9Gn^^ZrB0yXzDF(GfMqqe$0+Fg|8MEbl{BdIA+-gaH)n4|_00N~iG8u- zF&wx;^g^-#Y*?#V8XWXt(x&ct9hyC+DL?u*h5_v~uO!uHWb~dN(46Aa>b)Xj|Pr7vfJ#{dbL$?qM|s{C!MNm8lKfrACff8q9lF{Ke;Kkh+KLiOR=Pm zdj-%yA9Dk}3$Qpl=spM3xY*Of@-a96*PqW5Ls|vzqE`Xt1<K0rND^89$_mi!GW>iN_3xD+fv(riBi!}|jG7#cPJ diff --git a/test_conformance/spirv_new/spirv_bin/op_spec_constant_ushort_simple.spv64 b/test_conformance/spirv_new/spirv_bin/op_spec_constant_ushort_simple.spv64 index 8f7074280a9357028bcd79c17c20988e948303d9..00d01a7f44bab3e1b43dd1bccac3a6b50005c93d 100644 GIT binary patch delta 271 zcmXw!O%4H35QNLq{uyLqKlWA@Nc^lE!4;f?BM?_{C?^mbiLb^?CH4Asbyp?*if7yw z_tX0l>PTw9Gn^^ZrB0yXzDF(GfMqqe$0+Fg|8MEbl{BdIA+-gaH)n4|_00N~iG8u- zF&wx;^g^-#Y*?#V8XWXt(x&ct9hyC+DL?u*h5_v~uO!uHWb~dN(46Aa>b)Xj|Pr7vfJ#{dbL$?qM|s{C!MNm8lKfrACff8q9lF{Ke;Kkh+KLiOR=Pm zdj-%yA9Dk}3$Qpl=spM3xY*Of@-a96*PqW5Ls|vzqE`Xt1<K0rND^89$_mi!GW>iN_3xD+fv(riBi!}|jG7#cPJ diff --git a/test_conformance/spirv_new/test_op_spec_constant.cpp b/test_conformance/spirv_new/test_op_spec_constant.cpp index a6feddfd..a280a4f7 100644 --- a/test_conformance/spirv_new/test_op_spec_constant.cpp +++ b/test_conformance/spirv_new/test_op_spec_constant.cpp @@ -20,8 +20,7 @@ Agreement as executed between Khronos and the recipient. template int run_case(cl_device_id deviceID, cl_context context, cl_command_queue queue, const char *name, T init_buffer, T spec_constant_value, - T final_value, bool use_spec_constant, - bool (*notEqual)(const T &, const T &) = isNotEqual) + T final_value, bool use_spec_constant) { clProgramWrapper prog; cl_int err = CL_SUCCESS; @@ -54,7 +53,6 @@ int run_case(cl_device_id deviceID, cl_context context, cl_command_queue queue, err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &work_size, NULL, 0, NULL, NULL); SPIRV_CHECK_ERROR(err, "Failed to enqueue kernel"); - clFinish(queue); T device_results = 0; err = clEnqueueReadBuffer(queue, output_buffer, CL_TRUE, 0, bytes,