From 63e3eeba794b5b271c9b33f7edb033e24ac78adb Mon Sep 17 00:00:00 2001 From: unknown <> Date: Thu, 13 Jul 2023 10:34:54 -0400 Subject: [PATCH] JNI attempt 1 --- dependency-reduced-pom.xml | 28 +- pom.xml | 27 +- shared-folder/libfluidsim.dll | Bin 0 -> 72666 bytes src/main/c/compile.sh | 62 +++ src/main/c/fluidsim.c | 237 ++++++++++ src/main/c/includes/electrosphere_FluidSim.h | 31 ++ src/main/java/electrosphere/FluidSim.java | 128 ++---- src/main/java/electrosphere/Main.java | 46 +- src/main/java/electrosphere/MyTest.java | 89 ---- .../java/electrosphere/opencl/CLBuffer.java | 35 -- .../electrosphere/opencl/CLCommandQueue.java | 434 ------------------ .../java/electrosphere/opencl/CLContext.java | 164 ------- .../java/electrosphere/opencl/CLDevice.java | 68 --- .../java/electrosphere/opencl/CLEvent.java | 65 --- .../java/electrosphere/opencl/CLImage.java | 218 --------- .../java/electrosphere/opencl/CLKernel.java | 32 -- .../java/electrosphere/opencl/CLPlatform.java | 153 ------ .../java/electrosphere/opencl/CLProgram.java | 70 --- .../opencl/CLReadImageResult.java | 41 -- .../electrosphere/opencl/CLUtilities.java | 181 -------- .../electrosphere/opencl/LWJGLContext.java | 28 -- .../electrosphere/render/GLFWContext.java | 21 +- .../java/electrosphere/render/GridCell.java | 19 + src/main/java/electrosphere/render/Mesh.java | 164 +++---- .../render/TerrainChunkData.java | 15 + .../java/electrosphere/render/Triangle.java | 11 + src/main/resources/shader.fs | 15 - src/main/resources/shader.vs | 4 - 28 files changed, 549 insertions(+), 1837 deletions(-) create mode 100644 shared-folder/libfluidsim.dll create mode 100644 src/main/c/compile.sh create mode 100644 src/main/c/fluidsim.c create mode 100644 src/main/c/includes/electrosphere_FluidSim.h delete mode 100644 src/main/java/electrosphere/MyTest.java delete mode 100644 src/main/java/electrosphere/opencl/CLBuffer.java delete mode 100644 src/main/java/electrosphere/opencl/CLCommandQueue.java delete mode 100644 src/main/java/electrosphere/opencl/CLContext.java delete mode 100644 src/main/java/electrosphere/opencl/CLDevice.java delete mode 100644 src/main/java/electrosphere/opencl/CLEvent.java delete mode 100644 src/main/java/electrosphere/opencl/CLImage.java delete mode 100644 src/main/java/electrosphere/opencl/CLKernel.java delete mode 100644 src/main/java/electrosphere/opencl/CLPlatform.java delete mode 100644 src/main/java/electrosphere/opencl/CLProgram.java delete mode 100644 src/main/java/electrosphere/opencl/CLReadImageResult.java delete mode 100644 src/main/java/electrosphere/opencl/CLUtilities.java delete mode 100644 src/main/java/electrosphere/opencl/LWJGLContext.java create mode 100644 src/main/java/electrosphere/render/GridCell.java create mode 100644 src/main/java/electrosphere/render/TerrainChunkData.java create mode 100644 src/main/java/electrosphere/render/Triangle.java diff --git a/dependency-reduced-pom.xml b/dependency-reduced-pom.xml index 806f758..765e648 100644 --- a/dependency-reduced-pom.xml +++ b/dependency-reduced-pom.xml @@ -2,7 +2,7 @@ 4.0.0 electrosphere - sample-cli-java + fluid-sim 1.0-SNAPSHOT @@ -55,11 +55,31 @@ 17 17 - --add-modules - jdk.incubator.vector + -h + src/main/c/includes + + org.codehaus.mojo + exec-maven-plugin + 3.1.0 + + + compile-c-code + generate-resources + + exec + + + bash + + ./src/main/c/compile.sh + + + + + @@ -99,6 +119,6 @@ 17 UTF-8 1.10.5 - 3.3.2 + 3.2.3 diff --git a/pom.xml b/pom.xml index b1e600d..a923d85 100644 --- a/pom.xml +++ b/pom.xml @@ -9,7 +9,7 @@ UTF-8 17 17 - 3.3.2 + 3.2.3 1.10.5 @@ -133,7 +133,6 @@ - maven-assembly-plugin 3.0.0 @@ -164,11 +163,31 @@ 17 17 - --add-modules - jdk.incubator.vector + -h + src/main/c/includes + + org.codehaus.mojo + exec-maven-plugin + 3.1.0 + + + compile-c-code + generate-resources + + exec + + + bash + + ./src/main/c/compile.sh + + + + + diff --git a/shared-folder/libfluidsim.dll b/shared-folder/libfluidsim.dll new file mode 100644 index 0000000000000000000000000000000000000000..2fbdd2c431177fe768937fbeccd894c50452d0b5 GIT binary patch literal 72666 zcmeFa3w%^XmOp+w-5ok4;WiR5s7Rwh6XlWc(6oZ~O}gQ>bV$sDhayP`31kc*lXOE+ zgkY1AYuhvItnRwA?hZRYAG0&=3^P09fa9Z+AW6VS2tH68CnD-?+DAkN0)zJNd+OfX zkC5=#&;0)X-Tip^+*H-6Q>RXyI(6#Q<2D6%ZRhkH#~JX3!yLB@Aw4$s_xyiakUeVR z^P{-mCA~Ijm(KOtq|(*3O_qknwW}H{*I245>+9FLEi0-mjh=c-ZN0@_SYlbTc4hUn z)YOp~D(NBz$F2OqWlHLKj#~^y`nl^Y$q5OWTm!(5<4xhXWrSiTwBDB!D4;kZ#5X-g zO3<@i$8m`q_aC!#XjcUi76w+zQj`&i1QAFN$E~77AL`eLQhEo+JrbKX>$sH!ALiMy zo#Qfx$XoIu$EAat;ht&k>Si}Gch%z$vDknY?TCFUIId#a%1U=7fIBIp1ur#w9p0G7 z#&LmZjS6DJNMwK~k|EcOH|DXSo2NA>c&v>(kV$oJ#T)b32(MYen?tgo$_ejbyfKfB z@K!W65y+g3KLLuJ`?L^$)e1c(^B&+(JBZJi$42yO6+G4_eaNJGe~dTgu@Sxe;!*-F zp+Eh2i5}@r%wr=wcb&o?gQG(T?Wzc#|lv9e?rk02X(W`sK!DDuZfcr$ZY`q%W=rqbKVK#@N-R_|<c6W5M1SgO=vvTS%W*~NRpn-FZ}IMFw?_yEV{ge7Ay74Jd&!u$BvHiXrp&oG&^N%YNrk>M^exqOqaA^bqU z$p7YQH!oT{-284CGmj8{GbB55j)oG27f7d2!Q&|8(}%k}DZ3{EBKOsCTxgTd?UOA32v3q zn~Bjji1g`Sq4LjHQ-h@u;zKh-&q_U_caP=n<>mJX%Q~WhhFl&76{#oL_%O-{Qd)*6 zy)GSeNu6$+{O}DN=g$v&2XyNNr!*kFCrV$*i@ptq!O1ldzy^#+h-PefzXz=u$mu1l zY+&Vl9$Mx7tIlJTPh8J&5&WgVmtLU^3yNJ8DIx$KNn{`q`c3O$x6%7oy~h$MoT4Bx zR-{+t?r#_~k4du7X~+oUiD@Pz{V!gKKBwC*nSvuIQ`!gVZ2J#r%$2bG+qZ8=(g zk07@L=`zy4T+v^bku6FEP{guKm*mKHNd*8pGDXSJFZ%B3->s;M@x?XJyB<+$6MYUa zc}SGnWeK_`InMeVXTeME-S>2qiPFd7md|#uUM?_$X1ZN+SnN_kCW@J%2`=9@8ElCD zf|o^KSuR1HQWwYy;RC30U`RCEU^H+x1T}Fmnmi<0(O|T`A<>o$MmsblTE$?ro*~gz z4@T=660Kn{TEOKiIOOs9dWom?4S@&LAtqp3eOhpD4-7ap~vjd8x4q;48@(tJ0Ff}?hki~A3U;Aq~ ztn?9GRp-xqN2V*Zk6g>>B>dz}Nf@4l`0+!zS9GagMVnSU`V_ zBT*gUQ#^sim$X4+Ns}$UvUI2)(@<6u3Sbae3Zyoe5U_X)4?_n_WYgH&(-O~K3=`%8 z5!JBDUnQ>bXS9QAr9Yz`dpaFw$?yfY!B}-V23Wk7#ZR#K9f)h>U;xb+1Ork#ETIm6 zkfKp>Lf0ccF8b42L5N09wy5ht4JMM#p0D|3toB116%5FWqLLd!1-vR$-qfg^xJW9t z7%FJHLWQ_9ZeP!F{B*~mb*Yqk5GfJ}vDuB@EP(vY#ySRF!sK;`pWe3ain#0oC0l79 z@EwP~=BETOp}_Qb13JlZl!#ofRP?4&(Rj5YYDu&reN07fo?FpTtjE6Qm#~W7q>Au| zD*DXl_>5K5M-{yUEVQUcsc0uFGuqJ;SJ6)tN~}da=T>xoOhqIm2mVk+vd{8G zATiqUETnp?!@MyGUGJ44U5NR!H~_>U#{t0h(Kx_S%rJ3)KA+?LIKcbF#qZRyjVei3 zKO>FrkI{H}kj7^!S_WR(h<5DaGtrQxWor4pODNxterCM`9vF?C!|?F_@H|9lXnd$N z&f((%G}=@#yr9Y?woObXM=wq$C}$>_KwXyJ6DyO8EZ+f{FdC2vqY*2USQ_VPPlSfX zN32XPfd*u|)DgHy>tkVh!Nx7;g`HG1*XlUIKl03(aM%ljtAPgiN47E0*$C(?|41!B z(J4VO|3Ah?2+aRyEt>iNtO(|xd{AwAf(^E)PG&kn3({~^9O*;}?u#nWUl|3Vq zI>*MM@OQhNQV%xG;P1A{+-a(uk%A|M;tDCaE{=2;!F>hsq~Mvs)s;%gGb_uc2Im%QqC`Q^q)U}F>qHR!`qqTu%FSKFK0BZxdB5#nj z1zaM|+JYs3ugrojAQ=oxiWYMq_0=gx%7&scNT68`8$!5aY}!V=Eb~I7uskJFPlFVE z)J#Q+tT7NKm{eR*G+gmzX}bnhkRbXHYJ%uP8lCB&Rvkh}olUJegwTSJSnNTlvDkw& z)dj&tdQ538wIA3@kAa-hW2i6MW2zv3%vh#s;Go0`>>^000xVQWKu`slmr6nvT*^fX zZUp2^`qdE8FY@4G$B5?mWP&-(@kwHF#ql|YcEio_7twGeuP4UT4R(0aRrAmA!xxQj zqPY&UUoZTPGBdNa%lAReG0~zz^M%|j#kiQkdI*w`oG}hq9;kCw6KacpyQ#{V! zMQ#U#aZ%1Cxu~d-6&YlhM_`|0kYSdwC*g(53*wNp)d9my+Y}f(xd6pj%v11|0nLtw ze#oHA)=&;?jWHQ5j94zCPD^RV3))L&;!bG2rl+6B=;`Y*dira`L6kJ~LJFg)X-wCf zQzh&$VM;r2vAV7~_ZO?{DtS8LwI4y$MbFLqe{ z7n`W?dr3{yhFx3}JHMwjF)G;$4&Z1jhbc7F0RAsFRN-1_=%B%t$3|A17j*5pgUz4} zHk$Kku&FcB!_a&THj4^HipGpvfzkv+xaz#I9MHHW^Jm%dOXf~TKMlIzwLsK{Tt6z( zriKG7u1pQTWm7{`Zs##gYKNIFcC@HEY`)mhsOda8$d^11VB9OEE^WS1N=2I=r)xv0 zO+1ss6dY=ni7Ck3?rcB5yU785)-`O0W3>LPi2p<0ny)Ku3N zeDcavd75$Q#aW3Vy}h9;WQ^)b)1m`UliYbU?FTC(ns(9U;;I-{XnIdxEMkoH==UCj zBnP{iar$4Sa5ekCL}BNZ^Hg_y ziJ>~uY%GrDgf%$6BZrl!Srvh=OzCb+Tf^{9(-XKMc`;;=93I)&RynL_E`hVSrpK2% zuWFtnw)nEcOT(BH$2X0&6IlNZ)%%Lz94AP5NGL;VSX64m=+NU8ae|82`eAwjwZf#U zCi)`vCkB6L?MaK{fi>DMDhusMUk_sk21~(L(C-bRZ`v*_4oz^Z2)bmLBF1&)!$|02 zJ2hTr@ue=TMp5a-mpqRW?O7hnF)7Cd^@w3GiCjo<409-jG#12F8=3i8fuX5JC8Jp) zVsvO+Rg)vycOJ6Fj(cJSpp9Z}^X5DyiJ_PYToc{-If}8U=PzKqF`h3r*fj1gEP;r$ zG!`#hBu;1-Uz72{Wl;et-ZB+tJJnOPZD4P}fV~SaWgbc!qJ}3hTsGz{bLt#c({gsd zB07%^RV326MQ!s?b*HYlD^8H{%0A>o z5b0q}ro#)WOPhzNA4y!=SQx$tbviSqks?Dfk@}0JaeS-vFJVH`xWC8=W*8;8h?z`P zV`~0I)|L9&lDNkCMP(|Fs?ngGUDefOJtHd1z-YG*jjOdsnYb=?(xOtVhvGOIev(Vo z4{8ZHC0ftMS33M0dyc}z7aYF!xHevDt&ye<&+_m!rzS6s_8@IGnD3!$ZIK5jfjDI| zJ~oeb!`KbP6x?6)8odCRLLZ|*W(T{239Jb`jIZb$BlU#aJ)?7a{def-8sVJ%9PZJ1 zz7qYjIz;OSZc})i=>5VJ`WkoX{zbWgmwET{;PyjU`^>%d2UzHMQaX*(5urs-=-|% z�htttBF2x$^9}L_T%eA_TbNH!7cQjUhuVfpby$GN2F)$7r}G*(Tor9Qn08>?B66 zTNXj`bT2gLu;>k>JEd2BsTD|y(%ZOiY2E)ohUhapt#3A_xqLR01FuWhjx_(^m*6Yh zfkK&~LoVq<`Q}6_^$J?(;a$=ZaJI7ywM{4KR~<$zsZ>j^c<*6sDISt=l(pI-??fSA zS&FcW-Hr99Kt{%o=_aZl$9FhGOivGw9s4xWpn4_py6JjpnF++3Kpag=PnTcAE*IKp z_ZOY@9{H=g)L#_v9(mdA!X;GEhf_$=SDKM4?6QI1T?YKI^Y(zKpRkFlkrtSo(z~Km zVh)de>M3?-G%Y-~3Nch6pb7|f4d{oB3xF=2c1fRp{Uj8jNWHaMCX@-K$^~JOE@;}k zb0+|i<6CX$w@za!5~Mq_%LtN>i}waHAu(R`_Vvq;q3w`|Z?+58TN-=J3k=`bLL_im zRm?gdX6+#jz)_~Lm-4))NAN%VHK_R#MCl-s$(zi{d4}XX6VxfCH6Zei_q?}q?o?dE z%>osdzj`2K^4@RY+*3j!Xdv(XK9ycz@aY963F;k9Uj_(z>yi!-`|`cCLry%t2Dsg)_bo8leI*8`Zh(L8K!MNX z)SUqVcmDdTyk|`AhrMUa?uod&>`y&`>-kiP{0ePsoUwU=PG3UM4PFHP4+%B6rY)@^ z#!`v-;4*^y5(s)HqOPRApcA40j&Ok#avuE1Wj(YpO<*HaI^aCm=duR)NAjpV#vHBy z^q|IPQKP&VwF$d$(URjInoUGm8Sut$y^bYm+y<|tlD@fj{2w;J{GI{pgoFnX~m_s>0!;&e%S z?7#mE5#NM%O5=iZmM=Lx_BDb@d&6TNrAX^ZcPa)^lIY7fLBC*we0Q5exEFmI>P(lg z#qeY=5LMzn<0jCO-hdqdZwsDJM{C#xWBDI$BeEUdt}xX!Q(P5%NACowc9<@ zH`@B1XA+qL$Vz|#G4>)eWDN9QuVOEbn^@oUt$1VYlPGb$?Iv|8rUSi71g*nLGQ7snndfq#$+Lket|9j05g&L z$wMaQ1nG78<2hj1J78Mxa`}us81^pAB+XE4@Mv{b+2-VgUAh4OhoUt8535njGIP51 z%ccp+;1H$bXuSOIpaJxe;Z6+p%+Q~S@*$c~$&O+Y`3#pkQr)19LkgGFEss~oU*;kM z(e29GN4UZ30aZFuEn*Z@c7l);(I}ZP#1erUFn*b%wjb?MyPC#8i!*-b+^C`>ftw93Gfo<8R3p z#Z>$=q}3M9S2)DwTeU^q7(e&-8m9=ChVT>el;P|>b{Dv#M%?8es=xt8n+-a?cH=BI66qoK$Xk>Ulnvxdk zV@=Ui>hMPq)nY!`okikY`W8#{yE4)n$BSEnEagq&#HlRCa@T($>gNNuI}T$r|3@8&8XU<>gfy$juYb?$Gr)-GEL){;0;$pnL2#8$0zC_?P#EMzXJKc6YoS4j+ck z01=et1i*fU_l&_kLOL{=YXP^uM>BSSn3TZ$K=zf$&u*pB=gP=-p=(O*kV>WtERl)e=Bm#CmK+$md7%=?9|=`&||Z)i^p=}l46XE7i` zM-^qEPmF|ngylkc#JCAQ8)l8y$j|Kn|B2SUB45}~4d{$=Q zeyr4F{nmXM|B`Di6zx3`Ox{o6;LyyKYs@JYkW?6s-hLeTW9IOT!0O@Y1+X1eMP_pQ zGcg2Odz(kd0{rBVQ7*ojIlmd1Qaf_c#O%9+^raD~)1A68gaYaLP!*`{AD*NeArPQcPkP zgZ#AKBa8kbJz0b)S^MN&;GPz8xgBLtm3RkuL)tfJh$lCv*X1u5z_$kI$+7`?^G(Du z68?gd!Y+(zD7aa^XBO*VQ;o0eWQ|mkF6F}%m!#KV?36<-)11;L5r4TF-!;G|cENkh z3Ap?#X@THC1s}@okP18u?e?s8yEU+jB+M7IhY|rbhP$99 z=xHd@2w#cW+7nfww5%TFwf1gK#5gdy71Y*~o0Ay&=IoFiaI^cX8t?rU&V9^3FJZSH zjQ-Q9)8Q^38!d0Q_HOD8??W9WDWT6=KlfNp<(mSqLY~VJp|5}4m$!cIo=z>AilQMq zW5;tOUisdqd<)6o2qYOg3P+HcuaN}5k;h^$v= zU78I`%?z-MtqJU)C2znEvmJ92&0Hc>QP?uTwS>d;@gc6Xw#!w}vC#Ewso5_73wVwd zT-0Pa!5T~!zEvz7tz7q6f{BjCSY(-LAw!I>xVm$?3#2_x>7d{}z1D#_&Mf#gW|;WL z+H!imZF*N#o6~Q*-Fd1%r`OTiEAY-<{`pT+?C^0Cnl7&qyc;v1ZKLc`;E#&7{Z-Hf zoK?#*Qk+!*yR}2)3utYxtW-m|vkH5kwTIOub@7k2)0|)tynPmb+FNv0)g^sVfaL(b zR{B5;?-eG!E%@hU&#|r>f$(N~QefM>u=8Ljbl7SAizj)@r(sYOB=8H~Henq1+u@F8 z$ysf!-J3>>RrtX0l+?X5`|Dqby0=4VRIO)}aMGTE0hG8mF&kAdLMW}JK?4M^dP^2M z?HS7Slb_Wgpi@i&X2ujEqAWzA^~p^myrG2F!<(`|VnarZO0T1;-`$xlRvl!ud-sQ_ zqepT zd#-dzJ3K%)>DsHG zLd)F%+1IO)YnR%9z;uCUy6|UB7Y>!l@3j zkM0pd7~k0vFtA~^ogAx(>AKFhD1~{Of8Lqh+H2>X?csd|Rb8kX_WkQ$NQ>Yw??)3M z%#n>(`Q|a_Wf2TKc^Jo)uKo$c20J>>^z#HK_U=*WA?AXL z2V{FzTib^T;=v=hDco!nb%D?bZ>Ir1t$sH?zvHxQ8H)xxc3V7YkkV$}d%v0EfAmuz z1b>J>yJ@Co#0>(@6CAWF>CVU$ydS|#AD7y_M|kgPo#zUt^>||{g!N%^V_JobaAWT4(aM!&1B#(nEQ+jYCfZ7pX$klr-);A=9GSSA6x=R-jJcQmrtmAFzX;Z2o2J^Oo-A3| z@*!yt=5o)qvDSMW%{vwplzk^i9aP`fp7pW!l>|(e&A+Ky6?(P zsF>4wqVdzvH1Fvto{7GEOAUO?>#$Yk{#2WLUGTn7Q1Q1wP5uqV_Ct!=q^Y>sY_flh zRibdpzx+94wn*L#sP5hrk$?FO5ZJPS`75!$c<=_y(&}cD;J0Bm*oVmsUY1y;e912C zqTcu?;nV=@i?+}S@BLH4o)Mk-Q>c@D`BSWId_kLU0Y+wm)QSBl>w!%Msf)Q=F6lkY zD41`RZ8`Y{v>7a-5V(O7N?qVHmGyqNL&`VFZyd)a;C!7fKLrR-8LV!^GLk%xB+ zFxiu84RQ*w=WO@ebZj49F8}KMo`I8;CRz0iK7(IPv;Up=0-XG7n*9|SRYbj#B-JYk zB|#W@W@s=|^d-cGja3Q<#YQW{Wl;y^P`|>)GB^o1KZrOemnja4bXpFe$AUYUWZ~jr zV-utPn~hcQW(ByjccT+bu#DI#4_!w?imk}^-9X!ySd!bb0+8y46mnR5HyL&0^0GD^ z2j%}DO^&YqHr)mhDhv z&`0s>)h(a>9BpvEhpoC>j}I@Tyj1^!B1Yd633Ef^K=>B)TS^JHDXL4ow4%fB&WBNoM?R#nbP zD&q9l>ZGo%-R_tDMSU%2?&2TPp$z|Aud}N3G0vqs7RIqlX-32D$*4>|`z1Q`9v}!W zF{|ib^sP@%+XW3^r$uE)z&1av>MGW!kP)b0j(4b%1z0Oq$ge?hf@9*SFd9^FeZJ_M ziJfIyug?CKwh5~qWt3^n3a{=fwr2g0No>vfT&A*Swa}XNEBPOZyFVJKkCL$`CZC~t zeu*?Z=vi_tkXcjc)auWa^(eBotjuT*V~7aeZ}lFd%a>ZcL*G%H#vnfX?l zx#8MuND^AQTkMc(a}Dj}xv^IBTwTL<=q?Z9CuvYMX<2JvR|59&ct=mjShM9=`1Tsw zQlyi=ir@)`HRexxZ2i%8UeI#tkisjb8!!U8*xkdfD{=Kaek5x%@hVu`H;fVeV{f^I z%!B;Q7ZAf}tdgj{Ow5uooM9aXw*ckf9e@XJ%W0Ot0!c_ZKz~iWo^e7_w|lhI2E-Q9 ztg)Y1F(7_kGdaAOhTLj*M$KGq&gQH6mwR^WFoo4H>9t?Z2N-;!CSw}+iwj~4|fHOsgR*k8V z5`wc7K!STTziCBm5uxp*7cw9~PEu@4$W_o9=C>|F-l@IB1>d?9aj@>s$BOTAIDt|D02qs-X;{;C z?WtuM8-uw3##lg{WjS)f%tQvSW%zvSl|%?bu;o=#$kKX}Zygx{=t4XJ(9bG{frfN2 zh%~`JL`YL_MNWWKih)nY#{#4Z`TlFk1R_Nice%C;R%DVFO=8<%sWs@pX z)^5)ibU2wP>M+-gK}qcIbg(%M3n-uS?7<_F2=gkQm5NM3H>w!4<(&(c;4U~LuREm! z)}?9uBj18GoSpc;moECT$=EAa?H94Qf)Ad`;rLweIXH6pd^+J&Ke~E~AnmbRb&V@T zNp?zSFq1oS4m+ezeP#y^@6i~ix8306fBlY2cibtxE%iFHdYo-vSPHTZ6+j)OytB@> zkOhk@ozwbCQ@7WOtv=BkG>FzW8jaHOv``X(PKNh~f|$N&I{j=MX;AYbQ96^;?MtwO zA<_E^gw}4rQM{Pdg^3ho$pm^o1l7-*-V>`jM8Dx7Z!RM)gLtP#93V#A?pukrg!kB` zJ_s>)m_CdX-j>gHK~{X$?=p5Egbr9RbR51MNM#mMnMgUY>t33%5s*?5plpE1VPMx< z8VEuo)CHfa1r*3%T5Xs5y?t*8l8#02^J04UREKmz@E$dI#z2?1(m?lR?XsW%4qw)I z2Mq2Tb1BNoMExrVT{0`AnrPP-JRa ztPp~0;2c_C*_5aYgxf+VWeV5EzNDL-IPyU#GrF4@OGpV2LT z!#J0B?0!QljR4VlKv}`_dF^7`K!SIpH;waLCc+Vh4T16h0*56{X`w4< zEta3w)@Kx}difVi|32M&d?FSgb^7~_% zajc;FEj3Pl6W#K#>$)6LFDW1xS7yuO{0;l?m2_$oC~xV;%y~NwOPy_fmIC+&QV&Yj zw;$)18E-{C@YHaL7Pc`bft@Z2Spi{-JR9<(_5JAM0)GR(!(N&?iIrT^#xL$Z&QIyD z;iq@48-;3Lui>Y4-Isv7Sg)fZ@P(NfEqnX*s!X_Pbp2P!x zxj9wm&JbZxapbu&WX|c$a%b2D>^qibtQ1cD8GEy{lGcx=z_`imt;?8Y0O8Q(z$ef* z&`_l3p2EJX!2d&l@C@!!hVJ%Ikbu_M&=}4=kCJ-#Dm3{OFiLQ0rT2|6$uk+Jcf@-s z;6=1Jfs6v)fL>LZ(837*#*9+#Wh%)RA}YP0sMN-cB5+h1x=zExC@a5k)_xZ@^x zKQWkW_M|DoM(Y3H}j1U4ATO(Kgthi|11V<9(4q?zM0_tFxL z#(;Y+KG;t+;iM5g`ue)*m@B6Pr-dL2rZ20{rr z0k-Uhkh@i#T?*}6lG=fCH6p(WKS|Ps-z9-N3?=0_o1S#smDGm8JeBnO|0sC%-Qb{& zf9N52GyLMz#|R3A?Vc0s*B(HD4K)DjJ+2n zT`)tL$NA+40hn+D0I28ywe&g`pVEVqQ!so4n{C3jMd_;d@G7UJk z2KdZ(UGV!M>i3x3!BhwT5H0wrv8*{c(dIxPbIxrNLtfe*ZbQXv-yIvs{I=@>7OWo3 z{1XtvwY*H|hgh{25-%kY8`?Kq*Bah@jps7bkB_Fap4!Vl*ac0@>CX9_{rpMrUy)M# z!94a?K;jo4R<{B`6{38db;5Ee%M^&}hJF4#bjCayn*c5Q3fz8uUn0NtDU5X0Vp|71 z8^J^_=N?0WPIv*v8m{kyKnYENBu=8){p-h|X#agy&&@4IbnY1;1GU{H^*eFrBHY!I zYxRs4tuWtklVJWc!+ZCR$o`fSUa){x@Mf|rjBHGaZ1-P+S0e+g9S;Z8p%c6X4+hH|+8ra(I9s%OBbenq}7-iSs56paPxOrlTm)fC4_N+5HvY1k*YLM(` zragpP!7QTHp7kwLJ5_*Wbu2pAXuP7l5_^nfb1b;I8XC!pp_-h}9jzxfWuS7A+cW{$ zqF~)C@C)~5k!?p(ZNHGU*8y{}l!?G?$JFD-DFa;2AhT_`=u@ltxkaXt*&eXAL{gOL+noQO_yPgK^?plb2i(on@5 z{G{#UcK*e>2h#=bUnWxfk!csQUJ<;X8wBsi6YaWVcIjC*+Ogf2OdQuM90O6*D|P@m zJiCA^7qdKt(Pqv)n#e)o@5(o&Ft2qa=D|6bNs2y2EVQkZ zjZkc`F_k|>4%D1rB0LFV3U~G}mccn&`$K$la!&5$a$GoW2zoG|FuwS$pCX#?FP+UP z(g?~GtZQC@`%elo%~xO&{T{PimIyYg3es?w0>{2U5W*~bQoje@IJ|HlNgWrIb;+1JJvPA`w(?u);xO!82Ao9A78ODR#lznZq_|FkaJ>(KCGO<;Y_-oaA(TYZ6kBUM|OP-3}2D z;mso%v8^qXMpgp9wFgnHVFmtboooRA@bDHmwr$K%G$qJ%p#Y0l_S-d9h)hu)k+=6V zUFwwo1SL@Ahs!BgICjV*vEfg$3++%^_#@OLggX@*7^)-#AP^<6$;p4mMk`ZVR(KOe zB-AI_lP;X2&aoRWcHYrPzv)aS+p+BduE6in1^*f4o&E=O zRAV-4wmXZ6;lO*ZgR@dra1wC9oKyOiim7&b1%UQ#o9XU*kyY_!_^tm5(qKC>L_Q*y zSRF4|C!*z|^$Sm`ydQ8J1CZIfQl^!D%(wzv{v`vwq8d%7^bHJpu#tA(X~JNxou-hg zF{B$$zsF8336)qP)>O70z^QJOr6NYULnR#!gr3Ve$uz;8#+vWTPjgp=yEacJxrMHe zbV3S^Va0N=PB@J;I>F?*N$`HEbK}7IThYHcEId3fJjd!u7qJ1G={Xh|xEQv&akgw_oyL4HxA~$lC1>BNm&j;nqt9g zc&~gdK_P+$5$##c340ScDH{abQ3zBOzTiu3aqJY^RB-;8CqQzZV zhiFwRNJ1Li-^82%V(HqNx&t?wVH^eQvZQQBNnBPQL;Dj+-gXPC?{oeIPS-B=L@$40 z2Z^1}W%`V7E<F!|`^+2ttZGcE-HS42HYY{m}1 z9r_HI{8kSnr|v#^=gy)+*vC2bE?CQ@0w3Z_m^Sc~3>wOHO}3JG$T*UbOQ4&m6ZSST zwnhLH#uovp*CB!%4A%H74i+2^5-)LShv@M)4beaTkdCQhG!Q?U0glx)sEq z69#L)Hbig%GSlv>#Ma&Zf+`t!RlPpDFI~ikq%NH-A?m~@O5rwV+vgU+`s(H*qV%T% zjO#+XwJK%v%{Z>dB?tLeM3;TaZT0vD^MjFnt^Zm1?S({QGNC>yd->_rRb`a&ac7UAAZ%fB&MfBgB zUf@q@!IxYv{B(O(54Jo*qCan}=>71l)7nPcS-7{J*Y));?(lE?J*FXhMmoL-S?JIB zxtN7}r5)q3>o^{JHW=Z5cKN6HVY)eX7s$ou1zpy@raK{A7y{_lZ0uY*rDNFZ6nrBc zzFXNAr?&$KWWV0;(w%nczH&;~Y*~P<%`YrY>zT$w^QBptPO>7;IkC4X(3WOm_;uVf z?3WrL!Jb$N#$DY-7XCwkcf5Ui_cicU{bL`#hRt%~9<$fj%`mYK2t^2@lTNope zqykkI3K6s(_IzLIiFTZD0eYhfy=M3KF{)+OKaz1QL765ipN7z$o>?h39VRg}RklyayG?&Ax{De2g2?Q+e1XH`?S=PE|%U9}f+Rz)Atxi6yu z<4`}kc?sB|-8c}I`ZcL^#9h+1K;xyieNL0~mfhEs*7mItC+>poq{CmYvs-2CsJCH9 z{V(!C5OPX;m~fsZnfQy+@eRRAUhF4_wCbeT_DO>D6|TPd=gff4Lyhv&uvPg3c>V4b z!fp%7U`HEfJ1giLN0Uf9CXu~{gGaFFqocjNv%Yy}56Uq3jeo|Dyy#yrUd+N3%VSWZ zV+l@c7gHbn*7jqhOHjT8lHqbRHwBy~gc6^kf@g&AYveKDOg>G13mm)H{xe<3P|pxB zQRJoBeL@;mgr3NpLa1k*3!HZ`hhF-%5by3_!xu{vGYl@4IeNvDJ(tgHTx3<@sr{C#?!FHJwsmombc zb-!4kF_n?eTLI6xOgJ3-}OAE*(jFOvr1i8*|Co`pYiun)QLWE@F7*rV%@;v@xV_bdQa zQqPjua6C@GBWz`gI#*IweUSysBB&ww&I~y5LY~V@lZE_kN8}r=cFdR_)cM9&r zEif}?>^$QeHGzMj4gZ_M!5~yRjg1yTI+z8(M^c+3>D`0JaBV`6imw`vJ&ilkz3mC1 z(fkYh4j%Ih7J$>)asv^%i-*%#|DzoJn!hyed7MVz6S1i?FleMiv8r9*U$j4%)^^0` z?VpHSIb(3*t{4V%Yt;S1sH7NgO$i;F*rjX)N))*4KTv2hMScd zY)W!{=9a4T@Md@w-SSidWp#*ce>Kz|W>qXOGh)HRwCsx^{eM7O<6qh@)0Gl@j*Q=g zRJ5#X2+UEwj)y#sE&bDNp@jiH6ORZV;S__bq!DG;1z7agU#=&+u? zeZV3bm&D=mv+aZM*u|(DeM&o8E?mo%uSeyb+zsHJzNExeuZ~5i*CxFUcgSO;1L9k0 z6R}ACgt{C@it<~C(22?vT9Op;e|Q_ZC)2h){5S6&+w%AcAv@o*iS+NwbuPeOlV>#-dD==rd0NYb{;OXNj2VFjC2F%-K!8m>i-77e0OV%zB>E zY}mF^mMW)-4Bl7%fek`lrf0{08drQmRA8CjH}L{=9q(Jg$El^^Py|s z6@5)6c{I9D$(YWn3A-oI?-q#GZyWI)Ya-SAWw-}lZ{zOUY>hI_gPN|A#I*0U;1ppF z!+tq!sNnyqa$!!xN^KIpW8}9!1Q-l3PN=2P{BvWVQFFHTdQxgSF&1+I!TregzQ8W4 z>G&V~3MqW0F?(BS##}-9Jpril92|Yn`7Na8&ag01L$7O?Uaw_mGr_IM#WpT65FR@P zSt>VZ$30^BHBa)x`fVfV>!EN6mlHunp^i(48tu?CT1**ObRm^Gm0^d+%GGSWc7^;Y z7>SG@+{x(nbn?&XYotVFZ^BET2;$Sy#Wbksq72;z)JE2A=lq@m4YUV9QRrw8&e?_hZ8k*AQbrnzLc#K}ief!=X1=QPDs5 z85EUAfqu?m0i!YnF&wL4OM&0Ig_`KSjdR~m4&!Yu>(@=+w@WKCG7D(e{xrrS))ur4 zInJ4Nh<)pcVL6JbekDQF;i5xFXus!n8Zn8XT(!Nl#fRa28wxQ}paW_O9Vv)OUT%k< z3OXZb{C0%iQ3j6~YN=$yVo56;iN?Rto3~AXAkhB^!@w{B1h=(lLZBDWx?=eX& zCa;`HYLNDB-yWWP$4WBH|J_4=BB?htd-B6o3cvTM;rG<=E;YPE4R2Ay8`SU$H5{pi zXI3b5j;Y~$EDZkj@99@rBM(Lkds;TB|Ia7m0(_{Cn*1*K8%Y4eGVa-V0l6o{VdH!k zY5KM|aJ%wr0hEm`jmTqbi|@w=*DGm@8uqE7R={>G;%Dx5Vp5d@iL9vj6Ybq=v45f}h=gp>hhPg4+63>!+^2 zX_f{539-9&O|_+_vbN6CSbfV#PFT6JvAU_rGP%LBy0Xbqzt&Q_rgByF)F!mQc5OXp zu`H@>bbBi6?(|eQZeRtgA*HpAmdPtED>k^Rn=F-XOJ#&gv;c(YCLvN?D6+oRz1mWR z|M}RG-8{J&Xl3>H)~{V(Z)s?%_N-iMVgJ``CDlp)d9r(L)!I7Cy6VOz0!&^xZRB}N ztf+NcnriQ>RxpV(0L~>0PSMV_9=By}jis@&epU4pi@UOMRka%uLe><^hFWwrCETm4 zE$b@lJRsSC%Hy8QdQV+l7X2T`^_+*R!#^arg_R3AXFXgXF25@ru2A1!bt>_?b|sY6 zH03Wr{AD%WgK&--FODg%(Ye%c?;fST7Pa1)NUy~Em^EvVuSWgTklqM5(LRn>qxCAv6JBhcTAbk2 zF_3y~eXV-~>w;{8&@*s*W&YZ#XCCUk{Xe@Ze3=jLz5TBnep)ofe(-j~nFs!5+|)T?NjmT7a}-rJl+SitRcg}^q1Ngn(aqQxUEA8 zzc;|bmaWqS%I~?XiiKO>`9~I3zxEo1mhv45c?jio;>{HP?9nUMAhdn_ME1`RZohiZ z3vVNwb=_a?P08n4`flFPA>?yA4z(_xwJx9QTDWRW+b{FExqnx2bjgSLwnu*7JmJ@4 z?L}W6dHd)U3+z9Bt03`>{uVpvvgZZq&_DjR_lY?(pZWP)*Pe7*=RR}7{tf@sjDPvZ zN8T_!fA{44wa>21OzavoO0?1l)%q{5H~#+jyZ(OOUryg1)EDL!%$SgUT3r9&U2|S; z+We=&v7VR`anbCScssX!(EQx+llc6>>6Y#B`H}xWU9F`^wWnI88H4J}9+aOuC|&XB zAo_#svT0D-JSc4$l%AiTe~Tq^{=%{>OU}$`InyF4tupvs{&({EPCx$Vjt|Ywd>8&V zk${=wQt)y+mC%EeZ3E&IF1w85(h#Sx9&ZNX6h44A6Y&;=kKmnxIE8+^*@$mPcogp} z#3_7c9A;<4DSRET3vmk9T+VSz5T~$i0>`aHoWks@Ij#Y53b$K0ZX@Cp?zoQQwjoa8 zsZ5T06mbeyS^-160pX)~pGBO)5AgmDaSBi14Is|pJGZOw?m?VFH{Mqfr*J*qLx{H^ zJc{>y#3{US4no8!q+bL)f%p!D-{3X5ADBJ706^C;pJat@Ar0&xmI!22ZP6gJJ{xEBzo z@OiwiB2M9y`5gBq;uPvc$PIA{PdmX2;uLPb1N45YgpazQYlu_$BWx39AWq@;?u1Mb zr_fvsIUr8q5AiNRoWh+YXglH*uE2$l2E-c>K7n@|;uOAvcRS(~evS7D#3{V)9`pm^ z*$91g_=Qx&w;O8B3M zryx#YNdxo&aSAu!%|x8SfBHV^N1VcE8nF+GIE7C(p^b=BcnohH;uKDTXV`!^h4laH zwjfR+{qMn#AWq>SypJQ^hwv=kClKe>qkgxg(`w*wlbsuDfIE7E)J%RX6gsV5l5vMTu z0gg*YoWd%+QxI?2ggtJ&xrl%F{C_V29jergOyow+NX#}?7#j3N34wFcgcJKF8)hf# z6IUc;>xy*A24k`=jWg&GH|dS0859FP#T#Sd$7A9@j*0ii#AES`W8jymaf6OCz^>^M z^;}{eYOhfF8V+tu;EavBY}TG1$Fy&>$uU}=@}MbuL`7nQF`(}O?;{M{h#C6qgbH0l zTs*ceqBjHdG+CI8oGH(ct*=OEU~Sg;xKKEl4$-4}9#q>z`Msz|!!svw<_A-&jpYWnCeqeed6QhI?JEXj z1Ab)+cjaoLb@{opOo^PzE8~py*I>GuY-4%WVAM9PwOZS>JmM#B2!003b(D#VNU|s~ zVCYHc8*;4a^_(IZv}rJpnx1KO4-ToHe}-}dZ`4)jsLo=w_hRcD%v)@mDeu7;9ad;H zFj|d+Xi>RFwI9LTchCPW3Cy4|j3#MsCel;zX5%G4J`W*0liC$Eb)MRlO|@&Lt*oo# z?xSzcX-?e504O%1E78>`FbQGQA7nsVTH>MGsU|GVV$zo|90zG!P&j=~51kmG(H z13L;>0I=bnh9gSOHZ|O#hC9{pRW_rb3pZ24`-4Zgq;jx7a&n1QMxulm$Px@ z<>hr%<#h7bfP=;I8c%%{mru|&E6S@pjpb`9n-w5e3TW-x6>6%PE3aJBw5q(ixz>#} z8CPE2*jSHMT4Q2 z?kesD7B+DKeNAIEtno=bvMXV;*Vbdb_>F#zihqTnv091VVrX(VqE?dm9E#UfW3_ub z#pvv7lx~fLE zRs$;e?zyl8%!AH#)Dn^N`U{@hU$~0zFus$(cM|yjSpp5%vjhu|{|bLvus?^uyrW`A z!;a=1EjzaFcznl>9o%-T0Vo^spYPbYW7m!yPwjjv@Kn!JeNXj2#XW6)+VXVv)3&FJ zp00Sh;pvvAxqi?9e)^8=9l1LY`R@5|D}gECEP-)nQh3Z)?EZ{lIKGPE^9O`homM3@xFAb>?IDWR|=iIE`(6CHh`c@Xk=eniZZ}cFk-mHWF4<0llVv?bIfBW!1e? z>GB)a1MoQ~z^!C8FGHB1OROUAJi#=gZY}Pou|J;rrrK5YxLgMx6gkFd4g(~NXj%sxz@7+cOnx^1}$bV)YPr5baN)tO%P{N zLV{7@gX+@PLMG)TD|qoMYpbf8xI<{(B;6OO5b+Y^m6k1Zl$Yd-4ttpk2X0YVkx*Rf z6kO%Sj``(wN0DQp-LWu#DW{8Kl$1Eihrvo19&27vd4A!-MULXql32R>;Y$pKZiu4W z9VMmZ^8^>T0iWgb3X996oVKoDRMbWsWBy7Vl+o%f1XR2&sthiD!5$Sh$u!u z;i4s+DJJbGC@fxDUf?V#5K8mKa-~306k%Q&YRfMuDi;>o%kzqbg+Q7)gBy8%jGJb0 z$>+x?DI_LKxRgO?f~zd1DlRpeE#&8y6_mMzQb&1VS?T;j=fe5rLNQ2;isBR%7rF}P zmpRJIN+=txb#Qz%+fn2^C)*s&rlz_a0?Lez0*lH@#PR|rChoE*nCQEpFv!IgdPMc? zg=KjfI4ug+unUX0u~Ddo4CuHhv`Da*mxzVM(1Nms`K8Xnh1})wxy6psvf_m$TzV8d zuduAR{Ejm8lgnA)Eak>Wvmoix#X>Q7Eq2Uvpp!{5d5(pSdCq){2=0m~sw=-p)qpFb zz+wj&Dk!u&=Pl)~ih@a4&V~7QM?MzSk@hD z>e4vMXZmn$6i4lrlG0+x1ig?E&0@kTE?=;)aPdM6r6OudNgkwASdzz0j^bhD2rlRR zg$|O@A`DDM3$iXO*R-vKyDo}PVk;5EdEE8!uzXjkoy&}e<&|<-QHau@c*~Gi;N)(I zL9yRa2J2C_5W~R5-5AX%gjICTW0ne?R{o4W?~GMXJKgxHzS%oA39pB znCy~r#ppTFteH^^0+ub*kQ!(pKZF%~qm+3{#aLjD5uq^cJrsR-sx z@tCDciySUT2^b@KCEU&N*~CkBdEV0Uq9ujJTy8Xz21qfOhV>?WDJd^4EMJ5%WIYec zyyC(FhIGr2NXmp#0z*;4-5SMErii!!5HneEbE25kr{Dn%DUZdNJ0ylSeQ~!%v)L#t zDiO=(<#D%1!HQT7DqEmMCZKH z9BxSza+iZ_|B_O;nB_`c3+>#}Xg;Qffp}( zM{W-TE+PF@))*|ESF%_r;*?bc128s>wD40GdFramh(nTkH%2#?U>>2xBWpY^A;EyN zXDq*RZsHOQ`l_{S8tST>Em4>$0#jCL3FZhuTfkW&khm3|S))vKM{!ir0FI=?o$>}Z zmZ>+G>57nH5|+$!@g{R%n@vwHCk973H+=>zrpwFiG-au~j3;fG8T zGPvaA+miLz9nc$eded!_MkG%%ILF*FNq4K!Iw?^%20ZI@&M_1*&O)hMCMDdOY@IaP zIp+59bH;OA;w)}H(Dj5q(m5sp+lwPc>n5AVB)M{ zNd^@LO`ZgXQYLZEG1KXnRla-v#u8Wpdbk0^O+%>D>Haz@l>*uc!i&g6wmCQv=>)|2 z1YCLMbnFi{H?+MX{RvKY%{8l*W_uQ_T9rN5b!|04;|eja^VkN#@5q>M1_P^C@#!t9CN4g^3LWqvS07?s|l zuODR&2%}8xD&BX&|40c;#}@2VheSSYc~m^Oewr2k!1 z#rb>0EJ#ubJv)esl2^*$D-`5qUjXL=;Iv!-=LB#XE`US-{Z+*Ua4f(nx&RKH|MbC~ zK3}`;0geYRlPY+_H(*GaFF^Tm;4Fh^qEe;*G=oGsTwS<=%Wjg@M@mft1Xs)0LxuIa z2atD-3Q>Y-D=XcVDg>FpA!9-7ncV;n$F(}xT^CTR?&nu6}rEodX_T98zvIs}d?Rg$G5)pJke zHK~#;6{%`tFPNNi1IE$$WK8#vUi z=<~|jni@}2HTQZ1Qn_*+PKEKE=m-smN@Rlu{zo#k)Vk`r@+Nn6LpE(taAa+8m`p!p z!)=G4qa)=Ol-e;)gEUpU%U9H|3Xax;N-L-b-@K0FW>bZc=WpGEvp`wt!RAOE zT14x?QUH`5B%G_YY!*&yjE)_}0<4V7*AmprJr~i-WXF|pf%%gq^{q-uqN`GY_M%F6+p}y$9h`zw}T(vLi zE8VyyfqSfx5vR4CCU0uZP$N!jy9N0W9B_uP5~>X+*PTgKmx9>25Akfq7*up@c)R+q zD5~g>yRNCp(-7rYq4{^j2;o(~5qm2uZycbR}Q67P`u3>;*!XE+?nBz?}l2z}lOT-8sY z+xOMY)lrSs^rwYs)G)fCnrG^Z3+WHZS98J`Y?)hHU4ce06nQ1AgI*wjCoj{BJKq8bRki)b9-!!Sg#Ps_?A*oZ3|O zJ8jXv??q=SJX?me2~_ikz&WZ~8$E|#Rx8J^%QdY=4(nw?;Z6Z(WgG+_g>pAe(R>oZ znE*&d98M-6f%DiyDzy|iJr}_F0dV@ygLAHikAqkz6Zyz(XtQt{N)Km6sR(Jp_KVPj&w(46(cG&WS4Ne9)%+xw-2q)D zk7%gbfspnP4#H@^_X5%q2k8Uk8t&ziSbPV@88kPmk4&Y(V}LfoT1u@I&$$ zKr63d)U-dB(qAoF?evo#gC4CfOsL^PeUS?sr7y@1p6g*yU+lbyzBq-_k-k_%ANEv5 zxQetLRe;=(eQ|EviL3ZW^9K6>9?lw53F5iuB6!vTPvMzL+2bm_6rOD;O?okuj-AW1 zqCr59dx+=$i{N>Dgy)88`bEqbF*ABg(25;Q&K8{A#8nYV+zXsN8V-Zr2Z$~wR(pR0 z$oJzQ&jRvf9Apn5`{E$)05TDm)Y%$@QT`N={5XgKD*0dpa;}^gKtSXl{NLWLKSr*j zzQ40A^(xeGKuBE`sSB9MqQP8@?|enAvW|bmZk%LQUxXzBZ})cJ-L1VpUU&D-dj_ee z281g`?L`grtVH5yP-CeIoRR>B1bQt<7&R@HNEQEq!qFmPRZv-KXbqxiKi~PzcjnET zx9h#F$&>8c_kG{_er7)No$t(>_uh;>#5eK~-H7LF(COtH(k}La7j73BZAZ^|%tHrR zN1fL2|8L`;dr*Cijp91Q+ednjy59>u_Av~#!rA!Faa&_J#Y(49=bEqq-SnYV0;6O4 zWtI>+<-t8}He(-0&X_Uc#62Xc6C=zRMG-$wbbWjP#`4f(#%4>WBNvO#Zl{dkpgElF zcqq>GO%i+T)jF7VfX=P(Ff0_K}D1zZPg6@yy@@617lGspDqxOlFU3c=Dy~ zBlo0``1B#IqZl4R6ISiPRXweB-U9LVkrVDkeYU7|#*_}V0yiqjT->y&3x8%__{kf7 ze5m30$@zvLp?as<=6YX4wNX2F52|;nZLasn-HZB+I3(tn!yJC0kg0=wgFJ$LgZM#c zP90QVR7+f6&!?{PDfxV16!n19eo>-gc*Z1u)abO9Y?L5N4+D#N&W`GSUHIH=r7p&h zpBqKdW#}MVC|B>8unCkvQW*=`KK+5U6R@Ix)pJ)}>eG4hT4XMF&E z{wBpwjOD`Xr4GjQeEc=|S=anzjpAnGs+v8`LtTwPR5F?+{WhkF)%}H%{2et2OiNqI z_aJ%2k^C5v_GD7$XOMi&k-Qo1dB^^w&U+yl!PJtqd>14?|GlK-0Z3Lpn3P12Jch>% zHq}{3{sM0bS;=2Pa_T@*@-!r4cPAy^faK2|$#)^46$;7DHZvw+4#|I>`Ays%D^BD; z-+waZY4v>FAo1rh)G5)uedINPFn`5=cA&D@e>QN9qP{(3ex9CKaD09Q7C&qi^e>9x zmk01M+JuSQV5MGdm)!skTR#9}@x7Wc=w==N0i7pAM=87x$;*!9*U=UIr;$kHXUoLK zmti}*_mGj^E0*U>?;Qt+D;N@~_HOBOjQs(0 zKK2o#OkeHx|_zTe(L&?Vay-%i)RuQ91|%1z%5b-wfFG>7{iimrq9i2?!g~| zjdjgNR*NIpX<`*OyOay^h9D?GvgSw{kUaQNW|ei$LbB*c9)aXpE!h^(YBIHIyOVe9 zoQF;k2*|d^dJ>X6N^8--2+3S1H3-Gvd9v(Cz5~gLVp8XOkSsZpS0FiG5}kN_f_qQU zP9@kL#1n1&4e(7~68)!0_0)`{qU*nazdMpQL-Nijsq(=IGwalw#yEP3}8~wB@TX!fp9E=Q^E;XYK?! z?3twVK1WY_tGsYHaunS82YGK{0PlS}^25CqtsR7WjznbDl3-I9%RmwiS%)~xh*A0p zON?wXhkkyg-@*I%?1x!cx^+JRewagI$vuQa+F_45>>Ur|0r+{3;RmZIrTEbjG!^Yp zS$UQ>d6W+ElM#__GCykr@N;DVepUwHXVUP)cO10TRlRM=Ux!ZiOrZ6sG3bzQ*<&jnMNbIWqsWeXh)4a%FxI=jSRa6~BZjdOrpL{)V&)#? z9s_p^@bjGE2boJ9!p|=aKfPPrQ@Yg8 zqgL;GQ$4Bk8gxuwY3hRWL*;rF9zZjs_K><*hiGtrrnc-JGLj)j<(N6;D&S1&qG0M` z98%ww8e{jX3*CFhs_ax}UMT|cpE1MOo+@>yyY{aDooGTTmP~AD4>gi~7jWItnS2QK zt|c11I~4k39{ME@o$A6xZ+Ym;L!qyF=$Ae8iyrzp5B=Ow=udd)S3L9$5B)P9`pQt~ zYaaThhknUJr!jzA@2f+hf8ImC>Y-ou(7)iJuMLHM!9&01pJoJsB z&@X!EV;=f-5B(Jn{nAkAFL>xXJ#?y2r!FRc?4e&C3Vp*v&w1#h9y;yH?$-O2q0s5= z2d7*UyFGN8O?7l8|I0()918uChfed@4%WV%9{LWOO8flPq0nEnc61Lr<)P<1^jmBy zjecz?^vl+cM$dcbyFK)sHkC%dJ{0;(){aI$>YY*R9sq|NM??fBq5x)P1 z$vpN~ju0;|n3y@+nO|sNEWcqb>r&*T6u6Q;yzy$5jy&y>5p0$Ymvk@99x0yHJQ50! zIP@sjAJZnsy}IlV_bI#F9lwF)+E)-()%{H#l!MY3-dh)ye*JmSUq-Qq1 zQ0TFTzUrYbd+4`$==q`0YaaTFhu-zjZ}HHN4uwANp`Y{6V-Nkhv=lQlP#6lm;-N2l z=<^=>riXrVDD-46m9TE{kydz{d)pS_wUP63gjYv-0c3H z?5iH6Pye159qD%_5_#5T|7@SKOF#GqmTTXdJ#kTLxwp>#W$_L5?KX@i4ta;+ zSyOhbfFWOwNBEFf4$Po)*>jjF6`}q6M)E_bx0?K5#$$x$RQ*i73pzA@<{4M~$T;pK zJGewwx2vy2K}71vLub|aGkKbirJf!4nU%$tn`QcPlTyQ6fQE$l{`i|`a0`83CRVWLzX-P$tCR@Ecr)Bt~rvIA(38&b*@4pqvdS@o@J~3a`;X& zr}MJ|`5lEi=ZAOf2?CyZ6P>p}hYtF~mUH+yj)z6W?l$_4wMJjI?#3A)wa#78k;u=f zjzePkkv)0%4oI+y(4k2CzgX*oQ*|Ms?Ha7)QIg=G>#Xw>B+Kd^R_r{ZbaIJQFG51T zW9?jp#Le?BAdyI-%TCWkX&)WX1&Q?Loa$ak zjzY~^jv&e5qUA7#r0dl3!;nxCi}eH~dB@I|ASt57S)C0?%Ff+LUF@NUWbGZ z&Y|p67Br~;k-8Ye+m>$Ge+QCHQwmCd7bL6r%Xyal(0KMz?1a#v!+Ko{61iLFvOfZe z%-6BxJS0U2>nTXaanUmW93%v7Eq@ylxi4TluRN+At@LgHBNvXI;>~@GjzzdO^6ik29ZJEeeqX-`Y{DLR%Phs654e5W(2uyQCd?*s{N zaqRppQWcN`n1ap=INMW+W)|kcQmfRM!;UrjwAz?$%J0(5Oe4wh?bVUeaNEUP^&174$y-y`Ay(=lu_fWa>oaw>*7=w$r{$iS%`UreIq3(V^d;Jt4rNnmb8>1r5hAq7diE%m8FIn*-{r7(kNjD zyB(oq`hzkI4o`6I{iM2d80gT)Xd$|zcC!%-kdR-NEX8ryS&C6swf@Sn0|x}+?Af!( z0Ty_I%{jcn9koKPEf!X=kFz5lXq)Fs)M%?^xUjxrBWyQoixEI7&2X;PoGH~pd|1(H zI})`XqvXj}YI3=3VvRR^tx!_UC7DW`s?kTfI`pVJ1o zMw~eO;GqYw*8K4Y4)Rjpw61=*l>`~e)jI7aJ5y>$FojjJXf0N(NRtF)GR3C`b?p++<%ReVuW zBdQ}_A<~D<7K++Pq*Yg1rAoCM#yVqSR$b!d^(|KVyoRVQSH`zH14;0f;Z_ympze-6}%-pkqOhviySYAnRa!reBjF#=YI7jHZ0 zL92^sNEE?Hg>tOX>)7C$8Z>9Dx_2fa*FWRtnF&(Ur?q6}vCFd-(~6Ax=geNQS26bB zvSqDG@8*YUepruE^hxNvh>m>fSE$2s9AKVGy@U-!tzN?S@wbcu?WFJlywHoMr>NBd_-?KNsFLk-?U%WGD8^!U!?E^!#HKCfj-J^#a z`U^UHsy>#pr@Y+okCvvfwxUMmg6c_EGBQ59af@QjAl5Yp4)(7nxa#sSKU~TQJ7*n zK$x{#x1+Ij!nm#aPi+CVF|Y|KFqO)vC3Fr|bT!098~eM5%~?W%OJOGu3Pq-~3+v6w zLJiv}w#fixQMOJ`r=(?1TIixw($Xmevi8DEdkN#6dbo%M$lyfkT1a0?V477*#-+~Q zt3G>}!YTNEPwyAcyTgMsWxSd)p4j5us_@-3+tT+LmT?Rlgzw%14DlqV1x_cnvkybO z4U)=>*|dJTDLdJN)js~p!x+mp94Jm!6Zvwg`Z@Me8?SZR6>Rkx;@kK#_3EUoZbtxF zp9#wgEv)?64XKqUip=TFh*r8nZLq!2H15oUH#pocSa20MY zBalJ378b;;)E=@5PiwXlE{0A#YW2pztq0vW=;}2C*ESYjA{WJQX|7Snm_#{i_Dz9x zK!j3irC3FmM$E?!OzTM+?N*5>-KU1hk4zLwMg$XwsM-+xWhAY_k}!}Qz^9^%jc7n1 zSCllJYGcV%ow>1FESGDj8oBk6IxOpChO^8YK`wD7_lhEnIioJJk-CZG@MWQFRB$%( zvDk7~Le7}H!y}6N^j>aAM94!{D2cSe>85Fz`U?1_c?A$xol9^g+9A!l5xJINk1?4(oJ3l4_l53}aLEi-!6ksP-lV zG~a!bqZ07X)b<=FaAfz(q^d`c-EG$J6gv^sLp8sYJ*K#EU+Y+B_ +#include + + + +#define SWAP(x0,x) {float *tmp=x0;x0=x;x=tmp;} +#define IX(i,j,k) ((i)+(N)*(j)+(N*N)*(k)) + +#define LINEARSOLVERTIMES 20 + + + + + +void diffuse(int N, int b, float * x, float * x0, float diff, float dt); +void advect(int N, int b, float * d, float * d0, float * u, float * v, float * w, float dt); +void project(int N, float * u, float * v, float * w, float * p, float * div); +void set_bnd(int N, int b, float * x); +void dens_step(int N, float * x, float * x0, float * u, float * v, float * w, float diff, float dt); +void vel_step(int N, float * u, float * v, float * w, float * u0, float * v0, float * w0, float visc, float dt); +void lin_solve(int N, int b, float* x, float* x0, float a, float c); + +JNIEXPORT void JNICALL Java_electrosphere_FluidSim_simulate( + JNIEnv * env, + jobject this, + jint DIM_X, + jint DIM_Y, + jint DIM_Z, + jfloatArray jx, + jfloatArray jx0, + jfloatArray ju, + jfloatArray jv, + jfloatArray jw, + jfloatArray ju0, + jfloatArray jv0, + jfloatArray jw0, + jfloat DIFFUSION_RATE, + jfloat VISCOSITY_RATE, + jfloat timestep){ + jboolean isCopy; + float * x = (*env)->GetFloatArrayElements(env,jx,&isCopy); + float * x0 = (*env)->GetFloatArrayElements(env,jx0,&isCopy); + float * u = (*env)->GetFloatArrayElements(env,ju,&isCopy); + float * v = (*env)->GetFloatArrayElements(env,jv,&isCopy); + float * w = (*env)->GetFloatArrayElements(env,jw,&isCopy); + float * u0 = (*env)->GetFloatArrayElements(env,ju0,&isCopy); + float * v0 = (*env)->GetFloatArrayElements(env,jv0,&isCopy); + float * w0 = (*env)->GetFloatArrayElements(env,jw0,&isCopy); + vel_step(DIM_X, u, v, w, u0, v0, w0, VISCOSITY_RATE, timestep); + dens_step(DIM_X, x, x0, u, v, w, DIFFUSION_RATE, timestep); + (*env)->SetFloatArrayRegion(env,jx,0,DIM_X*DIM_X*DIM_X,x); + (*env)->SetFloatArrayRegion(env,jx0,0,DIM_X*DIM_X*DIM_X,x0); + (*env)->SetFloatArrayRegion(env,ju,0,DIM_X*DIM_X*DIM_X,u); + (*env)->SetFloatArrayRegion(env,jv,0,DIM_X*DIM_X*DIM_X,v); + (*env)->SetFloatArrayRegion(env,jw,0,DIM_X*DIM_X*DIM_X,w); + (*env)->SetFloatArrayRegion(env,ju0,0,DIM_X*DIM_X*DIM_X,u0); + (*env)->SetFloatArrayRegion(env,jv0,0,DIM_X*DIM_X*DIM_X,v0); + (*env)->SetFloatArrayRegion(env,jw0,0,DIM_X*DIM_X*DIM_X,w0); + // for(int i=0; iSetFloatArrayRegion(env,jx,0,DIM_X*DIM_X*DIM_X,x); + // // jx[IX(i,j,k)] = x[IX(i,j,k)]; + // }}} +} + + + +int main() { + printf("Hello World from Project Panama Baeldung Article"); + return 0; +} + + +void diffuse(int N, int b, float * x, float * x0, float diff, float dt){ + float a=dt*diff*N*N*N; + lin_solve(N, b, x, x0, a, 1+6*a); +} + +void advect(int N, int b, float * d, float * d0, float * u, float * v, float * w, float dt){ + int i, j, k, i0, j0, k0, i1, j1, k1; + float x, y, z, s0, t0, s1, t1, u1, u0, dtx,dty,dtz; + + dtx=dty=dtz=dt*N; + + for ( i=1 ; iN+0.5f) x=N+0.5f; i0=(int)x; i1=i0+1; + if (y<0.5f) y=0.5f; if (y>N+0.5f) y=N+0.5f; j0=(int)y; j1=j0+1; + if (z<0.5f) z=0.5f; if (z>N+0.5f) z=N+0.5f; k0=(int)z; k1=k0+1; + + s1 = x-i0; s0 = 1-s1; t1 = y-j0; t0 = 1-t1; u1 = z-k0; u0 = 1-u1; + if(i0 >= N){ + i0 = N - 1; + } + if(j0 >= N){ + j0 = N - 1; + } + if(k0 >= N){ + k0 = N - 1; + } + if(i1 >= N){ + i1 = N - 1; + } + if(j1 >= N){ + j1 = N - 1; + } + if(k1 >= N){ + k1 = N - 1; + } + d[IX(i,j,k)] = s0*(t0*u0*d0[IX(i0,j0,k0)]+t1*u0*d0[IX(i0,j1,k0)]+t0*u1*d0[IX(i0,j0,k1)]+t1*u1*d0[IX(i0,j1,k1)])+ + s1*(t0*u0*d0[IX(i1,j0,k0)]+t1*u0*d0[IX(i1,j1,k0)]+t0*u1*d0[IX(i1,j0,k1)]+t1*u1*d0[IX(i1,j1,k1)]); + }}} + set_bnd(N, b, d); +} + +void dens_step(int N, float * x, float * x0, float * u, float * v, float * w, float diff, float dt){ + // add_source ( N, x, x0, dt ); + SWAP ( x0, x ); diffuse ( N, 0, x, x0, diff, dt ); + SWAP ( x0, x ); advect ( N, 0, x, x0, u, v, w, dt ); +} + +void vel_step(int N, float * u, float * v, float * w, float * u0, float * v0, float * w0, float visc, float dt){ + // add_source ( N, u, u0, dt ); add_source ( N, v, v0, dt ); + // SWAP ( u0, u ); diffuse ( N, 1, u, u0, visc, dt ); + // SWAP ( v0, v ); diffuse ( N, 2, v, v0, visc, dt ); + // project ( N, u, v, w, u0, v0 ); + // SWAP ( u0, u ); SWAP ( v0, v ); + // advect ( N, 1, u, u0, u0, v0, dt ); advect ( N, 2, v, v0, u0, v0, dt ); + // project ( N, u, v, w, u0, v0 ); + + // add_source(N, v, x, dt); + SWAP(u0, u); + diffuse(N, 1, u, u0, visc, dt); + SWAP(v0, v); + diffuse(N, 2, v, v0, visc, dt); + SWAP(w0, w); + diffuse(N, 3, w, w0, visc, dt); + project(N, u, v, w, u0, v0); + SWAP(u0, u); + SWAP(v0, v); + SWAP(w0, w); + advect(N, 1, u, u0, u0, v0, w0, dt); + advect(N, 2, v, v0, u0, v0, w0, dt); + advect(N, 3, w, w0, u0, v0, w0, dt); + project(N, u, v, w, u0, v0); +} + +void project(int N, float * u, float * v, float * w, float * p, float * div){ + int i, j, k; + + for ( i=1 ; i +/* Header for class electrosphere_FluidSim */ + +#ifndef _Included_electrosphere_FluidSim +#define _Included_electrosphere_FluidSim +#ifdef __cplusplus +extern "C" { +#endif +#undef electrosphere_FluidSim_DIM +#define electrosphere_FluidSim_DIM 18L +#undef electrosphere_FluidSim_DIFFUSION_CONSTANT +#define electrosphere_FluidSim_DIFFUSION_CONSTANT 1.0E-4f +#undef electrosphere_FluidSim_VISCOSITY_CONSTANT +#define electrosphere_FluidSim_VISCOSITY_CONSTANT 1.0E-4f +#undef electrosphere_FluidSim_LINEARSOLVERTIMES +#define electrosphere_FluidSim_LINEARSOLVERTIMES 20L +#undef electrosphere_FluidSim_GRAVITY +#define electrosphere_FluidSim_GRAVITY -10000.0f +/* + * Class: electrosphere_FluidSim + * Method: simulate + * Signature: (III[F[F[F[F[F[F[F[FFFF)V + */ +JNIEXPORT void JNICALL Java_electrosphere_FluidSim_simulate + (JNIEnv *, jobject, jint, jint, jint, jfloatArray, jfloatArray, jfloatArray, jfloatArray, jfloatArray, jfloatArray, jfloatArray, jfloatArray, jfloat, jfloat, jfloat); + +#ifdef __cplusplus +} +#endif +#endif diff --git a/src/main/java/electrosphere/FluidSim.java b/src/main/java/electrosphere/FluidSim.java index 3eddf66..969721b 100644 --- a/src/main/java/electrosphere/FluidSim.java +++ b/src/main/java/electrosphere/FluidSim.java @@ -9,36 +9,21 @@ import java.util.List; import java.util.Random; import java.util.Set; -import jdk.incubator.foreign.MemorySegment; -import jdk.incubator.vector.FloatVector; -import jdk.incubator.vector.VectorMask; -import jdk.incubator.vector.VectorOperators; -import jdk.incubator.vector.VectorOperators.Associative; - import org.joml.Vector2i; import org.lwjgl.BufferUtils; import org.lwjgl.PointerBuffer; -import org.lwjgl.opencl.CL30; import org.lwjgl.system.MemoryUtil; -import electrosphere.opencl.CLBuffer; -import electrosphere.opencl.CLCommandQueue; -import electrosphere.opencl.CLContext; -import electrosphere.opencl.CLDevice; -import electrosphere.opencl.CLEvent; -import electrosphere.opencl.CLImage; -import electrosphere.opencl.CLKernel; -import electrosphere.opencl.CLPlatform; -import electrosphere.opencl.CLProgram; -import electrosphere.opencl.CLReadImageResult; -import electrosphere.opencl.CLUtilities; -import electrosphere.opencl.LWJGLContext; - /** * Simulates a fluid via opencl */ public class FluidSim { + static { + System.out.println(System.getProperty("user.dir")); + System.load(System.getProperty("user.dir") + "/shared-folder/libfluidsim.dll"); + } + public static final int DIM = 18; float[] x = new float[DIM * DIM * DIM]; @@ -118,10 +103,11 @@ public class FluidSim { } } + simulate(DIM, DIM, DIM, x, x0, u, v, w, u0, v0, w0, DIFFUSION_CONSTANT, VISCOSITY_CONSTANT, timestep); - vel_step(DIM, DIM, DIM, u, v, w, u0, v0, w0, VISCOSITY_CONSTANT, timestep); + // vel_step(DIM, DIM, DIM, u, v, w, u0, v0, w0, VISCOSITY_CONSTANT, timestep); - dens_step(DIM, DIM, DIM, x, x0, u, v, w, DIFFUSION_CONSTANT, timestep); + // dens_step(DIM, DIM, DIM, x, x0, u, v, w, DIFFUSION_CONSTANT, timestep); sum(); } @@ -144,7 +130,7 @@ public class FluidSim { void diffuse(int M, int N, int O, int b, float[] x, float[] x0, float diff, float dt){ int max = Math.max(Math.max(M, N), Math.max(N, O)); float a=dt*diff*max*max*max; - linSolveJacobian(M, N, O, b, x, x0, a, 1+6*a); + lin_solve(M, N, O, b, x, x0, a, 1+6*a); } void setBounds(float[] target, int b){ @@ -245,7 +231,7 @@ public class FluidSim { setBounds(div, 0); setBounds(p, 0); - linSolveJacobian(M, N, O, 0, p, div, 1, 6); + lin_solve(M, N, O, 0, p, div, 1, 6); for ( i=1 ; i mask = FloatVector.SPECIES_256.indexInRange(O, numLeft); - - // for ( i=1 ; i waitList, - Vector2i workersInEachDimension, - Vector2i workersInEachGroup - ){ - long[] dimensionArray = new long[]{ - workersInEachDimension.x, - workersInEachDimension.y - }; - long[] groupArray = new long[]{ - workersInEachGroup.x, - workersInEachGroup.y - }; - //dimensionality must be 2 as the work arrays are 2D vectors - CLEvent event = queueKernelWork(kernel,waitList,2,dimensionArray,groupArray); - return event; - } - - /** - * Queues kernel work with a default worker offset vector - * @param kernel The kernel to queue work on - * @param waitList The waitlist of events to wait for before executing this work - * @param workDimensions The dimensionality of the data to act upon - * @param workersInEachDimension The total number of workers in each dimension - * @param workersInEachGroup The total number of workers in each work group - * @return The opencl event for this work item - */ - public CLEvent queueKernelWork( - CLKernel kernel, - List waitList, - int workDimensions, - long[] workersInEachDimension, - long[] workersInEachGroup - ){ - return queueKernelWork(kernel, waitList, workDimensions, null, workersInEachDimension, workersInEachGroup); - } - - /** - * Queues work on a given opencl kernel - * @param kernel The kernel to queue work on - * @param waitList The waitlist of events to wait for before executing this work - * @param workDimensions The dimensionality of the data to act upon - * @param globalWorkerOffsetVector The vector specifying the amount of offset per dimension between given global workers - * @param workersInEachDimension The total number of workers in each dimension - * @param workersInEachGroup The total number of workers in each work group - * @return The opencl event for this work item - */ - public CLEvent queueKernelWork( - CLKernel kernel, - List waitList, - int workDimensions, - long[] globalWorkerOffsetVector, //basically a vector that describes the increment in each dimension. Unlikely to see use generally - long[] workersInEachDimension, //basically the number of sub arrays - long[] workersInEachGroup // basically the size of each subarray to work one - ){ - CLEvent rVal = null; - try(MemoryStack stack = MemoryStack.stackPush()){ - // - //construct and fill buffers for opencl call - - //global worker offsets - PointerBuffer globalWorkerOffset = null; - if(globalWorkerOffsetVector != null){ - globalWorkerOffset = stack.mallocPointer(workDimensions); - globalWorkerOffset.put(globalWorkerOffsetVector); - globalWorkerOffset.flip(); - } - - //global worker size - PointerBuffer globalWorkerSize = null; - globalWorkerSize = stack.mallocPointer(workDimensions); - globalWorkerSize.put(workersInEachDimension); - globalWorkerSize.flip(); - - //local worker size - PointerBuffer localWorkerSize = stack.mallocPointer(workDimensions); - localWorkerSize.put(workersInEachGroup); - localWorkerSize.flip(); - - //the events to wait for before executing this one - PointerBuffer eventWaitList = null; - if(waitList != null && waitList.size() > 0){ - eventWaitList = stack.mallocPointer(waitList.size()); - for(CLEvent toWaitEvent : waitList){ - eventWaitList.put(toWaitEvent.getPointer()); - } - } - PointerBuffer event = stack.mallocPointer(1); - - // - //queue work - CL30.clEnqueueNDRangeKernel( - this.pointer, - kernel.getPointer(), - workDimensions, - globalWorkerOffset, - globalWorkerSize, - localWorkerSize, - eventWaitList, - event - ); - - //create returned event - rVal = new CLEvent(event.get()); - } - return rVal; - } - - /** - * Sets an argument for a given kernel - * @param kernel The kernel - * @param argumentIndex The index of the argument (0-index'd) - * @param argumentValue The value to set the argument to - * @return The result of the argument set operation - */ - public int setKernelArgument(CLKernel kernel, int argumentIndex, long argumentValue){ - return CL30.clSetKernelArg(kernel.getPointer(),argumentIndex,argumentValue); - } - - /** - * Sets an argument for a given kernel to be equal to a given image - * @param kernel The kernel - * @param argumentIndex The index of the argument (0-index'd) - * @param image The value to set the argument to - * @return The result of the argument set operation - */ - public int setKernelArgument(CLKernel kernel, int argumentIndex, CLImage image){ - int rVal = 0; - try(MemoryStack stack = MemoryStack.stackPush()){ - PointerBuffer pointerBuffer = stack.mallocPointer(1); - pointerBuffer.put(image.getPointer()); - pointerBuffer.flip(); - rVal = CL30.clSetKernelArg(kernel.getPointer(),argumentIndex,pointerBuffer); - } - return rVal; - } - - /** - * Reads an image from opencl global memory - * @param image The image - * @param waitList The wait list of events to make this query pend on - * @return The byte buffer containing the data of the image - */ - public CLReadImageResult readImage(CLImage image, List waitList){ - CLReadImageResult rVal = null; - - //the cpu buffer to write into from gpu memory - int bufferSize = (int)(image.getWidth() * image.getByteWidth()); - if(image.getHeight() > 0){ - bufferSize = (int)(bufferSize * image.getHeight()); - } - if(image.getDepth() > 0){ - bufferSize = (int)(bufferSize * image.getDepth()); - } - ByteBuffer outputBuffer = MemoryUtil.memAlloc(bufferSize); - - try(MemoryStack stack = MemoryStack.stackPush()){ - boolean blocking = false; //should the call block execution until the image is read? - - //defines the offset in outputBuffer to start writing (should be a vec) - PointerBuffer origin = stack.mallocPointer(3); - origin.put(0); - origin.put(0); - origin.put(0); - origin.flip(); - - //The dimensions of the image - PointerBuffer region = stack.mallocPointer(3); - region.put(image.getWidth()); - if(image.getHeight() == 0){ - region.put(1); //If it's a 1D image, this should be set to 1 - } else { - region.put(image.getHeight()); - } - if(image.getDepth() == 0){ - region.put(1); //If it's a 2D image, this should be set to 1 - } else { - region.put(image.getDepth()); - } - region.flip(); - - //how many bytes make up 1 line of the data - long rowPitch = image.getByteWidth() * image.getWidth(); - - //how many bytes make up a single 2d plane of pixels. Should be 0 if image is 1d or 2d - long slicePitch = 0; - if(image.getDepth() > 0){ - slicePitch = image.getByteWidth() * image.getWidth() * image.getHeight(); - } - - - //the event management buffers - PointerBuffer eventWaitList = null; - if(waitList != null && waitList.size() > 0){ - eventWaitList = stack.mallocPointer(waitList.size()); - for(CLEvent toWaitEvent : waitList){ - eventWaitList.put(toWaitEvent.getPointer()); - } - eventWaitList.flip(); - } - PointerBuffer event = stack.mallocPointer(1); - - //actually send to opencl - int errorCode = CL30.clEnqueueReadImage( - this.getPointer(), - image.getPointer(), - blocking, - origin, - region, - rowPitch, - slicePitch, - outputBuffer, - eventWaitList, - event - ); - - //error check - CLUtilities.checkCLError(errorCode); - - //construct return object - CLEvent returnEvent = new CLEvent(event.get()); - rVal = new CLReadImageResult(outputBuffer,returnEvent); - } - return rVal; - } - - /** - * Writes an array of bytes to a device-side image in opencl - * @param image The image to write to - * @param data The data to write - * @param waitList The waitlist of events to pend this write operation on - * @return The CLEvent encapsulating this write operation - */ - public CLEvent writeToDeviceImage(CLImage image, List waitList, int[] data){ - CLEvent rVal = null; - try(MemoryStack stack = MemoryStack.stackPush()){ - boolean blocking = true; - - //defines the offset in outputBuffer to start writing (should be a vec) - PointerBuffer origin = stack.mallocPointer(3); - origin.put(0); - origin.put(0); - origin.put(0); - origin.flip(); - - //The dimensions of the image - PointerBuffer region = stack.mallocPointer(3); - region.put(image.getWidth()); - if(image.getHeight() == 0){ - region.put(1); //If it's a 1D image, this should be set to 1 - } else { - region.put(image.getHeight()); - } - if(image.getDepth() == 0){ - region.put(1); //If it's a 2D image, this should be set to 1 - } else { - region.put(image.getDepth()); - } - region.flip(); - - //how many bytes make up 1 line of the data - long rowPitch = image.getByteWidth() * image.getWidth(); - - //how many bytes make up a single 2d plane of pixels. Should be 0 if image is 1d or 2d - long slicePitch = 0; - if(image.getDepth() > 0){ - slicePitch = image.getByteWidth() * image.getWidth() * image.getHeight(); - } - - //the event management buffers - PointerBuffer eventWaitList = null; - if(waitList != null && waitList.size() > 0){ - eventWaitList = stack.mallocPointer(waitList.size()); - for(CLEvent toWaitEvent : waitList){ - eventWaitList.put(toWaitEvent.getPointer()); - } - eventWaitList.flip(); - } - PointerBuffer event = stack.mallocPointer(1); - - int errorCode = CL30.clEnqueueWriteImage(this.getPointer(), image.getPointer(), blocking, origin, region, rowPitch, slicePitch, data, eventWaitList, event); - - //error check - CLUtilities.checkCLError(errorCode); - - //construct return event - rVal = new CLEvent(event.get()); - } - return rVal; - } - - /** - * Writes a buffer of bytes to a device-side image in opencl - * @param image The image to write to - * @param data The data to write - * @param waitList The waitlist of events to pend this write operation on - * @return The CLEvent encapsulating this write operation - */ - public CLEvent writeToDeviceImage(CLImage image, List waitList, ByteBuffer data){ - CLEvent rVal = null; - - try(MemoryStack stack = MemoryStack.stackPush()){ - boolean blocking = false; - - //defines the offset in outputBuffer to start writing (should be a vec) - PointerBuffer origin = stack.mallocPointer(3); - origin.put(0); - origin.put(0); - origin.put(0); - origin.flip(); - - //The dimensions of the image - PointerBuffer region = stack.mallocPointer(3); - region.put(image.getWidth()); - if(image.getHeight() == 0){ - region.put(1); //If it's a 1D image, this should be set to 1 - } else { - region.put(image.getHeight()); - } - if(image.getDepth() == 0){ - region.put(1); //If it's a 2D image, this should be set to 1 - } else { - region.put(image.getDepth()); - } - region.flip(); - - //how many bytes make up 1 line of the data - long rowPitch = image.getByteWidth() * image.getWidth(); - - //how many bytes make up a single 2d plane of pixels. Should be 0 if image is 1d or 2d - long slicePitch = 0; - if(image.getDepth() > 0){ - slicePitch = image.getByteWidth() * image.getWidth() * image.getHeight(); - } - - //the event management buffers - PointerBuffer eventWaitList = null; - if(waitList != null && waitList.size() > 0){ - eventWaitList = stack.mallocPointer(waitList.size()); - for(CLEvent toWaitEvent : waitList){ - eventWaitList.put(toWaitEvent.getPointer()); - } - eventWaitList.flip(); - } - PointerBuffer event = stack.mallocPointer(1); - - int errorCode = CL30.clEnqueueWriteImage(this.getPointer(), image.getPointer(), blocking, origin, region, rowPitch, slicePitch, data, eventWaitList, event); - - //error check - CLUtilities.checkCLError(errorCode); - - //construct return event - rVal = new CLEvent(event.get()); - } - - return rVal; - } - - /** - * Blocks thread execution until the event list has completed or otherwise resolved - * @param eventList The event list - */ - public void waitForEventList(List waitList){ - try(MemoryStack stack = MemoryStack.stackPush()){ - PointerBuffer eventWaitList = null; - if(waitList != null && waitList.size() > 0){ - eventWaitList = stack.mallocPointer(waitList.size()); - for(CLEvent toWaitEvent : waitList){ - eventWaitList.put(toWaitEvent.getPointer()); - } - eventWaitList.flip(); - } - - //wait - int errorCode = CL30.clWaitForEvents(eventWaitList); - - //error check - CLUtilities.checkCLError(errorCode); - } - } - - - -} diff --git a/src/main/java/electrosphere/opencl/CLContext.java b/src/main/java/electrosphere/opencl/CLContext.java deleted file mode 100644 index 9175435..0000000 --- a/src/main/java/electrosphere/opencl/CLContext.java +++ /dev/null @@ -1,164 +0,0 @@ -package electrosphere.opencl; - -import java.nio.ByteBuffer; -import java.nio.IntBuffer; - -import org.lwjgl.BufferUtils; -import org.lwjgl.PointerBuffer; -import org.lwjgl.opencl.CL30; -import org.lwjgl.opencl.CLContextCallback; -import org.lwjgl.system.MemoryUtil; - -/** - * Wrapper object for an opencl context - */ -public class CLContext { - - //the context pointer - long pointer; - - //the context callback - CLContextCallback clContextCB; - //the error code buffer - IntBuffer errorCodeBuffer = BufferUtils.createIntBuffer(1); - int[] errorCodeArray = new int[1]; - - //the device this context is created on - CLDevice device; - - //creates an opencl context - public CLContext(CLDevice device, PointerBuffer ctxProps){ - pointer = CL30.clCreateContext(ctxProps, device.getId(), clContextCB = CLContextCallback.create( - (errinfo, private_info, cb, user_data) -> System.out.println(String.format("clCreateContext\n\tInfo: %s", MemoryUtil.memUTF8(errinfo))) - ), MemoryUtil.NULL, errorCodeBuffer); - this.device = device; - } - - /** - * Gets the pointer for this context - * @return The pointer - */ - public long getPointer(){ - return pointer; - } - - /** - * Gets the opencl error code storage buffer - * @return The opencl error code storage buffer - */ - protected IntBuffer getErrorCodeBuffer(){ - return errorCodeBuffer; - } - - /** - * Gets an array used for storing error codes in relevant api calls - * @return The array - */ - protected int[] getErrorCodeArray(){ - return errorCodeArray; - } - - /** - * Creates an opencl buffer with a specified size - * @param flags The flags for the buffer - * @param size The size of the opencl buffer in bytes - * @return The opencl buffer - */ - public CLBuffer createBuffer(long flags, long size){ - CLBuffer buffer = new CLBuffer(this, flags, size); - CLUtilities.checkCLError(errorCodeBuffer); - return buffer; - } - - /** - * Creates an opencl buffer based on a cpu buffer - * @param flags The flags for the buffer - * @param hostBuffer The cpu buffer to base the opencl buffer on - * @return The opencl buffer - */ - public CLBuffer createBuffer(long flags, IntBuffer hostBuffer){ - CLBuffer buffer = new CLBuffer(this, flags, hostBuffer); - CLUtilities.checkCLError(errorCodeBuffer); - return buffer; - } - - /** - * Creates a 2d opencl image buffer - * @param flags The flags for creation of the image - * @param width The width of the image - * @param height The height of the image - * @return The opencl image wrapper object - */ - public CLImage createImage(long flags, long width, long height){ - CLImage image = new CLImage(this, flags, width, height); - CLUtilities.checkCLError(errorCodeBuffer); - return image; - } - - /** - * Creates a 3d opencl image buffer - * @param flags The flags for creation of the image - * @param width The width of the image - * @param height The height of the image - * @param depth The depth of the image - * @return The opencl image wrapper object - */ - public CLImage createImage(long flags, long width, long height, long depth){ - CLImage image = new CLImage(this, flags, width, height, depth); - CLUtilities.checkCLError(errorCodeBuffer); - return image; - } - - /** - * Creates a 3d opencl image buffer - * @param flags The flags for creation of the image - * @param width The width of the image - * @param height The height of the image - * @param depth The depth of the image - * @param data The data to init the image with - * @return The opencl image wrapper object - */ - public CLImage createImage(long flags, long width, long height, long depth, ByteBuffer data){ - CLImage image = new CLImage(this, flags, width, height, depth, data); - CLUtilities.checkCLError(errorCodeBuffer); - return image; - } - - /** - * Creates a command queue for this context - * @return The command queue - */ - public CLCommandQueue createCommandQueue(LWJGLContext lwjglContext){ - return new CLCommandQueue(lwjglContext, this); - } - - /** - * Gets the opencl device this context was created on - * @return The device - */ - public CLDevice getDevice(){ - return device; - } - - /** - * Creates an opencl program with a provided source - * @param source The source - * @return The program - */ - public CLProgram createProgram(String source){ - CLProgram program = CLProgram.createProgram(this, source); - return program; - } - - /** - * Creates an opencl kernel from the specified program - * @param program The program - * @param kernelName The name of the kernel in the program source - * @return The kernel object - */ - public CLKernel createKernel(CLProgram program, String kernelName){ - CLKernel kernel = new CLKernel(this, program, kernelName); - return kernel; - } - -} diff --git a/src/main/java/electrosphere/opencl/CLDevice.java b/src/main/java/electrosphere/opencl/CLDevice.java deleted file mode 100644 index cd58076..0000000 --- a/src/main/java/electrosphere/opencl/CLDevice.java +++ /dev/null @@ -1,68 +0,0 @@ -package electrosphere.opencl; - -import org.lwjgl.opencl.CL30; - -/** - * Wrapper object for an opencl device - */ -public class CLDevice { - - //the id of the device - long id; - - //the name of the device - String deviceName; - - //true if the device is a gpu - boolean isGpu = false; - - //true if the device is available - boolean isAvailable = false; - - /** - * Creates a CL device wrapper object - * @param id the id of the device - * @param isGpu true if the device is a gpu - */ - protected CLDevice(long id){ - this.id = id; - //get device name - deviceName = CLUtilities.getDeviceInfoStringUTF8(id, CL30.CL_DEVICE_NAME); - //get if device is gpu - long deviceType = CLUtilities.getDeviceInfoLong(id, CL30.CL_DEVICE_TYPE); - if (deviceType == CL30.CL_DEVICE_TYPE_GPU) { - isGpu = true; - } - //check if is available - long isAvailableRaw = CLUtilities.getDeviceInfoLong(id, CL30.CL_DEVICE_AVAILABLE); - if(isAvailableRaw == CL30.CL_TRUE){ - isAvailable = true; - } - - } - - /** - * Gets the id of the device - * @return the id of the device - */ - public long getId(){ - return id; - } - - /** - * True if the device is a gpu - * @return True if the device is a gpu - */ - public boolean isGpu(){ - return isGpu; - } - - /** - * Gets the name of the device - * @return The name of the device - */ - public String getDeviceName(){ - return deviceName; - } - -} diff --git a/src/main/java/electrosphere/opencl/CLEvent.java b/src/main/java/electrosphere/opencl/CLEvent.java deleted file mode 100644 index 4e13222..0000000 --- a/src/main/java/electrosphere/opencl/CLEvent.java +++ /dev/null @@ -1,65 +0,0 @@ -package electrosphere.opencl; - -import java.nio.ByteBuffer; - -import org.lwjgl.BufferUtils; -import org.lwjgl.PointerBuffer; -import org.lwjgl.opencl.CL30; -import org.lwjgl.system.MemoryStack; -import org.lwjgl.system.MemoryUtil; - -/** - * Wrapper for an opencl event - */ -public class CLEvent { - - //The pointer for the event - long pointer; - - /** - * Constructs a wrapper object around an opencl event - * @param pointer The pointer for the event - */ - protected CLEvent(long pointer){ - this.pointer = pointer; - } - - /** - * Gets the pointer for this event - * @return The pointer as a long - */ - public long getPointer(){ - return pointer; - } - - - /** - * Queries the status of the event from opencl - * @return The event status code - */ - public int getStatus(){ - int rVal = 0; - try(MemoryStack stack = MemoryStack.stackPush()){ - //create buffers for call - ByteBuffer returnValueBuffer = stack.malloc(1); - PointerBuffer sizeBuffer = stack.mallocPointer(1); - - //submit request to opencl - int errorCode = CL30.clGetEventInfo(this.getPointer(),CL30.CL_EVENT_COMMAND_EXECUTION_STATUS,returnValueBuffer,sizeBuffer); - - //check error - CLUtilities.checkCLError(errorCode); - - //get return value - rVal = returnValueBuffer.getInt(); - - // //free buffers - // MemoryUtil.memFree(returnValueBuffer); - // sizeBuffer.free(); - } - - //return - return rVal; - } - -} diff --git a/src/main/java/electrosphere/opencl/CLImage.java b/src/main/java/electrosphere/opencl/CLImage.java deleted file mode 100644 index fd8e4ad..0000000 --- a/src/main/java/electrosphere/opencl/CLImage.java +++ /dev/null @@ -1,218 +0,0 @@ -package electrosphere.opencl; - -import java.nio.ByteBuffer; -import java.nio.IntBuffer; - -import org.lwjgl.BufferUtils; -import org.lwjgl.opencl.CL30; -import org.lwjgl.opencl.CLImageDesc; -import org.lwjgl.opencl.CLImageFormat; -import org.lwjgl.system.MemoryUtil; - -/** - * Wrapper object for an opencl image - */ -public class CLImage { - - //the pointer to the image - long pointer; - - //The dimensions of the image - long width; - long height; - long depth; - - //true if device can read from the image - boolean canRead; - //true if device can write to the image - boolean canWrite; - - //the number of bytes per pixel. This is used by the command queue fetching tasks to allocate buffer sizes - long byteWidth; - - /** - * Creates an empty 2d image - * @param context The context to create the image in - * @param flags The flags for the image - * @param width The width of the image - * @param height The height of the image - */ - protected CLImage(CLContext context, long flags, long width, long height){ - //set initial values - this.width = width; - this.height = height; - this.depth = 0; - - //parse flags - this.parseFlags(flags); - - // - //Get the image format - CLImageFormat imageFormat = CLImageFormat.create(); - imageFormat.image_channel_data_type(CL30.CL_SNORM_INT8); - byteWidth = 4; - imageFormat.image_channel_order(CL30.CL_RGBA); - - // - //Get the image description - CLImageDesc imageDesc = CLImageDesc.create(); - imageDesc.image_type(CL30.CL_MEM_OBJECT_IMAGE2D); - imageDesc.image_width(width); - imageDesc.image_height(height); - imageDesc.image_depth(0); //only used for 3d image - imageDesc.image_array_size(1); //only used for 1d and 2d images - imageDesc.image_row_pitch(0); //must be 0 if host ptr is null - imageDesc.image_slice_pitch(0); //must be 0 if host ptr is null - - // - //host pointer (the data to prepopulate with) - IntBuffer hostPointerBuffer = null; //BufferUtils.createIntBuffer(1); - - // - //Create the image object - pointer = CL30.clCreateImage(context.getPointer(), flags, imageFormat, imageDesc, hostPointerBuffer, context.getErrorCodeBuffer()); - - //error check - CLUtilities.checkCLError(context.getErrorCodeBuffer()); - } - - /** - * Creates an empty 3d image - * @param context The context to create the image in - * @param flags The flags for the image - * @param width The width of the image - * @param height The height of the image - * @param depth The depth of the image - */ - protected CLImage(CLContext context, long flags, long width, long height, long depth){ - //set initial values - this.width = width; - this.height = height; - this.depth = depth; - - //parse flags - this.parseFlags(flags); - - // - //Get the image format - CLImageFormat imageFormat = CLImageFormat.create(); - imageFormat.image_channel_data_type(CL30.CL_SNORM_INT8); - byteWidth = 4; - imageFormat.image_channel_order(CL30.CL_RGBA); - - // - //Get the image description - CLImageDesc imageDesc = CLImageDesc.create(); - imageDesc.image_type(CL30.CL_MEM_OBJECT_IMAGE3D); - imageDesc.image_width(width); - imageDesc.image_height(height); - imageDesc.image_depth(depth); //only used for 3d image - imageDesc.image_array_size(1); //only used for 1d and 2d images - imageDesc.image_row_pitch(0); //must be 0 if host ptr is null - imageDesc.image_slice_pitch(0); //must be 0 if host ptr is null - - // - //host pointer (the data to prepopulate with) - IntBuffer hostPointerBuffer = null; //BufferUtils.createIntBuffer(1); - - // - //Create the image object - pointer = CL30.clCreateImage(context.getPointer(), flags, imageFormat, imageDesc, hostPointerBuffer, context.getErrorCodeBuffer()); - } - - /** - * Creates an empty 3d image - * @param context The context to create the image in - * @param flags The flags for the image - * @param width The width of the image - * @param height The height of the image - * @param depth The depth of the image - */ - protected CLImage(CLContext context, long flags, long width, long height, long depth, ByteBuffer data){ - //set initial values - this.width = width; - this.height = height; - this.depth = depth; - - //parse flags - this.parseFlags(flags); - - // - //Get the image format - CLImageFormat imageFormat = CLImageFormat.create(); - imageFormat.image_channel_data_type(CL30.CL_SNORM_INT8); - byteWidth = 4; - imageFormat.image_channel_order(CL30.CL_RGBA); - - // - //Get the image description - CLImageDesc imageDesc = CLImageDesc.create(); - imageDesc.image_type(CL30.CL_MEM_OBJECT_IMAGE3D); - imageDesc.image_width(width); - imageDesc.image_height(height); - imageDesc.image_depth(depth); //only used for 3d image - imageDesc.image_array_size(1); //only used for 1d and 2d images - imageDesc.image_row_pitch(width * byteWidth); //must be 0 if host ptr is null - imageDesc.image_slice_pitch(width * height * byteWidth); //must be 0 if host ptr is null - - // - //Create the image object - pointer = CL30.clCreateImage(context.getPointer(), flags, imageFormat, imageDesc, data, context.getErrorCodeBuffer()); - - CLUtilities.checkCLError(context); - } - - /** - * Parses the flags passed in on image creation to set variables (eg read/write status) - * @param flags The flags passed in - */ - private void parseFlags(long flags){ - if((CL30.CL_MEM_READ_ONLY & flags) != 0){ - this.canRead = true; - } - if((CL30.CL_MEM_WRITE_ONLY & flags) != 0){ - this.canWrite = true; - } - } - - /** - * Gets the pointer for the image - * @return The image pointer - */ - public long getPointer(){ - return pointer; - } - - /** - * Gets the width of the image - * @return The width of the image - */ - public long getWidth(){ - return width; - } - - /** - * Gets the height of the image - * @return The height of the image - */ - public long getHeight(){ - return height; - } - - /** - * Gets the depth of the image - * @return The depth of the image - */ - public long getDepth(){ - return depth; - } - - /** - * Gets the byte width of the image - * @return The byte width - */ - public long getByteWidth(){ - return byteWidth; - } - -} diff --git a/src/main/java/electrosphere/opencl/CLKernel.java b/src/main/java/electrosphere/opencl/CLKernel.java deleted file mode 100644 index 7efc2d2..0000000 --- a/src/main/java/electrosphere/opencl/CLKernel.java +++ /dev/null @@ -1,32 +0,0 @@ -package electrosphere.opencl; - -import org.lwjgl.opencl.CL30; - -/** - * A wrapper object for an opencl kernel - */ -public class CLKernel { - - //The pointer to the kernel - long pointer; - - /** - * Creates a cl kernel from a context and a program - * @param context The opencl context - * @param program The opencl program - * @param kernelName The name of the kernel - */ - protected CLKernel(CLContext context, CLProgram program, String kernelName){ - pointer = CL30.clCreateKernel(program.getPointer(), kernelName, context.getErrorCodeBuffer()); - CLUtilities.checkCLError(context.getErrorCodeBuffer()); - } - - /** - * Gets the pointer for this kernel - * @return The pointer as a long - */ - public long getPointer(){ - return pointer; - } - -} diff --git a/src/main/java/electrosphere/opencl/CLPlatform.java b/src/main/java/electrosphere/opencl/CLPlatform.java deleted file mode 100644 index f309dea..0000000 --- a/src/main/java/electrosphere/opencl/CLPlatform.java +++ /dev/null @@ -1,153 +0,0 @@ -package electrosphere.opencl; - -import java.nio.ByteBuffer; -import java.nio.CharBuffer; -import java.nio.IntBuffer; -import java.nio.LongBuffer; -import java.util.ArrayList; -import java.util.Collections; -import java.util.HashSet; -import java.util.LinkedList; -import java.util.List; -import java.util.Set; - -import org.lwjgl.PointerBuffer; -import org.lwjgl.opencl.CL; -import org.lwjgl.opencl.CL30; -import org.lwjgl.opencl.CLCapabilities; -import org.lwjgl.system.MemoryStack; - -/** - * An object that encapsulates an openCL platform - */ -public class CLPlatform { - - //the id of the platform - long id; - - //the capabilities of the platform - CLCapabilities caps; - - //Information about the platform - String vendor; - String name; - String version; - - // the list of all cl devices on this platform - List clDevices; - - /** - * Creates a platform object - * @param id The id of the platform - */ - protected CLPlatform(long id, CLCapabilities caps, String vendor, String name, String version, List clDevices){ - this.id = id; - this.caps = caps; - this.vendor = vendor; - this.name = name; - this.version = version; - this.clDevices = clDevices; - } - - - /** - * Gets a set containing all currently available platforms - * @return The set containing all currently available platforms - */ - public static Set getPlatforms(){ - Set platforms = new HashSet(); - - try(MemoryStack stack = MemoryStack.stackPush()){ - //check if there are any platforms - IntBuffer pi = stack.mallocInt(1); - CLUtilities.checkCLError(CL30.clGetPlatformIDs(null, pi)); - if (pi.get(0) == 0) { - throw new IllegalStateException("No OpenCL platforms found."); - } - //if there are platforms, get all platform IDs - PointerBuffer platformIDs = stack.mallocPointer(pi.get(0)); - CLUtilities.checkCLError(CL30.clGetPlatformIDs(platformIDs, (IntBuffer)null)); - - //add all platforms to the returned list - for (int i = 0; i < platformIDs.capacity(); i++) { - //get basic identification - long platformId = platformIDs.get(i); - CLCapabilities caps = CL.createPlatformCapabilities(platformId); - String vendor = CLUtilities.getPlatformInfoStringUTF8(platformId, CL30.CL_PLATFORM_VENDOR); - String name = CLUtilities.getPlatformInfoStringUTF8(platformId, CL30.CL_PLATFORM_NAME); - String version = CLUtilities.getPlatformInfoStringUTF8(platformId, CL30.CL_PLATFORM_VERSION); - - - //get devices on the platform - int errcode = CL30.clGetDeviceIDs(platformId, CL30.CL_DEVICE_TYPE_GPU, null, pi); - List devices = new LinkedList(); - if(errcode != CL30.CL_DEVICE_NOT_FOUND){ - CLUtilities.checkCLError(errcode); - PointerBuffer deviceIDs = stack.mallocPointer(pi.get(0)); - CLUtilities.checkCLError(CL30.clGetDeviceIDs(platformId, CL30.CL_DEVICE_TYPE_ALL, deviceIDs, (IntBuffer)null)); - for (int j = 0; j < deviceIDs.capacity(); j++) { - long deviceId = deviceIDs.get(j); - devices.add(new CLDevice(deviceId)); - } - } - - platforms.add(new CLPlatform(platformId, caps, vendor, name, version, devices)); - } - - if (platforms.isEmpty()) { - throw new IllegalStateException("No OpenCL platform found."); - } - } - - return platforms; - } - - /** - * Gets the id of the platform - * @return the id of the platfomr - */ - public long getId(){ - return id; - } - - /** - * Gets the capabilities object of the platform - * @return The capabilities - */ - public CLCapabilities getCapabilities(){ - return caps; - } - - /** - * Gets the name of the vendor the platform - * @return The name of the vendor of the platform - */ - public String getVendor(){ - return vendor; - } - - /** - * Gets the name of the platform - * @return The name of the platform - */ - public String getName(){ - return name; - } - - /** - * Gets the version of opencl on the platform - * @return The version of opencl on the platform - */ - public String getVersion(){ - return name; - } - - /** - * Gets the list of cl devices on this platform - * @return The list of devices - */ - public List getCLDevices(){ - return clDevices; - } - -} diff --git a/src/main/java/electrosphere/opencl/CLProgram.java b/src/main/java/electrosphere/opencl/CLProgram.java deleted file mode 100644 index f30797a..0000000 --- a/src/main/java/electrosphere/opencl/CLProgram.java +++ /dev/null @@ -1,70 +0,0 @@ -package electrosphere.opencl; - -import java.util.concurrent.CountDownLatch; - -import org.lwjgl.opencl.CL30; -import org.lwjgl.opencl.CLProgramCallback; -import org.lwjgl.system.MemoryUtil; - -/** - * Wrapper object for an opencl program - */ -public class CLProgram { - - //the pointer to the program - long pointer; - - /** - * Creates an opencl program from a context and a source - * @param context The context - * @param source The source - * @return The program - */ - protected static CLProgram createProgram(CLContext context, String source){ - CLProgram program = new CLProgram(); - - //create object - program.pointer = CL30.clCreateProgramWithSource(context.getPointer(), source, context.getErrorCodeBuffer()); - CLUtilities.checkCLError(context.getErrorCodeBuffer()); - - //latch for blocking execution until the program finishes compiling - CountDownLatch latch = new CountDownLatch(1); - - //build the program - String options = ""; - CLProgramCallback buildCallback; - int errcode = CL30.clBuildProgram(program.pointer, context.getDevice().getId(), options, buildCallback = CLProgramCallback.create((programId, user_data) -> { - System.out.println(String.format( - "The cl_program [0x%X] was built %s", - programId, - CLUtilities.getProgramBuildInfoInt(programId, context.getDevice().getId(), CL30.CL_PROGRAM_BUILD_STATUS) == CL30.CL_SUCCESS ? "successfully" : "unsuccessfully" - )); - String log = CLUtilities.getProgramBuildInfoStringASCII(programId, context.getDevice().getId(), CL30.CL_PROGRAM_BUILD_LOG); - if (!log.isEmpty()) { - System.out.println(String.format("BUILD LOG:\n----\n%s\n-----", log)); - } - - latch.countDown(); - }), MemoryUtil.NULL); - CLUtilities.checkCLError(errcode); - - //Block execution until the program finishes building - try { - latch.await(); - } catch (InterruptedException e) { - throw new RuntimeException(e); - } - buildCallback.free(); - - return program; - } - - /** - * Gets the pointer for this program object - * @return The pointer as a long - */ - public long getPointer(){ - return pointer; - } - -} diff --git a/src/main/java/electrosphere/opencl/CLReadImageResult.java b/src/main/java/electrosphere/opencl/CLReadImageResult.java deleted file mode 100644 index ab3081b..0000000 --- a/src/main/java/electrosphere/opencl/CLReadImageResult.java +++ /dev/null @@ -1,41 +0,0 @@ -package electrosphere.opencl; - -import java.nio.ByteBuffer; - -/** - * A pair of an opencl event and the buffer from the read operation. Used so we can pass back both from the call. - */ -public class CLReadImageResult { - - //the buffer - ByteBuffer buffer; - //the event - CLEvent event; - - /** - * Constructs the pair - * @param buffer The buffer - * @param event The event - */ - protected CLReadImageResult(ByteBuffer buffer, CLEvent event){ - this.buffer = buffer; - this.event = event; - } - - /** - * Gets the buffer - * @return The buffer - */ - public ByteBuffer getBuffer(){ - return buffer; - } - - /** - * Gets the event - * @return The event - */ - public CLEvent getEvent(){ - return event; - } - -} diff --git a/src/main/java/electrosphere/opencl/CLUtilities.java b/src/main/java/electrosphere/opencl/CLUtilities.java deleted file mode 100644 index faa4875..0000000 --- a/src/main/java/electrosphere/opencl/CLUtilities.java +++ /dev/null @@ -1,181 +0,0 @@ -package electrosphere.opencl; - -import java.nio.ByteBuffer; -import java.nio.IntBuffer; -import java.nio.LongBuffer; - -import org.lwjgl.BufferUtils; -import org.lwjgl.PointerBuffer; -import org.lwjgl.opencl.CL30; -import org.lwjgl.system.MemoryStack; -import org.lwjgl.system.MemoryUtil; - -import static org.lwjgl.opencl.CL30.*; -import static org.lwjgl.system.MemoryStack.*; -import static org.lwjgl.system.MemoryUtil.*; - -/** - * Utilities for all opencl files (for instance error checking functions) - */ -public class CLUtilities { - - /** - * Checks an opencl error code - * @param errcode The error code - */ - public static void checkCLError(IntBuffer errcode) { - checkCLError(errcode.get(errcode.position())); - } - - /** - * Checks an opencl error code - * @param errcode The error code - */ - public static void checkCLError(int errcode) { - if (errcode != CL_SUCCESS) { - throw new RuntimeException(String.format("OpenCL error [%d]", errcode)); - } - } - - /** - * Checks an opencl error code - * @param context The context containing the error code buffer - */ - public static void checkCLError(CLContext context) { - int code = context.getErrorCodeBuffer().get(); - if (code != CL_SUCCESS) { - throw new RuntimeException(String.format("OpenCL error [%d]", code)); - } - context.getErrorCodeBuffer().flip(); - } - - /** - * Gets a platform info as a utf string - * @param platformId The platform id - * @param paramName The parameter name - * @return The string containing the information - */ - public static String getPlatformInfoStringUTF8(long platformId, int paramName) { - String rVal = ""; - try(MemoryStack stack = MemoryStack.stackPush()){ - PointerBuffer resultSizeBuffer = stack.mallocPointer(1); - checkCLError(CL30.clGetPlatformInfo(platformId, paramName, (ByteBuffer)null, resultSizeBuffer)); - int bytes = (int)resultSizeBuffer.get(0); - - ByteBuffer buffer = stack.malloc(bytes); - checkCLError(CL30.clGetPlatformInfo(platformId, paramName, buffer, null)); - rVal = memUTF8(buffer, bytes - 1); - } - - return rVal; - } - - /** - * Gets an opencl device parameter as a utf8 string - * @param deviceId The device id - * @param paramName The parameter name - * @return The string - */ - public static String getDeviceInfoStringUTF8(long deviceId, int paramName) { - String rVal = ""; - try(MemoryStack stack = MemoryStack.stackPush()){ - PointerBuffer pp = stack.mallocPointer(1); - checkCLError(clGetDeviceInfo(deviceId, paramName, (ByteBuffer)null, pp)); - int bytes = (int)pp.get(0); - - ByteBuffer buffer = stack.malloc(bytes); - checkCLError(clGetDeviceInfo(deviceId, paramName, buffer, null)); - rVal = memUTF8(buffer, bytes - 1); - } - return rVal; - } - - /** - * Gets a device parameter's value as an int - * @param deviceId The id of the device - * @param paramName The parameter name - * @return The int - */ - public static int getDeviceInfoInt(long deviceId, int paramName) { - int rVal = 0; - try(MemoryStack stack = MemoryStack.stackPush()){ - IntBuffer pl = stack.mallocInt(1); - checkCLError(clGetDeviceInfo(deviceId, paramName, pl, null)); - rVal = pl.get(0); - } - return rVal; - } - - /** - * Gets a device's parameter's value as a long - * @param deviceId The device's id - * @param paramName The parameter name - * @return The value as a long - */ - public static long getDeviceInfoLong(long deviceId, int paramName) { - long rVal = 0; - try(MemoryStack stack = MemoryStack.stackPush()){ - LongBuffer pl = stack.mallocLong(1); - checkCLError(clGetDeviceInfo(deviceId, paramName, pl, null)); - rVal = pl.get(0); - } - return rVal; - } - - /** - * Gets a device's parameter's value as a boolean - * @param deviceId The device's id - * @param paramName The parameter name - * @return The value as a boolean - */ - public static boolean getDeviceInfoBoolean(long deviceId, int paramName) { - byte rVal = 0; - try(MemoryStack stack = MemoryStack.stackPush()){ - ByteBuffer pl = stack.malloc(1); - checkCLError(clGetDeviceInfo(deviceId, paramName, pl, null)); - rVal = pl.get(0); - } - return rVal == 0; - } - - /** - * Gets an opencl program build parameter as an int - * @param programId The program id - * @param deviceId The device id - * @param parameterName The parameter name - * @return THe value of the parameter as an int - */ - public static int getProgramBuildInfoInt(long programId, long deviceId, int parameterName) { - int rVal = 0; - try(MemoryStack stack = MemoryStack.stackPush()){ - IntBuffer pl = stack.mallocInt(1); - checkCLError(clGetProgramBuildInfo(programId, deviceId, parameterName, pl, null)); - rVal = pl.get(0); - } - return rVal; - } - - /** - * Gets an opencl program build parameter as a string - * @param programId The program id - * @param deviceId The device id - * @param parameterName The parameter name - * @return The value of the parameter as a string - */ - public static String getProgramBuildInfoStringASCII(long programId, long deviceId, int parameterName) { - String rVal = ""; - try(MemoryStack stack = MemoryStack.stackPush()){ - PointerBuffer pp = stack.mallocPointer(1); - checkCLError(clGetProgramBuildInfo(programId, deviceId, parameterName, (ByteBuffer)null, pp)); - int bytes = (int)pp.get(0); - - ByteBuffer buffer = stack.malloc(bytes); - checkCLError(clGetProgramBuildInfo(programId, deviceId, parameterName, buffer, null)); - - rVal = memASCII(buffer, bytes - 1);; - } - - return rVal; - } - -} diff --git a/src/main/java/electrosphere/opencl/LWJGLContext.java b/src/main/java/electrosphere/opencl/LWJGLContext.java deleted file mode 100644 index 5a92439..0000000 --- a/src/main/java/electrosphere/opencl/LWJGLContext.java +++ /dev/null @@ -1,28 +0,0 @@ -package electrosphere.opencl; - -import org.lwjgl.system.MemoryStack; - -/** - * A context for the lwjgl framework. Contains stack used by opencl api. - */ -public class LWJGLContext { - - //The stack for memory allocation - MemoryStack stack; - - /** - * Constructor - */ - public LWJGLContext(){ - stack = MemoryStack.create(1000 * 1000 * 4 * 4); - } - - /** - * Gets the memory stack for the current context - * @return The lwjgl context's memory stack - */ - public MemoryStack getStack(){ - return stack; - } - -} diff --git a/src/main/java/electrosphere/render/GLFWContext.java b/src/main/java/electrosphere/render/GLFWContext.java index b2f3b1e..b2c7f77 100644 --- a/src/main/java/electrosphere/render/GLFWContext.java +++ b/src/main/java/electrosphere/render/GLFWContext.java @@ -39,14 +39,19 @@ public class GLFWContext { //Maximize it GLFW.glfwMaximizeWindow(window); //grab actual framebuffer - IntBuffer xBuffer = BufferUtils.createIntBuffer(1); - IntBuffer yBuffer = BufferUtils.createIntBuffer(1); - GLFW.glfwGetFramebufferSize(window, xBuffer, yBuffer); - MemoryUtil.memFree(xBuffer); - MemoryUtil.memFree(yBuffer); - - int bufferWidth = xBuffer.get(); - int bufferHeight = yBuffer.get(); + int bufferWidth = 0; + int bufferHeight = 0; + + try(MemoryStack stack = MemoryStack.stackPush()){ + IntBuffer xBuffer = MemoryUtil.memAllocInt(1); + IntBuffer yBuffer = MemoryUtil.memAllocInt(1); + GLFW.glfwGetFramebufferSize(window, xBuffer, yBuffer); + MemoryUtil.memFree(xBuffer); + MemoryUtil.memFree(yBuffer); + + bufferWidth = xBuffer.get(); + bufferHeight = yBuffer.get(); + } // // Attack controls callbacks diff --git a/src/main/java/electrosphere/render/GridCell.java b/src/main/java/electrosphere/render/GridCell.java new file mode 100644 index 0000000..033d0cb --- /dev/null +++ b/src/main/java/electrosphere/render/GridCell.java @@ -0,0 +1,19 @@ +package electrosphere.render; + +import org.joml.Vector3f; + +public class GridCell { + Vector3f[] points = new Vector3f[8]; //array of size 8 + double[] val = new double[8]; //array of size 8 + public void setValues( + Vector3f p1, Vector3f p2, Vector3f p3, Vector3f p4, + Vector3f p5, Vector3f p6, Vector3f p7, Vector3f p8, + double val1, double val2, double val3, double val4, + double val5, double val6, double val7, double val8 + ){ + points[0] = p1; points[1] = p2; points[2] = p3; points[3] = p4; + points[4] = p5; points[5] = p6; points[6] = p7; points[7] = p8; + val[0] = val1; val[1] = val2; val[2] = val3; val[3] = val4; + val[4] = val5; val[5] = val6; val[6] = val7; val[7] = val8; + } +} diff --git a/src/main/java/electrosphere/render/Mesh.java b/src/main/java/electrosphere/render/Mesh.java index 1e6fe62..bda1c20 100644 --- a/src/main/java/electrosphere/render/Mesh.java +++ b/src/main/java/electrosphere/render/Mesh.java @@ -38,7 +38,7 @@ public class Mesh { static Matrix4f projectionMatrix = new Matrix4f(); static Matrix4f model = new Matrix4f().identity(); - static MemoryStack stack = MemoryStack.create(16 * 1000 * 1000); + static MemoryStack stack = MemoryStack.create(32 * 1000 * 1000); public static void meshInitially(FluidSim sim){ //create and bind vao @@ -55,7 +55,7 @@ public class Mesh { GL44.glBindBuffer(GL44.GL_ARRAY_BUFFER, vertBufferPtr); try { int vertexCount = data.vertices.size(); - FloatBuffer vertData = BufferUtils.createFloatBuffer(vertexCount); + FloatBuffer vertData = MemoryUtil.memAllocFloat(vertexCount); for(float vertValue : data.vertices){ vertData.put(vertValue); } @@ -63,7 +63,7 @@ public class Mesh { GL44.glBufferData(GL44.GL_ARRAY_BUFFER,vertData,GL44.GL_STATIC_DRAW); GL44.glVertexAttribPointer(0, 3, GL44.GL_FLOAT, false, 0, 0); GL44.glEnableVertexAttribArray(0); - // MemoryUtil.memFree(vertData); + MemoryUtil.memFree(vertData); } catch (NullPointerException ex){ ex.printStackTrace(); } @@ -77,13 +77,13 @@ public class Mesh { GL44.glBindBuffer(GL44.GL_ELEMENT_ARRAY_BUFFER, faceBufferPtr); elementCount = data.elements.size(); try { - IntBuffer elementArrayBufferData = BufferUtils.createIntBuffer(elementCount); + IntBuffer elementArrayBufferData = MemoryUtil.memAllocInt(elementCount); for(int element : data.elements){ elementArrayBufferData.put(element); } elementArrayBufferData.flip(); GL44.glBufferData(GL45.GL_ELEMENT_ARRAY_BUFFER,elementArrayBufferData,GL44.GL_STATIC_DRAW); - // MemoryUtil.memFree(elementArrayBufferData); + MemoryUtil.memFree(elementArrayBufferData); } catch (NullPointerException ex){ ex.printStackTrace(); } @@ -100,7 +100,7 @@ public class Mesh { int normalCount = data.normals.size() / 3; FloatBuffer NormalArrayBufferData; if(normalCount > 0){ - NormalArrayBufferData = BufferUtils.createFloatBuffer(normalCount * 3); + NormalArrayBufferData = MemoryUtil.memAllocFloat(normalCount * 3); float[] temp = new float[3]; for(float normalValue : data.normals){ NormalArrayBufferData.put(normalValue); @@ -202,67 +202,66 @@ public class Mesh { //generate verts TerrainChunkData data = generateTerrainChunkData(sim.getData()); - try(MemoryStack stackLocal = stack.push()){ - // - //Buffer data to GPU - // - GL44.glBindBuffer(GL44.GL_ARRAY_BUFFER, vertBufferPtr); - try { - int vertexCount = data.vertices.size(); - FloatBuffer vertData = stackLocal.mallocFloat(vertexCount); - for(float vertValue : data.vertices){ - vertData.put(vertValue); - } - vertData.flip(); - GL44.glBufferData(GL44.GL_ARRAY_BUFFER,vertData,GL44.GL_STATIC_DRAW); - GL44.glVertexAttribPointer(0, 3, GL44.GL_FLOAT, false, 0, 0); - GL44.glEnableVertexAttribArray(0); - } catch (NullPointerException ex){ - ex.printStackTrace(); + // + //Buffer data to GPU + // + GL44.glBindBuffer(GL44.GL_ARRAY_BUFFER, vertBufferPtr); + try { + int vertexCount = data.vertices.size(); + FloatBuffer vertData = MemoryUtil.memAllocFloat(vertexCount); + for(float vertValue : data.vertices){ + vertData.put(vertValue); } - + vertData.flip(); + GL44.glBufferData(GL44.GL_ARRAY_BUFFER,vertData,GL44.GL_STATIC_DRAW); + GL44.glVertexAttribPointer(0, 3, GL44.GL_FLOAT, false, 0, 0); + GL44.glEnableVertexAttribArray(0); + MemoryUtil.memFree(vertData); + } catch (NullPointerException ex){ + ex.printStackTrace(); + } + - - // - // FACES - // - GL44.glBindBuffer(GL44.GL_ELEMENT_ARRAY_BUFFER, faceBufferPtr); - elementCount = data.elements.size(); - try { - IntBuffer elementArrayBufferData = stackLocal.mallocInt(elementCount); - for(int element : data.elements){ - elementArrayBufferData.put(element); - } - elementArrayBufferData.flip(); - GL44.glBufferData(GL45.GL_ELEMENT_ARRAY_BUFFER,elementArrayBufferData,GL44.GL_STATIC_DRAW); - } catch (NullPointerException ex){ - ex.printStackTrace(); + // + // FACES + // + GL44.glBindBuffer(GL44.GL_ELEMENT_ARRAY_BUFFER, faceBufferPtr); + elementCount = data.elements.size(); + try { + IntBuffer elementArrayBufferData = MemoryUtil.memAllocInt(elementCount); + for(int element : data.elements){ + elementArrayBufferData.put(element); } - - - - - // - // NORMALS - // - GL44.glBindBuffer(GL44.GL_ARRAY_BUFFER, normBufferPtr); - try { - int normalCount = data.normals.size() / 3; - FloatBuffer NormalArrayBufferData; - if(normalCount > 0){ - NormalArrayBufferData = stackLocal.mallocFloat(normalCount * 3); - for(float normalValue : data.normals){ - NormalArrayBufferData.put(normalValue); - } - NormalArrayBufferData.flip(); - GL44.glBufferData(GL44.GL_ARRAY_BUFFER,NormalArrayBufferData,GL44.GL_STATIC_DRAW); - GL44.glVertexAttribPointer(1, 3, GL44.GL_FLOAT, false, 0, 0); - GL44.glEnableVertexAttribArray(1); + elementArrayBufferData.flip(); + GL44.glBufferData(GL45.GL_ELEMENT_ARRAY_BUFFER,elementArrayBufferData,GL44.GL_STATIC_DRAW); + MemoryUtil.memFree(elementArrayBufferData); + } catch (NullPointerException ex){ + ex.printStackTrace(); + } + + + + // + // NORMALS + // + GL44.glBindBuffer(GL44.GL_ARRAY_BUFFER, normBufferPtr); + try { + int normalCount = data.normals.size() / 3; + FloatBuffer NormalArrayBufferData; + if(normalCount > 0){ + NormalArrayBufferData = MemoryUtil.memAllocFloat(normalCount * 3); + for(float normalValue : data.normals){ + NormalArrayBufferData.put(normalValue); } - } catch (NullPointerException ex){ - ex.printStackTrace(); + NormalArrayBufferData.flip(); + GL44.glBufferData(GL44.GL_ARRAY_BUFFER,NormalArrayBufferData,GL44.GL_STATIC_DRAW); + GL44.glVertexAttribPointer(1, 3, GL44.GL_FLOAT, false, 0, 0); + GL44.glEnableVertexAttribArray(1); + MemoryUtil.memFree(NormalArrayBufferData); } + } catch (NullPointerException ex){ + ex.printStackTrace(); } } @@ -613,33 +612,6 @@ public class Mesh { (byte)0x13 }; - static class Triangle { - int[] indices = new int[3]; //array of size 3 - - public Triangle(int index0, int index1, int index2){ - indices[0] = index0; - indices[1] = index1; - indices[2] = index2; - } - } - - static class GridCell { - Vector3f[] points = new Vector3f[8]; //array of size 8 - double[] val = new double[8]; //array of size 8 - public void setValues( - Vector3f p1, Vector3f p2, Vector3f p3, Vector3f p4, - Vector3f p5, Vector3f p6, Vector3f p7, Vector3f p8, - double val1, double val2, double val3, double val4, - double val5, double val6, double val7, double val8 - ){ - points[0] = p1; points[1] = p2; points[2] = p3; points[3] = p4; - points[4] = p5; points[5] = p6; points[6] = p7; points[7] = p8; - val[0] = val1; val[1] = val2; val[2] = val3; val[3] = val4; - val[4] = val5; val[5] = val6; val[6] = val7; val[7] = val8; - } - } - - protected static int polygonize( @@ -897,20 +869,6 @@ public class Mesh { return rVal; } - static class TerrainChunkData { - - List vertices; - List normals; - List elements; - - TerrainChunkData(List vertices, List normals, List elements){ - this.vertices = vertices; - this.normals = normals; - this.elements = elements; - } - - } - //TODO: more optimal key creation private static String getVertKeyFromPoints(float x, float y, float z){ return x + "_" + y + "_" + z; @@ -934,5 +892,7 @@ public class Mesh { rVal = rVal.mul(proportion0).add(new Vector3f(normal1).mul(proportion1)); return rVal; } + + static TerrainChunkData data = generateTerrainChunkData(new float[FluidSim.DIM * FluidSim.DIM * FluidSim.DIM]); } diff --git a/src/main/java/electrosphere/render/TerrainChunkData.java b/src/main/java/electrosphere/render/TerrainChunkData.java new file mode 100644 index 0000000..a099e62 --- /dev/null +++ b/src/main/java/electrosphere/render/TerrainChunkData.java @@ -0,0 +1,15 @@ +package electrosphere.render; + +import java.util.List; + +public class TerrainChunkData { + List vertices; + List normals; + List elements; + + TerrainChunkData(List vertices, List normals, List elements){ + this.vertices = vertices; + this.normals = normals; + this.elements = elements; + } +} diff --git a/src/main/java/electrosphere/render/Triangle.java b/src/main/java/electrosphere/render/Triangle.java new file mode 100644 index 0000000..bede795 --- /dev/null +++ b/src/main/java/electrosphere/render/Triangle.java @@ -0,0 +1,11 @@ +package electrosphere.render; + +public class Triangle { + int[] indices = new int[3]; //array of size 3 + + public Triangle(int index0, int index1, int index2){ + indices[0] = index0; + indices[1] = index1; + indices[2] = index2; + } +} diff --git a/src/main/resources/shader.fs b/src/main/resources/shader.fs index 401f626..905d839 100644 --- a/src/main/resources/shader.fs +++ b/src/main/resources/shader.fs @@ -1,7 +1,5 @@ #version 330 core -#define NR_POINT_LIGHTS 10 - out vec4 FragColor; @@ -26,19 +24,7 @@ void main(){ vec3 norm = normalize(Normal); vec3 viewDir = normalize(viewPos - FragPos); - // //grab light intensity - // float lightIntensity = calcLightIntensityTotal(norm); - - // //get color of base texture - // vec3 textureColor = vec3(1); - - - // //calculate final color - // vec3 finalColor = textureColor * lightIntensity; vec3 lightAmount = CalcDirLight(norm, viewDir); - // for(int i = 0; i < NR_POINT_LIGHTS; i++){ - // lightAmount += CalcPointLight(i, norm, FragPos, viewDir); - // } vec3 color = vec3(0.3,0.7,0.9) * lightAmount; @@ -54,7 +40,6 @@ vec3 CalcDirLight(vec3 normal, vec3 viewDir){ vec3 reflectDir = reflect(-lightDir, normal); float spec = pow(max(dot(viewDir, reflectDir), 0.0), 0.6); // combine results - // vec3 texColor = texture(material.diffuse, TexCoord).rgb; vec3 diffuse = dLDiffuse * diff; vec3 specular = spec * color; diff --git a/src/main/resources/shader.vs b/src/main/resources/shader.vs index 5222219..a2b344b 100644 --- a/src/main/resources/shader.vs +++ b/src/main/resources/shader.vs @@ -1,10 +1,6 @@ //Vertex Shader #version 330 core -//defines -#define TEXTURE_MAP_SCALE 3.0 - - //input buffers layout (location = 0) in vec3 aPos; layout (location = 1) in vec3 aNormal;