From 4e1282b4ba46a059f20ad6e15c7b18565ad7d6c9 Mon Sep 17 00:00:00 2001 From: castano Date: Fri, 16 Nov 2007 11:05:17 +0000 Subject: [PATCH] Add nvassemble tool to create cubemaps, volumes, and texture arrays. --- trunk/project/vc8/nvassemble/nvassemble.rc | Bin 0 -> 3130 bytes .../project/vc8/nvassemble/nvassemble.vcproj | 353 ++++++++++++++++++ trunk/project/vc8/nvassemble/nvidia.ico | Bin 0 -> 15086 bytes trunk/project/vc8/nvassemble/resource.h | 16 + trunk/project/vc8/nvddsinfo/nvddsinfo.vcproj | 149 ++++---- trunk/project/vc8/nvtt.sln | 43 ++- trunk/src/nvimage/DirectDrawSurface.cpp | 2 +- trunk/src/nvimage/DirectDrawSurface.h | 5 +- trunk/src/nvtt/cuda/CompressKernel.cu | 2 + trunk/src/nvtt/cuda/CudaCompressDXT.cpp | 6 +- trunk/src/nvtt/tools/assemble.cpp | 65 +++- 11 files changed, 542 insertions(+), 99 deletions(-) create mode 100644 trunk/project/vc8/nvassemble/nvassemble.rc create mode 100644 trunk/project/vc8/nvassemble/nvassemble.vcproj create mode 100644 trunk/project/vc8/nvassemble/nvidia.ico create mode 100644 trunk/project/vc8/nvassemble/resource.h diff --git a/trunk/project/vc8/nvassemble/nvassemble.rc b/trunk/project/vc8/nvassemble/nvassemble.rc new file mode 100644 index 0000000000000000000000000000000000000000..842ded193e7d9fc52ad5c664186995b6ea53e2a2 GIT binary patch literal 3130 zcmds(+l~@J5Qgj8#CK@f8#fwI_X!qQk;tN1VAX_#5O7=)PT{a>vd^x5f6p`yFkpf= zqv_1_bahpC9sa7CpFc`gGH2H|vVkpZW=o#PR?u#tEp2Um8?&m|GlsJ-_SvR93+OB6 zA+!av}xXJ6WT zYgyYmR<*jlv!31971S!5XEC@tRCR)*5xJKQfF@;8EZVRlx-FS|aD@Y2jC*uTO% z@tawJlV>mA@Y~@8Vto5)FN9t=CfvqcZ3ARaQEIAFWNth7iH5eo_LF1)qk#XWY>g%AkBJS;-73z z=l9FlUh@$Xw@XHLo$lI;q*g2bvB$za7bo7+r4Dt$dfWrH&)tr4teaw;mHYd#44UOD z!_(SGL#@N?h{J%&KiyE$HvHao_7&u-H#jQ9jFoz_`gonFIJ@_KeC<2O9FD&4)au^q ztc!$qg{1R&q|UD{;^3S!RCNZv=a0NL&acyY#a(9tk2?Gk$yo0tYiXjMCVeEMQ^ySb zDSS_MC(Lo5lE!tiw@ximPZwT~jlg$|C%sodgL|}6t>|;7ESn{0JB+VWs&6{7`T7nh s3bX#pBK-2dR>r^kr&1+5{!bPBn6D_OwULI0(LE&tl%4y3P}x=V7etkE?*IS* literal 0 HcmV?d00001 diff --git a/trunk/project/vc8/nvassemble/nvassemble.vcproj b/trunk/project/vc8/nvassemble/nvassemble.vcproj new file mode 100644 index 0000000..b953801 --- /dev/null +++ b/trunk/project/vc8/nvassemble/nvassemble.vcproj @@ -0,0 +1,353 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + diff --git a/trunk/project/vc8/nvassemble/nvidia.ico b/trunk/project/vc8/nvassemble/nvidia.ico new file mode 100644 index 0000000000000000000000000000000000000000..6aa721e600b7daf49a87573a5505b6032edc6c42 GIT binary patch literal 15086 zcmeI3ca+sdw(rmBoO9EG?$e2erV)@F1w@i4QBaZyq8I=JO-2#LfB{4W(})5}xJ1Q4 zQ9v042?lTg1rsU`i7TB z@4D-*3-<5dKe1=ep4nZxbSW7$Xi%MNuDPb=)mLBLsZXCiy|23Js%x*j^2%X)U(vmL z_i^32b(_?+Yu9@&yX><2I(6zawSD{c4|VL=@nJyE4_$icr4I`3@7%fbl*=!_{4Vt$ zf5jD7RP^fAYq)R?7S3Lx*HLtu3van-hSZ)mWXKTB%75=C93`ULQ+!m2#^a*yLfRWb@e^>UieRgLoyy0 zZgBoxbYTJa};A?@4mU9d~5(=+R?Ny?XVI)T&j>fb#Nk zt6jUcgF1EU*o7BfXm#t>HD9N!U%$RJY0|_RHELuH8#c5SEm~O1mM!g~i!M3`ty{OY zX3d(pK24i8btU~9H*V}eG#zg6K@U2fE#*RwQ6Mmyi>wIvb1T_#^9fJp1oYu8ISB<4xZ6L zWj(<^yc87`Ih{gpz<>c8o__l2bMr7|%9I+C>t}RcqecxUFEW>vm07iF)vTnX#Es9( z%d@Igt6FYuu4QLuTa_wR9AsoFEYA#shf3 z2VNjQKi>)p3Jg6$8=l~WvU>IECVRHAW5>4ib8zdew+==>^ksZvVxq;z$6IV{>^a~$ zE-ubNLPEm1GASv^^@neG_nWU;c0j>k=D?jEr=;XS^TlaZ!Sc68g~f^e6$^eh)mql)kU;@A2^tnmqH4 zXX><}OZdm;SdYDX_g;JBjW-76Ys2JsHqu5v_^0%ED$}J7ZmI?fxhNs7Kbab?9XD!GN?;bsRbo9uPBV**FKcNl#_O?U+%Cx}a`C^`+ z506<{Sq@l(JTrzi)}t?}`}K(p(+^+_xcoC?J>JS~$i}?;HNhIzp77ThGiGEy{`ljW z@;!gWM@2mR#FE3 zWAUwwK?dsVMerF4gkj_V;jCG+9A5S}*^1*2{?VuQ!PCQr4ZG^DyYA|(J?#{2@XT8B{If^C_10T%k7Zuq z4L%=x>@kOneILH2PoHipR;+M)3^+8mwqe5tM@KQ5!{K%Rzyl9B*{~DEXzm@j9^j8SL;h#Vs8OTtrcN90=*N6Q3p+S?@SvL$ z?G>&KyI8(_xm!!{`26$FyK(d9&v$$x7qRA#KmO=&YM-$me)z%J8#wSkhYufi@sMIJ z+rNLmTPM(d~y<2bK#izkLGO$+Q75^sQSIgdBBX*#TcV14$F&DIX06)Ca2N`Jx zH~xY-=DA4zwW|ClG8PmTSawdftMklz=08q#_8Z31j*l%YE^>aCKD0p#ocI-Z0icV2 zfA|AG>%z;;UwCuQ+B)>vN4yL?!;9B5Jj;J=*V*+$ z>d4JI&p-Tm9G++RrVhZfheFj0^FtP!mSW-YF&32|Uz;e17bIxT5-UMG^|S9O)i+%I zB2yI~XXjYxtbY7i!E(_auzXdi4)=sj(@-a@9u$OMPpm+~10B#fL8UB}zv0U+AaF%~` zlDO24;SV8=9z zlRYGJ~+nqt8( zpR$yD2U5A4i!7??r`HkezEu z(rs|Wr55-E!F7FjLNB68C$Qa&h2de6+) zp0HB7{>MvN(_<}ID`bhXuLAkcnsI3s63ys7tu3(bZVSA1w`CWGtSBSZ0OsgIlZ80y8w}kf}wWJSbSmL2Y7Q6Q;3)U;JD%gK{bxY9Lq{GWCeA1N` zOwk$bp1zi_^9f7b@`%NJG{dq>wU4A#`HlXW2lNSe{yFbz{qB?h{DirHf9(H!{tGpy z>5);EG=H?^9b0WN<9b}NSg*I7g(<;KDyv_7@xF2B=?eqL+o(}rr!immv^ zH!bhbN((ktY@j|x1%;NgZ?+X3dD$|zPP5RyxmNX~CoSi_$1Q9BJn^qFsq&$UNBn-{ zZNT&IV>tMi|NLC%_#OW2UYUPqd+dc#NtWBWg*80!xfSfyddx_)G_Av08=kPnY71*u z%Yw3l2*v$%|FO#I{J75Y(=x4EQ2sSZYxTQDR_o|iOKw%i$|QHBY%YJ#bJpOeRaWcM zd3M3)FIw#{7g)`suUc-65-Uv2JjXx#d*wBN44!|^p4xl&DDK@voQnRL8!snvql2p2 zQ*x_hS;NgwTdz|GEq3-W3&clTf_QE?bCh*H`i&KCoMGX^I$M=~?JPSc+L9$-AVTNZ z#(CEH_}A9q;2JB4OK~y=^AfG@7xS(4sg2g`$SQ00(`(l9mpARA6Yp7eS+Rw*w$VQ^ zAv}2geTj_|KftqMpN|y7y^buzZ|J`=|BNY+&kxBaGqq>bz56DceCm+3_;{fO3KC2g zBbJ^RZ=L2%wa&X%SRf_Nf(bDe-n__Ktes`EPyJ-QHqN(1* z!)=G%Xdl^N`_Klee*Niht!j;GmajdXTrITSJ})^u@5tl%=RB&j^}C9pU&Ma=x&3_p z!Qu2#5Yn2BvLe~_^zT2mnMeL^HRevRu#n`Ct^?sg8+Bltz4r4VTX*6oTX66@tKXuz zh3U){h>5WH@x5&E@x!*sPS_&*%9hx-w$S$5OAZd$!c#}BYN_~^?Ujf(^pE|VzYg&6 z=$iOjvGeOXe@%yfVz z%`_|R*4A3LX=|}bl1cKwcXaQDHgxwN?Nd8$@7d4xM>}Az+u!Up`%dL|_Nriw=40Kd zW0qB{xJ~v-t_6C~{XPDPld&J2313nkdKCE_Vh7?M_-3Av9U$KW9&iF>XYUenlF}?$ zYof z#?Cx22473;N6b-MP^+|t6{csoIhlHJx41lIpV+Umkkj^=>YG(xCs;3dTed@a22R+Y zR9~yMcT~Q1`j}-GB%5qHAnM`HztkCQN9d^ZNyFY2##9As(RHwAMb)EfyjoZ4)ZrQNVMr>GMx2#|2pkn=h*vR!u zY}AIOHu@dGMqs&(-?+@is;t<&(uQnVX4Pw#Tivp9t5xR$m$O4}&_|E(1pnw585HxJ zmhJRbZlOT=pF^w*`1kr}4En~lI3L3^c_MOhbn}bO z@QvOdVb|~o|I*`e621poAB9OM1{9QRJ}3XmTmH>4b}%T;mT$h~4$ z<=VO`_tT?$PwT0CT6e(}y?VJi`85ySyLIR6=f;rhx=eVq7dgD>qIIh?_kGZ2&e0=0 z!9TtS{`KzQJ@?#`r897t;-EjlKkI_{lX)X1C9ehF+`;Ky)A0xFb0$RwyT^}~eTEkyrL!N)x z{4V9jqPT<6+&_c;AQ$?_Mz9UV7fvVOBo{o4 zTJp3v-+Z$v-|9;GkbCn$ZVS2ko|OYxm>=c@S(tO=fH&-iwJM#jfOmeglM_qzNl`7m;S*em+ye2)K-d^exrJ`j+8 zBBa~7$b}to&cxnWv&#x7=`p=y^*YQqnTz;H8Q8!o6(KiK7kK~Y9hu(jp zf7$kC*|EO(_}5Q4{NOXsJkv|E%$D73k==hI{*TD^PT*@<&-gOpYv$1d&*%mGmG?Q; z95%~Zq7E<~ppAK^KYF4Lus+c*xS@G=uQ@4SbBwc!?0AiGrsFg}>Hp&Vw>9O}Bb291 z`SYLuT&!Gk1L^q^*>Nx361YR*7tJR#iHA8yKSy)^jX~nJ^t{+4?m#!oG2Q@ zwRZa7efQm$%J!OS9;=C;}<4)Oe7mAw1^ zKc|7o{QVt0t*2f3{DrbWLHSv|^PN_qpp78tYSQ(;t!%ycAM}%7&(wa^Q~q$Y?s*=S zub(4Wp_pKe_KJ7)ZoA5zdfub;{+Zr=srpy)p{I;Hd2h zVlsg7-~k^o8~#rHevogt%rl@E;Z^S!@q_q206&QT1n_}AR`j`4O5zg#T-o>f{)_>S z2Yf8?v-bOU4T4_?L-Ke2yPa>j;to+Xv#yJ+lNcgc?_?wQRV z4}HcWH?bP)h_S>n@JDRT-8*-;(3jn~IF&d|IwWUfOP4NnYmJzXwZ^vtydy>=#zO|L zQ`U#}>G_(26~rp&i#bPr?xE?+Hzq4ru5>eZ`FF@?>MUWgTm z57Eo4C!chEYsq&VIB>x6k1zZB>#yBdd>ZxwEp&{XC|>!K=4CT=+K?F<=nNSE?3A+< z@-n`?;`Wgh6^dC0S+^eDbys+mEm^!ovGqc`uzo#aIB|MspNv5iM$2^z`d2X2x|E$=%o12QBcy!xq?Zue)mr{89HB+-+=mz|t;j zVny-k?%vVk13z|#{&fcUSnGK&@?)>O1K`!&4Hat-(!Fwm?qY+Zx>(pE-64+XVo|MX zTVykRSC*ykKQeR|-M*0}&KPRJS9J%sYJvsl>E2F!q;zUxC5aivT>&zIm(u%0`H#Ii z|9uX9c!D;30wMN4<#yvUQ{0_c!jyrQDq4|kB>#~17CpR+g>ygogmj@k`4Lf8Tvlo+ zW3RHH?oqh6N$S$vN)pb{cQT!|0q+w<|5MTb6#DQ4efSF%=**H@#S+qzEU;^uC5`T7 zDXIr0Q}iU=(_P=u!n-!M_$u0uZs=@5<==xhb+R=1)q;R>i}{%rcyFpD^=)e%WuENSEw7FU>K;quM#7uB|Ihxb@9m4~Fyx5y!!R^|QHdm(4d-_+6i`QLxu?_iRgLQiQS&Poo{Xk89 zpV7aawSVY#yMD(?8~gb-OP21#^o>HFt;!+TVSCsq5z=MB}-zk*838m}%=j0UYdD{(k%jP9^#frzQ?2bX!zG8@#wrXTiT7Omh zTx_#{{Kj@EU%Ezl-#3&4UM+ZEdF?K@U9WFP;?JF{oqnYU}`;^PwE!becC~y0Va=5^8 z*LL~nn=CUS(MlEbI+@PKvwr<6PIyJ}*EIH1=zE{beouVj^RkVUw<4CVTdZ7Vi>5Yx z-!_}^=M6Sx<4PO7W}%Hd6@90N7M$AucioKm!pM9?rxx}{ZtVvmIYg)UGHNBv&H3RC@bI`a}9cxrx z+ZxoYW%Y>jn>5459WG=6H#D$4=0G`u35vB^(T4v==k!BX=*m9b|5NbY8}Scu#8rK+ z5}e5m_Eo-afZ~mQefwK~)rk}O_3dX@svr4LVl>6}=kjdii=n~(ho2F?uKLD2UGx2S z_DJ%B(!j(nKT5bhgHa#pnS{=WO~cd-faBJmk` z$QhCgmCtc;oqX*v#aG;l1xfmCFg8xayIppbx?TCE< zzF8wyAzw!R4<3~NcQO!*lAA@J%qw!U&ZL_M{rQVh<1bVUJ5)Ku`^o)iUY3&I;S4EW zc5)}Hx%^h}H=W-OYVY__a9HpYpxE$`XdKkLzsX+yqH&x>cd6W_9MM|kk(Malep>pP zC?4BYrhR^W$BrG*iaRs)y=bb3y?gg|(D#|$wr<_pM|r=&$_EbBH>x)YhU?q9oArEy z+OO67e)`t2$5&r{)xnSTAiYJmI^lgpq>mvrWv*=FsEvj2ni*2p{^HA(lWk_cW+x(E?rbjKUMIB4h z_`3BjJW*b&{A5a4l%)l9|N6;emQbU>qVK%Y0ylK9V82!tb*0WwZ%@+ufMp2+%PTBd zdt18p`LgnwC(EkWIF+n=><81uS3OHxHQAInwa~sfmif89J0H@)aukQeEA9#0b)_X7 zT3}Urw6auLPK_+JqPJy^>0_~1C~g;?_O{2J;Q|-cuvYfH zHP8+gy1JvjfjTq3W_kIq$q~_(BmM_$oNE_9aJ%Iy9&0vguwAianj2eubqCw1^|VB{ zoo}jaIk>llX}rdsRKEP@V@7u##INF)r`$Wm9+~!t`&NFc{PRSe9oaj2>l=9JVCV9I z_v;(zSlwT=)%Wxd=#FKD{?BZK?k7H!zx`Ny$LHElzEuAGQ$6pJf83_Kn>Tb{^R)gi SpzA;X`OnnJ&IA0P|M^e08BNFl literal 0 HcmV?d00001 diff --git a/trunk/project/vc8/nvassemble/resource.h b/trunk/project/vc8/nvassemble/resource.h new file mode 100644 index 0000000..e765787 --- /dev/null +++ b/trunk/project/vc8/nvassemble/resource.h @@ -0,0 +1,16 @@ +//{{NO_DEPENDENCIES}} +// Microsoft Visual C++ generated include file. +// Used by nvcompress.rc +// +#define IDI_ICON1 101 + +// Next default values for new objects +// +#ifdef APSTUDIO_INVOKED +#ifndef APSTUDIO_READONLY_SYMBOLS +#define _APS_NEXT_RESOURCE_VALUE 102 +#define _APS_NEXT_COMMAND_VALUE 40001 +#define _APS_NEXT_CONTROL_VALUE 1000 +#define _APS_NEXT_SYMED_VALUE 101 +#endif +#endif diff --git a/trunk/project/vc8/nvddsinfo/nvddsinfo.vcproj b/trunk/project/vc8/nvddsinfo/nvddsinfo.vcproj index 7528865..df62cbb 100644 --- a/trunk/project/vc8/nvddsinfo/nvddsinfo.vcproj +++ b/trunk/project/vc8/nvddsinfo/nvddsinfo.vcproj @@ -43,7 +43,11 @@ Optimization="0" AdditionalIncludeDirectories="..;..\..\..\src;..\..\..\gnuwin32\include" PreprocessorDefinitions="WIN32;_DEBUG;_CONSOLE;NVTT_SHARED;__SSE2__;__SSE__;__MMX__" + BasicRuntimeChecks="3" RuntimeLibrary="3" + WarningLevel="3" + Detect64BitPortabilityProblems="true" + DebugInformationFormat="4" /> + + + + + + + + + + + + + + + + + + + + @@ -158,78 +235,6 @@ Name="VCPostBuildEventTool" /> - - - - - - - - - - - - - - - - - - - - +#include namespace nv { @@ -102,10 +102,11 @@ namespace nv bool hasDX10Header() const; }; + NVIMAGE_API Stream & operator<< (Stream & s, DDSHeader & header); /// DirectDraw Surface. (DDS) - class DirectDrawSurface + class NVIMAGE_CLASS DirectDrawSurface { public: DirectDrawSurface(const char * file); diff --git a/trunk/src/nvtt/cuda/CompressKernel.cu b/trunk/src/nvtt/cuda/CompressKernel.cu index 784a7bf..4db1141 100644 --- a/trunk/src/nvtt/cuda/CompressKernel.cu +++ b/trunk/src/nvtt/cuda/CompressKernel.cu @@ -294,6 +294,7 @@ __device__ float evalPermutation4(const float3 * colors, float3 color_sum, uint uint akku = 0; // Compute alpha & beta for this permutation. + #pragma unroll for (int i = 0; i < 16; i++) { const uint bits = permutation >> (2*i); @@ -329,6 +330,7 @@ __device__ float evalPermutation3(const float3 * colors, float3 color_sum, uint uint akku = 0; // Compute alpha & beta for this permutation. + #pragma unroll for (int i = 0; i < 16; i++) { const uint bits = permutation >> (2*i); diff --git a/trunk/src/nvtt/cuda/CudaCompressDXT.cpp b/trunk/src/nvtt/cuda/CudaCompressDXT.cpp index d49633e..45bc395 100644 --- a/trunk/src/nvtt/cuda/CudaCompressDXT.cpp +++ b/trunk/src/nvtt/cuda/CudaCompressDXT.cpp @@ -109,12 +109,14 @@ void nv::cudaCompressDXT1(const Image * image, const OutputOptions & outputOptio uint imageSize = w * h * 16 * sizeof(Color32); uint * blockLinearImage = (uint *) malloc(imageSize); - convertToBlockLinear(image, blockLinearImage); + convertToBlockLinear(image, blockLinearImage); // @@ Do this on the GPU! const uint blockNum = w * h; const uint compressedSize = blockNum * 8; const uint blockMax = 32768; // 49152, 65535 + clock_t start = clock(); + // Allocate image in device memory. uint * d_data = NULL; cudaMalloc((void**) &d_data, min(imageSize, blockMax * 64U)); @@ -125,8 +127,6 @@ void nv::cudaCompressDXT1(const Image * image, const OutputOptions & outputOptio setupCompressKernel(compressionOptions.colorWeight.ptr()); - clock_t start = clock(); - // TODO: Add support for multiple GPUs. uint bn = 0; while(bn != blockNum) diff --git a/trunk/src/nvtt/tools/assemble.cpp b/trunk/src/nvtt/tools/assemble.cpp index b8dea45..0e91311 100644 --- a/trunk/src/nvtt/tools/assemble.cpp +++ b/trunk/src/nvtt/tools/assemble.cpp @@ -24,17 +24,17 @@ #include #include -#include -#include +#include +#include #include +#include #include "cmdline.h" // @@ Add decent error messages. // @@ Add option to resize images. -// @@ Output DDS header according to surface type. -// @@ Output images. +// @@ Add support for reading DDS files with 2D images and possibly mipmaps. int main(int argc, char *argv[]) { @@ -46,6 +46,7 @@ int main(int argc, char *argv[]) bool assembleTextureArray = false; nv::Array files; + nv::Path output = "output.dds"; // Parse arguments. for (int i = 1; i < argc; i++) @@ -69,7 +70,14 @@ int main(int argc, char *argv[]) assembleVolume = false; assembleTextureArray = true; }*/ - + else if (strcmp("-o", argv[i]) == 0) + { + i++; + if (i < argc && argv[i][0] != '-') + { + output = argv[i]; + } + } else if (argv[i][0] != '-') { files.append(argv[i]); @@ -83,7 +91,13 @@ int main(int argc, char *argv[]) return 1; } - if (assembleCubeMap && files.count() != 0) + if (nv::strCaseCmp(output.extension(), ".dds") != 0) + { + //output.stripExtension(); + output.append(".dds"); + } + + if (assembleCubeMap && files.count() != 6) { printf("*** error, 6 files expected, but got %d\n", files.count()); return 1; @@ -96,6 +110,8 @@ int main(int argc, char *argv[]) bool hasAlpha = false; const uint imageCount = files.count(); + images.resize(imageCount); + for (uint i = 0; i < imageCount; i++) { if (!images[i].load(files[i])) @@ -130,9 +146,40 @@ int main(int argc, char *argv[]) return 1; } - // @@ Output DDS header. - - // @@ Output images. + // Output DDS header. + nv::DDSHeader header; + header.setWidth(w); + header.setHeight(h); + + if (assembleCubeMap) + { + header.setTextureCube(); + } + else if (assembleVolume) + { + header.setTexture3D(); + header.setDepth(imageCount); + } + else if (assembleTextureArray) + { + //header.setTextureArray(imageCount); + } + + // @@ It always outputs 32 bpp. + header.setPitch(4 * w); + header.setPixelFormat(32, 0xFF0000, 0xFF00, 0xFF, hasAlpha ? 0xFF000000 : 0); + + stream << header; + + // Output images. + for (uint i = 0; i < imageCount; i++) + { + const uint pixelCount = w * h; + for (uint p = 0; p < pixelCount; p++) + { + stream << images[i].pixel(p).b << images[i].pixel(p).g << images[i].pixel(p).r << images[i].pixel(p).a; + } + } return 0; }