From 32e702efbc2c31e2cba11c743c318abbf0ffeea3 Mon Sep 17 00:00:00 2001 From: Grzegorz Wawiorko <35483345+gwawiork@users.noreply.github.com> Date: Wed, 22 May 2019 16:56:43 +0200 Subject: [PATCH] =?UTF-8?q?cl21:=20Fix=20Issue=2038=20-=20Test=20spirv=5Fn?= =?UTF-8?q?ew=20-=20new=20test=20for=20cl=5Fkhr=5Fspirv=5Fno=5Finteger=5F?= =?UTF-8?q?=E2=80=A6=20(#119)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * Fix Issue 38 - Test spirv_new - new test for cl_khr_spirv_no_integer_wrap_decoration SPV_KHR_no_integer_wrap_decoration extension * Remove not used variable. --- ..._no_integer_wrap_decoration_fadd_int.spv32 | Bin 0 -> 732 bytes ..._no_integer_wrap_decoration_fadd_int.spv64 | Bin 0 -> 820 bytes ...no_integer_wrap_decoration_fadd_uint.spv32 | Bin 0 -> 732 bytes ...no_integer_wrap_decoration_fadd_uint.spv64 | Bin 0 -> 820 bytes ..._no_integer_wrap_decoration_fmul_int.spv32 | Bin 0 -> 732 bytes ..._no_integer_wrap_decoration_fmul_int.spv64 | Bin 0 -> 820 bytes ...no_integer_wrap_decoration_fmul_uint.spv32 | Bin 0 -> 732 bytes ...no_integer_wrap_decoration_fmul_uint.spv64 | Bin 0 -> 820 bytes ..._integer_wrap_decoration_fnegate_int.spv32 | Bin 0 -> 700 bytes ..._integer_wrap_decoration_fnegate_int.spv64 | Bin 0 -> 772 bytes ...teger_wrap_decoration_fshiftleft_int.spv32 | Bin 0 -> 768 bytes ...teger_wrap_decoration_fshiftleft_int.spv64 | Bin 0 -> 856 bytes ...eger_wrap_decoration_fshiftleft_uint.spv32 | Bin 0 -> 768 bytes ...eger_wrap_decoration_fshiftleft_uint.spv64 | Bin 0 -> 856 bytes ..._no_integer_wrap_decoration_fsub_int.spv32 | Bin 0 -> 732 bytes ..._no_integer_wrap_decoration_fsub_int.spv64 | Bin 0 -> 820 bytes ...no_integer_wrap_decoration_fsub_uint.spv32 | Bin 0 -> 732 bytes ...no_integer_wrap_decoration_fsub_uint.spv64 | Bin 0 -> 820 bytes ..._integer_wrap_decoration_fadd_int.spvasm32 | 47 ++++ ..._integer_wrap_decoration_fadd_int.spvasm64 | 53 +++++ ...integer_wrap_decoration_fadd_uint.spvasm32 | 47 ++++ ...integer_wrap_decoration_fadd_uint.spvasm64 | 53 +++++ ..._integer_wrap_decoration_fmul_int.spvasm32 | 47 ++++ ..._integer_wrap_decoration_fmul_int.spvasm64 | 53 +++++ ...integer_wrap_decoration_fmul_uint.spvasm32 | 47 ++++ ...integer_wrap_decoration_fmul_uint.spvasm64 | 53 +++++ ...teger_wrap_decoration_fnegate_int.spvasm32 | 46 ++++ ...teger_wrap_decoration_fnegate_int.spvasm64 | 51 ++++ ...er_wrap_decoration_fshiftleft_int.spvasm32 | 49 ++++ ...er_wrap_decoration_fshiftleft_int.spvasm64 | 55 +++++ ...r_wrap_decoration_fshiftleft_uint.spvasm32 | 49 ++++ ...r_wrap_decoration_fshiftleft_uint.spvasm64 | 55 +++++ ..._integer_wrap_decoration_fsub_int.spvasm32 | 47 ++++ ..._integer_wrap_decoration_fsub_int.spvasm64 | 53 +++++ ...integer_wrap_decoration_fsub_uint.spvasm32 | 47 ++++ ...integer_wrap_decoration_fsub_uint.spvasm64 | 53 +++++ ...l_khr_spirv_no_integer_wrap_decoration.cpp | 223 ++++++++++++++++++ 37 files changed, 1128 insertions(+) create mode 100644 test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fadd_int.spv32 create mode 100644 test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fadd_int.spv64 create mode 100644 test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fadd_uint.spv32 create mode 100644 test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fadd_uint.spv64 create mode 100644 test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fmul_int.spv32 create mode 100644 test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fmul_int.spv64 create mode 100644 test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fmul_uint.spv32 create mode 100644 test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fmul_uint.spv64 create mode 100644 test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fnegate_int.spv32 create mode 100644 test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fnegate_int.spv64 create mode 100644 test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fshiftleft_int.spv32 create mode 100644 test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fshiftleft_int.spv64 create mode 100644 test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fshiftleft_uint.spv32 create mode 100644 test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fshiftleft_uint.spv64 create mode 100644 test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fsub_int.spv32 create mode 100644 test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fsub_int.spv64 create mode 100644 test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fsub_uint.spv32 create mode 100644 test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fsub_uint.spv64 create mode 100644 test_conformance/spirv_new/spirv_txt/ext_cl_khr_spirv_no_integer_wrap_decoration_fadd_int.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_txt/ext_cl_khr_spirv_no_integer_wrap_decoration_fadd_int.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_txt/ext_cl_khr_spirv_no_integer_wrap_decoration_fadd_uint.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_txt/ext_cl_khr_spirv_no_integer_wrap_decoration_fadd_uint.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_txt/ext_cl_khr_spirv_no_integer_wrap_decoration_fmul_int.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_txt/ext_cl_khr_spirv_no_integer_wrap_decoration_fmul_int.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_txt/ext_cl_khr_spirv_no_integer_wrap_decoration_fmul_uint.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_txt/ext_cl_khr_spirv_no_integer_wrap_decoration_fmul_uint.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_txt/ext_cl_khr_spirv_no_integer_wrap_decoration_fnegate_int.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_txt/ext_cl_khr_spirv_no_integer_wrap_decoration_fnegate_int.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_txt/ext_cl_khr_spirv_no_integer_wrap_decoration_fshiftleft_int.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_txt/ext_cl_khr_spirv_no_integer_wrap_decoration_fshiftleft_int.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_txt/ext_cl_khr_spirv_no_integer_wrap_decoration_fshiftleft_uint.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_txt/ext_cl_khr_spirv_no_integer_wrap_decoration_fshiftleft_uint.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_txt/ext_cl_khr_spirv_no_integer_wrap_decoration_fsub_int.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_txt/ext_cl_khr_spirv_no_integer_wrap_decoration_fsub_int.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_txt/ext_cl_khr_spirv_no_integer_wrap_decoration_fsub_uint.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_txt/ext_cl_khr_spirv_no_integer_wrap_decoration_fsub_uint.spvasm64 create mode 100644 test_conformance/spirv_new/test_cl_khr_spirv_no_integer_wrap_decoration.cpp diff --git a/test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fadd_int.spv32 b/test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fadd_int.spv32 new file mode 100644 index 0000000000000000000000000000000000000000..0c127d85e47582d1aa6623e88d236851b707c199 GIT binary patch literal 732 zcma)4O-sW-5S=Er(Z;scT6^(O5d<$DRgj7(7^8w;;B|>@w1G4U`KVWawim(o)+~5* zm&we$eKYftjcx5Y(rAc?BO^Sy`eMmS?VGi4!MoBFGVil6xx70p%9%)4Y>ipKd(y%d zRExYVf^x2=F+wo7Mb106w6UKbUexXNlS)z*C1v;?=ISF4GBpdQX&#hOng}Mjx<;uY zi}IzqsG_*^6E99*gV;}&=~NSaNC4N?def@ZF*V2YLdVpTBlf^>5bwi?`B2H6IJLR3 z9i9Jb-PUnOpB*6nj}Cqt4;8EMRbS+$o=H8NIYhn&KZbFQkhkH9QNtKHb#M_MK;t?* z@eVxukh_a+U3i`mH+EXrvlrJ9D#+n|aIxkXlgQ`&nwWh}Or2k-=)q6GqK*w1dp@Ec z=i}8h1}9g**l!cYK0KR1A9({9an9#F!`(v6Ie)oBU_XW8K?48) literal 0 HcmV?d00001 diff --git a/test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fadd_int.spv64 b/test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fadd_int.spv64 new file mode 100644 index 0000000000000000000000000000000000000000..a2c4ea1616e92dc7b3e71c2f88c1366f5ba09111 GIT binary patch literal 820 zcma))-7f=C5XFb0sG`2V9*9KZ#Um1_G$g7>d?j9Qs){CSx3=4Q^=EsLIKOLe;?dnq z=giqTGiP^O7~U@0(2$uGZ9~7DiVfRHo=5XMmgkak$wm#ub+;X#UhVa=C$niQNM$}| zY)scLetv6audWrTR0o$?tz584UH3N+?(+Q3#fj_0S^K4(x{tKkbx-Y9k~XtW5}P)a zyvDg+x04RsNxu_ijkq2qkIkqN50aK|Y{-Io3*L9q&pgI^^xE?nKdEqC9Yx_C1zbap zsp0X)^YI?)s;HMdE_=Ky{2w2@8}#GK{fgN|z7s#}TvR=wTvKqUCzUrTbe2`)gC0C= zs-9Bjz|YHbSs8v>nYrkf~$$hk9V*te))4}4Yr(eTdU@rR$FT^43u(5?vM q3qEL9h2ewtQ`4Erw_q;!!@M=&zw_2rvk(4QI%;%b*0=ukhl(F1*hGB* literal 0 HcmV?d00001 diff --git a/test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fadd_uint.spv32 b/test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fadd_uint.spv32 new file mode 100644 index 0000000000000000000000000000000000000000..0cb6df83d73b6b1f1749e6e818757e3ac470b2dc GIT binary patch literal 732 zcma)4O-sW-5S=Er(Z;scT6^(O5d<$DRgj7(7^8w;;B|>@w1G4UNgA*IY%hZEty%Es zE|Zyg`)1}P8{67*q|p!&M@D#Z^~I8v+Ba+8f_J4UWZq|Ca(Q=Hlrxd8*c!8d_oRg_ zs1|u!1m#>!V}xLEi=20CX=6V>yr|pjCzYfsO3Ls(%+*I8WNH>p(>y4nG!aa4b&XO* z7Uh+?SVnQ_CtjSs2C<)1=~NSaNC4N?dedd8V``4)g^sBwN9=*&Al`=&^P!SCacXm6 zJ39Z>x~=1mK084CA07NQ9x7JhtG>uhJ(GGkbBKHmehlLpA#cMIqlPhZ>fj7XFZGGl6`n?>UlTp$NDqy On3HE}zwt+&zd?j9Qs){CSx3=4Q^=EsLIKOLe;?dnq z=giqTGiP^O7~U@0(2$uGZ9~7DiVfRHo=5XMmgkak$wm#ub+;X#UhVa=C$niQNM$}| zY)scLetv6audWrTR0o$?tz584UH3N+?(+Q3#fj_0S^K4(x{tKkbx-Y9k~XtW5}P)a zyvDg+x04RsNxu_ijkq2qkIkqN50aK|Y{-Io3*L9q&pgI^^xE?nKdEqC9Yx_C1zbap zsp0X)^YI?)s;HMdE_=Ky{2w2@8}#GKgNoTjz7s#}TvR=wTvKqUCzUrTbe2`)gC0C= zs-9Bjz|YHbSs8v>nYrkf~$$hk9V*te))4}4Yr(eTdU@rR$FT^43u(5?vM q3qEL9h2ewtQ`4Erw_q;!!@M=&zw_2rvk(4QI%;%b*0=ukhl(F3tVDhQ literal 0 HcmV?d00001 diff --git a/test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fmul_int.spv32 b/test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fmul_int.spv32 new file mode 100644 index 0000000000000000000000000000000000000000..0244f5fbb2c3fc650dc53eab4fd1d2a7a58d76f4 GIT binary patch literal 732 zcma)4O-sW-5S=Er(Z;scT6^(O5d<$DRgj7(7^8w;;B|>@w1G4U`KW)uAMHi(y)_FS z-DNT}Z{N(kWMf-ZY3uD)2ZQu}7@Tkx(lh0OabOfK&Zi*hE?6W|wcfNUbxh6iywEZAZ#EqTS_3XuUgbH$aA6%?C#w7B2zb0m16I16GDthn}u&84L#-5Mp z$N6|QjlszkF!tMou@BEC&_~_?Mx66G&v3U8bIxDx5IJ$f{j6v4Te1(&O+D|1{aAkn O9&_?c?Kl3&6W9-i@<9gx literal 0 HcmV?d00001 diff --git a/test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fmul_int.spv64 b/test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fmul_int.spv64 new file mode 100644 index 0000000000000000000000000000000000000000..bbed3eb40cfb4238c8745168812233c7e44db9ed GIT binary patch literal 820 zcma))-7f=C5XFb0sG`2V9*9KZ#Um1_G$g7>d?j9Qs){CSx3=5*2mH}qB+l>Jn|O3L z(>Zf?&dk}}7KXQrHZ){rMcdFXr((l4lIPJpkL9_fT(VI^aouglr&oKu?8$7}3R0Pm z85`5}i=W?`*{f?sD%HVdRx1~5QrG>>gS$Mxb8+H2an^omr|u(dcHL9Em88wAlfe8ym8q-h%g?^fQm~9=-NF#!o6-S4UBJM*-K6 zV`_N3@qE08x+>}=kINqK3jfCk?*{$2a=&7Bk?+J0I~P??DAyDm>Ph8I3Y}%u_@D<5 zo2sXjIq>uHTvmpkR%R~xElrA{&b}1%@Vz)_4z!7WzI(vT4H(~FQp_q> z)rAi|n)i3Bd3c@6IVF2n51!@~@TqEg7IN-P8ul$J*aKgce>A*vc>LifXqSbV7qlzF r_<|4GRblv`{nT`3@-3Ll{V;D$`0u=R)$D^mmW~=-nDwoH{h{IqCn!XL literal 0 HcmV?d00001 diff --git a/test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fmul_uint.spv32 b/test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fmul_uint.spv32 new file mode 100644 index 0000000000000000000000000000000000000000..1614552d2b8e70c2944d09ac20a26d36ae16cabe GIT binary patch literal 732 zcma)4O-sW-5S=Er(Z;scT6^(O5d<$DRgj7(7^8w;;B|>@w1G4UNgDrvKiZ4ndutXv zy31r{-oBZ6$;P(!9BDK}#E}u6Tz#=*rS{F*x8PlA3Yqs=m|Wf+7UfK&E4Icg;5}(! z3#vul7C|{z(-E|yVT`iU2(uR-i5RXWu~9}>W|wcd1D>X@42d7)$K$q{>CIEeRQ#C)h^PMq3Y z*pAMBwQlRUqt6Zy|3?SEjfaX=_^K~*Q_rLx&Kx3NgCE1VM#$Un#HeA6oI1FO51?@! zo_GhIeaPKKw=O)-h#Nbt>)DIz2o>b;KDbzOj7jA4eof52CZ^6WRP^8{U{S{gj6EOG zkMr?r8iSK7VC=UEV;`PPppU!(j5z0Wp5bmG=A6IWA#&n|`&rN8w`3oln|j_2`?3BE OJm%z?+Hd@kC$Jxgc0mXL literal 0 HcmV?d00001 diff --git a/test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fmul_uint.spv64 b/test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fmul_uint.spv64 new file mode 100644 index 0000000000000000000000000000000000000000..bcf14c846daec74f189801eab47e3c220a6f0094 GIT binary patch literal 820 zcma))-7f=C5XFb0sG`2V9*9KZ#Um1_G$g7>d?j9Qs){CSx3=5*2mH}qB+l>Jn|O3L z(>Zf?&dk}}7KXQrHZ){rMcdFXr((l4lIPJpkL9_fT(VI^aouglr&oKu?8$7}3R0Pm z85`5}i=W?`*{f?sD%HVdRx1~5QrG>>gS$Mxb8+H2an^omr|u(dcHL9Em88wAlfe8ym8q-h%g?^fQm~9=-NF#!o6-S4UBJM*-K6 zV`_N3@qE08x+>}=kINqK3jfCk?*{$2@}OdNk?+J0I~P??DAyDm>Ph8I3Y}%u_@D<5 zo2sXjIq>uHTvmpkR%R~xElrA{&b}1%@Vz)_4z!7WzI(vT4H(~FQp_q> z)rAi|n)i3Bd3c@6IVF2n51!@~@TqEg7IN-P8ul$J*aKgce>A*vc>LifXqSbV7qlzF r_<|4GRblv`{nT`3@-3Ll{V;D$`0u=R)$D^mmW~=-nDwoH{h{IqDEvf% literal 0 HcmV?d00001 diff --git a/test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fnegate_int.spv32 b/test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fnegate_int.spv32 new file mode 100644 index 0000000000000000000000000000000000000000..5bd6eb570a3dad29d95806f66a50a0d11b1fd3fb GIT binary patch literal 700 zcma)4O-sW-5S=Er)z-AuTD^Fv2!aSbsvs3nFh&Kx!0Qs*Xag|`$+q77-ChLW+ivd8 zGMTS8GmqW4&VEyBH4$mbHolaeII?1G-P#6hTk1mOLms79_oGrzMLOb|oP&H<8aRS( zk@wF~&(tJ=3j~kIh385O=jHKL-Q7H^G*fY^qmQUiUqzU!X*9`-P{&y+*wi*XN|kwB zEY#&ZPIQp^N%j^dLAuB$1{i<>u&&{o&9(8#IhmElC!Yeo4@49G0hBo(J*%+uo=5=3 ztPw+=ZRoCvm%Mo|zIdi~2>(Bvcu%)>C;4#a0Pz~^5K1|+acddM4r1;mhCGId*KLjY zK6rXI-+|>^V)rn;3(Fewm7dvt&Y~QngBZRKg*_+OL_FWS@;SHi$@>Q#8?a-`4_9XR zJz*aAZ$g=)TqDPMeJJO!HU>Z9wxG;&Klgp(%RJwKwW@ajpLO=0V+vzps=ig8m^beS9N#utl~t4aUQ$}nfl7SRE+{Z$-F#FV!@`m<|s8y z!)&fDXJM4P@i0o>yvU8`iLV=7NIhHsS*~YeR6BfwPu~ z3-vr>4|VcccXhtdy{ThIpB*Cpk0!pSTeFjTxU++N1vr4ynAdlt$PbWH6U9PZ;g<{sAj&`(Ym&N}r`&$pSy f@5ZxY))r#UFJ^5cWu!I|?8 zeDsdZ?#|8bPA+lE+pd&KBI3$Ao?LAy%bf8QsR)_3lQ22I*`MYkk)}8rD+6yy z6OHsg263v+Hj?gBysEIx#fa_?zX_4!gn&a_Q$JCP{_P}rv@4$%pP|2J) zwbfx8I{($Ws^f+}+eQ2z9lRS474z^#Tja{jq#n-fB42>-!?^ZLUPHVI>mjEOa;T|~ zd=Z|Dcn6vr@Wh+&>`U$zy0_qYM%>!5sb?>)15}X1@59BKLrfx{-`mFQYh&vELd6pN z04(ajt7(1r=*RiX27{9eVC>g{u@BD%&_~`1j5z0Wp5=BCbIxBb@1MBkewtZ5U_9fU eeHe`Q&E7n-J=YL(278`>$DBOd?;C&Z5$p&06GGkq literal 0 HcmV?d00001 diff --git a/test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fshiftleft_int.spv64 b/test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fshiftleft_int.spv64 new file mode 100644 index 0000000000000000000000000000000000000000..6aea920b3618e8c1eddd0095d288ac762613f6ff GIT binary patch literal 856 zcma)*(N7aW5XJ{eTcxy6v4V;MAt50qKB|eSA(~K%Ca5I7Zb~WilH;yiuk{}!@n7PF z#P7H4CO*2C>CAjP-^{mr*UI8v-4+(itZuLM%WB)AEd^f=2+p-PiEAnvAxTCx&&jMf98Cvq-tMbgD->#;u$+HI@?PK86i}g+nYIuJv zIPW!ydcK1g(>KPP<4z16&$_BY#}gSnf2bc?*8?VJH>yF?QyF-noSII+j>KUe_stye zuJps%kb$EQJ(0Vqn0}G_Trs-9Blm@3@W}n6?(}>qV~;z&4H)l>S?t9bd<*~1c%>W- R;Ge`JhZp!0yK_7@pgMs5HA literal 0 HcmV?d00001 diff --git a/test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fshiftleft_uint.spv32 b/test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fshiftleft_uint.spv32 new file mode 100644 index 0000000000000000000000000000000000000000..d164e1baa29117c377c988d6dab9b6176eccd976 GIT binary patch literal 768 zcma)4T}#725ZpGl(bl%sTB{&f5d;x@R6#1DV2lcWfv-nwqYcEQq-p$vf`7>u!I|?8 zeDsdZ?#|8bPA+lE+pd&KBI3$Ao?LAy%bf8QsR)_3lQ22I*`MYkk)}8rD+6yy z6OHsg263v+Hj?cJ!)I=W=z;(3Vbe8Lwn&a_Q$JCP{_P}rv@4$%pP|2J) zwbfx8I{($Ws^f+}+eQ2z9lRS474z^#Tja{jq#n-fB42>-!?^ZLUPHVI>mjEOa;T|~ zd=Z|Dcn6vr@Wh+&>`U$zy0_qYM%>!5sb?>)15}X1@59BKLrfx{-`mFQYh&vELd6pN z04(ajt7(1r=*RiX27{9eVC>g{u@BD%&_~`1j5z0Wp5=BCbIxBb@1MBkewtZ5U_9fU eeHe`Q&E7n-J=YL(278`>$DBOd?;C&Z5$p&1yF%Xp literal 0 HcmV?d00001 diff --git a/test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fshiftleft_uint.spv64 b/test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fshiftleft_uint.spv64 new file mode 100644 index 0000000000000000000000000000000000000000..3ab2482228a39b0ceff19825977821755dec0feb GIT binary patch literal 856 zcma)*(N7aW5XJ{eTcxy6v4V;MAt50qKB|eSA(~K%Ca5I7Zb~WilH;yiuk{}!@n7PF z#P7H4CO*2C>CAjP-^{mr*UI8v-4+(itZuLM%WB)AEd^f=2+p-PiEAnvAxTCx&&jMf98Cvq-tMbgD->#;u$+HI@?PK86i}g+nYIuJv zIPW!ydcK1g(>KPP<4z16&$_BY#}gSnf2bc?*8?VJH>yF?QyF-noSII+j>KUe_stye zuJps%kb$EQJ(0Vqn0}G_Trs-9Blm@3@W}n6?(}>qV~;z&4H)l>S?t9bd<*~1c%>W- R;Ge`JhZp!0yK_7@vaMsENB literal 0 HcmV?d00001 diff --git a/test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fsub_int.spv32 b/test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fsub_int.spv32 new file mode 100644 index 0000000000000000000000000000000000000000..fd6c17cd3af6ccde1b7f46c0d0ec606953057f60 GIT binary patch literal 732 zcma)4O-sW-5S=Er(Z;scT6^(O5d<$DRgj7(7^8w;;B|>@w1G4U`KWh)wHLwn)+~5* zm&we$eKYftjcx5Y(rAc?BO^Sy`eMmS?VGi4!MoBFGVil6xx70p%9%)4Y>ipKd(y%d zRExYVf^x2=F+wo7Mb106w6UKbUexXNlS)z*C1v;?=ISF4GBpdQX&#hOng}Mjx<;uY zi}IzqsG_*^6E99*gV;}&=~NSaNC4N?def@ZF*V2YLdVpTBlf^>5bwi?`B2H6IJLR3 z9i9Jb-PUnOpB*6nj}CsDn|cm4aE^x#tMFA{WQemh#A6uO2zeWx7&VNMQwJCE0W_|| z6Ys#YC%L=m+=b^Eabu@-J$rE-p@JOV2N!FOF^PQMuZh{$#MJzSiXQv~Eb7>RvF9WD zaXwx}V{mc>jQuuY?8CDO^pQ7!5$Al)Gu$o2ob#7EL{8jrKkHficI?A*Q_s6$Kh~dt O$DBM<`;9;H1oi`cYC#16 literal 0 HcmV?d00001 diff --git a/test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fsub_int.spv64 b/test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fsub_int.spv64 new file mode 100644 index 0000000000000000000000000000000000000000..ded47aeae9e6cbcdac73ffe7592c90ab5c966ae7 GIT binary patch literal 820 zcma))-7iBy5XFb0sG`2V9*9KZ#Um1_G$g7>d?j8tRYjBQ-nzZ5cYn1PiSt``6OVQ? zoinp@X3p+yVR*Y}Lqld(v<>~TDmH8*xgX8_Snf;mB^xyu*WGq}dbQWfp3J7LAeMQX zu`%sm{QTC;UR^6vs0y}Otz584?faVtcX@v2;>300to_nX-ACH&x~Fz4Nt;aK5ilXq2j9f#C ziOHjl`|%ESMZ`-Ump$GU{*Mmc<|3bi2IkOjT)tm1Th!Tva81S{o|NAt&{|fE4r<(UD0OOs~gIU^tJDCg6QwJ!!Wd@mN915Kiy?;bFH1Lh82V$8}{ zRfP^+#p}COKeW!}IXQb*4Vva<@w1G4UNgD6|YA=HCty%Es zE|Zyg`)1}P8{67*q|p!&M@D#Z^~I8v+Ba+8f_J4UWZq|Ca(Q=Hlrxd8*c!8d_oRg_ zs1|u!1m#>!V}xLEi=20CX=6V>yr|pjCzYfsO3Ls(%+*I8WNH>p(>y4nG!aa4b&XO* z7Uh+?SVnQ_CtjSs2C<)1=~NSaNC4N?dedd8V``4)g^sBwN9=*&Al`=&^P!SCacXm6 zJ39Z>x~=1mK084CA07NQH}xE9;2aMfR^h9@$Pj01h{rIl5%M-XF=`kirw%US187`_ zC*FZ)PjYwBxeL!T;>J$vdiLTvLIpXz4=&amV-oqiUlX&hiK+Pu6+QR~Sk$orW6wwQ z<9xh|#^B@%82fF)*oS8m=p%0cBhLApXSiF4Ip;5Th@80Le%7=2?bwIsrk;1heyl$O Ok2!g!_8WiX3G4@c-a!Tc literal 0 HcmV?d00001 diff --git a/test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fsub_uint.spv64 b/test_conformance/spirv_new/spirv_bin/ext_cl_khr_spirv_no_integer_wrap_decoration_fsub_uint.spv64 new file mode 100644 index 0000000000000000000000000000000000000000..4ed0e5d5bbdc3b86a6063cbeea7a526b25342e8d GIT binary patch literal 820 zcma))-7iBy5XFb0sG`2V9*9KZ#Um1_G$g7>d?j8tRYjBQ-nzZ5cYn1PiSt``6OVQ? zoinp@X3p+yVR*Y}Lqld(v<>~TDmH8*xgX8_Snf;mB^xyu*WGq}dbQWfp3J7LAeMQX zu`%sm{QTC;UR^6vs0y}Otz584?faVtcX@v2;>300to_nX-ACH&x~Fz4Nt;aK5ilXq2j9f#C ziOHjl`|%ESMZ`-Ump$GU{*Mmc<|3bi2IkOjTz*h7Th!Tva81S{o|NAt&{|fE4r<(UD0OOs~gIU^tJDCg6QwJ!!Wd@mN915Kiy?;bFH1Lh82V$8}{ zRfP^+#p}COKeW!}IXQb*4Vva< +#include +#include + + +template +int test_ext_cl_khr_spirv_no_integer_wrap_decoration(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + const char *spvName, + const char *funcName, + const char *Tname) +{ + + cl_int err = CL_SUCCESS; + const int num = 10; + std::vector h_lhs(num); + std::vector h_rhs(num); + std::vector expected_results(num); + std::vector h_ref(num); + if (!is_extension_available(deviceID, "cl_khr_spirv_no_integer_wrap_decoration")) { + log_info("Extension cl_khr_spirv_no_integer_wrap_decoration not supported; skipping tests.\n"); + return 0; + } + + /*Test with some values that do not cause overflow*/ + if (std::is_signed::value == true) { + h_lhs.push_back((T)-25000); + h_lhs.push_back((T)-3333); + h_lhs.push_back((T)-7); + h_lhs.push_back((T)-1); + h_lhs.push_back(0); + h_lhs.push_back(1); + h_lhs.push_back(1024); + h_lhs.push_back(2048); + h_lhs.push_back(4094); + h_lhs.push_back(10000); + } else { + h_lhs.push_back(0); + h_lhs.push_back(1); + h_lhs.push_back(3); + h_lhs.push_back(5); + h_lhs.push_back(10); + h_lhs.push_back(100); + h_lhs.push_back(1024); + h_lhs.push_back(2048); + h_lhs.push_back(4094); + h_lhs.push_back(52888); + } + + h_rhs.push_back(0); + h_rhs.push_back(1); + h_rhs.push_back(2); + h_rhs.push_back(3); + h_rhs.push_back(4); + h_rhs.push_back(5); + h_rhs.push_back(6); + h_rhs.push_back(7); + h_rhs.push_back(8); + h_rhs.push_back(9); + size_t bytes = num * sizeof(T); + + clMemWrapper lhs = clCreateBuffer(context, CL_MEM_READ_ONLY, bytes, NULL, &err); + SPIRV_CHECK_ERROR(err, "Failed to create lhs buffer"); + + err = clEnqueueWriteBuffer(queue, lhs, CL_TRUE, 0, bytes, &h_lhs[0], 0, NULL, NULL); + SPIRV_CHECK_ERROR(err, "Failed to copy to lhs buffer"); + + clMemWrapper rhs = clCreateBuffer(context, CL_MEM_READ_ONLY, bytes, NULL, &err); + SPIRV_CHECK_ERROR(err, "Failed to create rhs buffer"); + + err = clEnqueueWriteBuffer(queue, rhs, CL_TRUE, 0, bytes, &h_rhs[0], 0, NULL, NULL); + SPIRV_CHECK_ERROR(err, "Failed to copy to rhs buffer"); + + std::string kernelStr; + + { + std::stringstream kernelStream; + kernelStream << "#define spirv_fadd(a, b) (a) + (b) \n"; + kernelStream << "#define spirv_fsub(a, b) (a) - (b) \n"; + kernelStream << "#define spirv_fmul(a, b) (a) * (b) \n"; + kernelStream << "#define spirv_fshiftleft(a, b) (a) << (b) \n"; + kernelStream << "#define spirv_fnegate(a, b) (-a) \n"; + + kernelStream << "#define T " << Tname << "\n"; + kernelStream << "#define FUNC spirv_" << funcName << "\n"; + kernelStream << "__kernel void fmath_cl(__global T *out, \n"; + kernelStream << "const __global T *lhs, const __global T *rhs) \n"; + kernelStream << "{ \n"; + kernelStream << " int id = get_global_id(0); \n"; + kernelStream << " out[id] = FUNC(lhs[id], rhs[id]); \n"; + kernelStream << "} \n"; + kernelStr = kernelStream.str(); + } + + size_t kernelLen = kernelStr.size(); + const char *kernelBuf = kernelStr.c_str(); + + for (int i = 0; i < num; i++) { + if (std::string(funcName) == std::string("fadd")) { + expected_results[i] = h_lhs[i] + h_rhs[i]; + } else if (std::string(funcName) == std::string("fsub")) { + expected_results[i] = h_lhs[i] - h_rhs[i]; + } else if (std::string(funcName) == std::string("fmul")) { + expected_results[i] = h_lhs[i] * h_rhs[i]; + } else if (std::string(funcName) == std::string("fshiftleft")) { + expected_results[i] = h_lhs[i] << h_rhs[i]; + } else if (std::string(funcName) == std::string("fnegate")) { + expected_results[i] = 0 - h_lhs[i]; + } + } + + { + // Run the cl kernel for reference results + clProgramWrapper prog; + err = create_single_kernel_helper_create_program(context, &prog, 1, &kernelBuf, NULL); + SPIRV_CHECK_ERROR(err, "Failed to create cl program"); + + err = clBuildProgram(prog, 1, &deviceID, NULL, NULL, NULL); + SPIRV_CHECK_ERROR(err, "Failed to build program"); + + clKernelWrapper kernel = clCreateKernel(prog, "fmath_cl", &err); + SPIRV_CHECK_ERROR(err, "Failed to create cl kernel"); + + clMemWrapper ref = clCreateBuffer(context, CL_MEM_READ_WRITE, bytes, NULL, &err); + SPIRV_CHECK_ERROR(err, "Failed to create ref buffer"); + + err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &ref); + SPIRV_CHECK_ERROR(err, "Failed to set arg 0"); + + err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &lhs); + SPIRV_CHECK_ERROR(err, "Failed to set arg 1"); + + err = clSetKernelArg(kernel, 2, sizeof(cl_mem), &rhs); + SPIRV_CHECK_ERROR(err, "Failed to set arg 2"); + + size_t global = num; + err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0, NULL, NULL); + SPIRV_CHECK_ERROR(err, "Failed to enqueue cl kernel"); + + err = clEnqueueReadBuffer(queue, ref, CL_TRUE, 0, bytes, &h_ref[0], 0, NULL, NULL); + SPIRV_CHECK_ERROR(err, "Failed to read from ref"); + } + + for (int i = 0; i < num; i++) { + if (expected_results[i] != h_ref[i]) { + log_error("Values do not match at index %d expected = %d got = %d\n", i, expected_results[i], h_ref[i]); + return -1; + } + } + + clProgramWrapper prog; + err = get_program_with_il(prog, deviceID, context, spvName); + SPIRV_CHECK_ERROR(err, "Failed to build program"); + + clKernelWrapper kernel = clCreateKernel(prog, "fmath_cl", &err); + SPIRV_CHECK_ERROR(err, "Failed to create spv kernel"); + + clMemWrapper res = clCreateBuffer(context, CL_MEM_READ_WRITE, bytes, NULL, &err); + SPIRV_CHECK_ERROR(err, "Failed to create res buffer"); + + err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &res); + SPIRV_CHECK_ERROR(err, "Failed to set arg 0"); + + err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &lhs); + SPIRV_CHECK_ERROR(err, "Failed to set arg 1"); + + err = clSetKernelArg(kernel, 2, sizeof(cl_mem), &rhs); + SPIRV_CHECK_ERROR(err, "Failed to set arg 2"); + + size_t global = num; + err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0, NULL, NULL); + SPIRV_CHECK_ERROR(err, "Failed to enqueue cl kernel"); + + std::vector h_res(num); + err = clEnqueueReadBuffer(queue, res, CL_TRUE, 0, bytes, &h_res[0], 0, NULL, NULL); + SPIRV_CHECK_ERROR(err, "Failed to read from ref"); + + for (int i = 0; i < num; i++) { + if (expected_results[i] != h_res[i]) { + log_error("Values do not match at location %d expected = %d got = %d\n", i, expected_results[i], h_res[i]); + return -1; + } + } + + return 0; +} + +#define TEST_FMATH_FUNC(TYPE, FUNC) \ + TEST_SPIRV_FUNC(ext_cl_khr_spirv_no_integer_wrap_decoration_##FUNC##_##TYPE) \ + { \ + return test_ext_cl_khr_spirv_no_integer_wrap_decoration(deviceID, context, queue, \ + "ext_cl_khr_spirv_no_integer_wrap_decoration_"#FUNC"_"#TYPE, \ + #FUNC, \ + #TYPE \ + ); \ + } + +TEST_FMATH_FUNC(int, fadd) +TEST_FMATH_FUNC(int, fsub) +TEST_FMATH_FUNC(int, fmul) +TEST_FMATH_FUNC(int, fshiftleft) +TEST_FMATH_FUNC(int, fnegate) +TEST_FMATH_FUNC(uint, fadd) +TEST_FMATH_FUNC(uint, fsub) +TEST_FMATH_FUNC(uint, fmul) +TEST_FMATH_FUNC(uint, fshiftleft) \ No newline at end of file