From 9bc3b5eb109f0e7efedba57017e0b0a39ccd3f93 Mon Sep 17 00:00:00 2001 From: "erwin.coumans" Date: Sun, 4 Mar 2012 22:12:12 +0000 Subject: [PATCH] Update btOpenCLUtils to allow caching of precompiled program binaries (save/load). See Bullet/Demos/SerializeDemo/AMD for an example use Fix in btBulletWorldImporter: load friction/restitution and patch radius of btCapsuleShape (it needs to embed the margin) Partly apply a modified patch to make the SerializeDemo_AMD work, but avoid breaking the MiniCL version. See Issue 594 --- Demos/OpenCLClothDemo/cl_cloth_demo.cpp | 2 +- Demos/SerializeDemo/AMD/CMakeLists.txt | 19 +- Demos/SerializeDemo/SerializeDemo.cpp | 60 +++- Demos/SerializeDemo/testFileCloth.bullet | Bin 0 -> 98740 bytes Demos/SharedOpenCL/btOpenCLUtils.cpp | 312 ++++++++++++------ Demos/SharedOpenCL/btOpenCLUtils.h | 9 +- .../btBulletWorldImporter.cpp | 36 +- .../btBulletWorldImporter.h | 4 + .../DX11/btSoftBodySolver_DX11.cpp | 6 +- ...ollisionsAndUpdateVelocitiesSIMDBatched.cl | 113 ++++--- .../OpenCL/btSoftBodySolver_OpenCL.cpp | 4 +- .../btSoftBodySolver_OpenCLSIMDAware.cpp | 7 +- src/MiniCL/MiniCL.cpp | 26 +- src/MiniCL/cl.h | 4 +- 14 files changed, 422 insertions(+), 180 deletions(-) create mode 100644 Demos/SerializeDemo/testFileCloth.bullet diff --git a/Demos/OpenCLClothDemo/cl_cloth_demo.cpp b/Demos/OpenCLClothDemo/cl_cloth_demo.cpp index 8a354e37b..5a48709b7 100644 --- a/Demos/OpenCLClothDemo/cl_cloth_demo.cpp +++ b/Demos/OpenCLClothDemo/cl_cloth_demo.cpp @@ -20,7 +20,7 @@ subject to the following restrictions: #ifndef USE_MINICL -//#define USE_SIMDAWARE_SOLVER +#define USE_SIMDAWARE_SOLVER #endif #if !defined (__APPLE__) diff --git a/Demos/SerializeDemo/AMD/CMakeLists.txt b/Demos/SerializeDemo/AMD/CMakeLists.txt index abe46ebac..88644d5d9 100644 --- a/Demos/SerializeDemo/AMD/CMakeLists.txt +++ b/Demos/SerializeDemo/AMD/CMakeLists.txt @@ -115,7 +115,24 @@ IF (NOT INTERNAL_CREATE_DISTRIBUTABLE_MSVC_PROJECTFILES AND NOT INTERNAL_UPDATE_ ADD_CUSTOM_COMMAND( TARGET AppSerializeDemo_AMD POST_BUILD - COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/Demos/SerializeDemo/testFile.bullet ${CMAKE_CURRENT_BINARY_DIR}/testFile.bullet + COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/Demos/SerializeDemo/testFileCloth.bullet ${CMAKE_CURRENT_BINARY_DIR}/testFile.bullet + COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/OpenCLC10/ApplyForces.cl ${CMAKE_CURRENT_BINARY_DIR}/OpenCLC10/ApplyForces.cl + COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/OpenCLC10/ComputeBounds.cl ${CMAKE_CURRENT_BINARY_DIR}/OpenCLC10/ComputeBounds.cl + COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/OpenCLC10/Integrate.cl ${CMAKE_CURRENT_BINARY_DIR}/OpenCLC10/Integrate.cl + COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/OpenCLC10/OutputToVertexArray.cl ${CMAKE_CURRENT_BINARY_DIR}/OpenCLC10/OutputToVertexArray.cl + COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/OpenCLC10/PrepareLinks.cl ${CMAKE_CURRENT_BINARY_DIR}/OpenCLC10/PrepareLinks.cl + COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/OpenCLC10/SolveCollisionsAndUpdateVelocities.cl ${CMAKE_CURRENT_BINARY_DIR}/OpenCLC10/SolveCollisionsAndUpdateVelocities.cl + COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/OpenCLC10/SolveCollisionsAndUpdateVelocitiesSIMDBatched.cl ${CMAKE_CURRENT_BINARY_DIR}/OpenCLC10/SolveCollisionsAndUpdateVelocitiesSIMDBatched.cl + COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/OpenCLC10/SolvePositions.cl ${CMAKE_CURRENT_BINARY_DIR}/OpenCLC10/SolvePositions.cl + COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/OpenCLC10/SolvePositionsSIMDBatched.cl ${CMAKE_CURRENT_BINARY_DIR}/OpenCLC10/SolvePositionsSIMDBatched.cl + COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/OpenCLC10/UpdateConstants.cl ${CMAKE_CURRENT_BINARY_DIR}/OpenCLC10/UpdateConstants.cl + COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/OpenCLC10/UpdateFixedVertexPositions.cl ${CMAKE_CURRENT_BINARY_DIR}/OpenCLC10/UpdateFixedVertexPositions.cl + COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/OpenCLC10/UpdateNodes.cl ${CMAKE_CURRENT_BINARY_DIR}/OpenCLC10/UpdateNodes.cl + COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/OpenCLC10/UpdateNormals.cl ${CMAKE_CURRENT_BINARY_DIR}/OpenCLC10/UpdateNormals.cl + COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/OpenCLC10/UpdatePositions.cl ${CMAKE_CURRENT_BINARY_DIR}/OpenCLC10/UpdatePositions.cl + COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/OpenCLC10/UpdatePositionsFromVelocities.cl ${CMAKE_CURRENT_BINARY_DIR}/OpenCLC10/UpdatePositionsFromVelocities.cl + COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/OpenCLC10/VSolveLinks.cl ${CMAKE_CURRENT_BINARY_DIR}/OpenCLC10/VSolveLinks.cl + ) ENDIF () diff --git a/Demos/SerializeDemo/SerializeDemo.cpp b/Demos/SerializeDemo/SerializeDemo.cpp index 7d131c676..704cf8b3b 100644 --- a/Demos/SerializeDemo/SerializeDemo.cpp +++ b/Demos/SerializeDemo/SerializeDemo.cpp @@ -13,7 +13,6 @@ subject to the following restrictions: 3. This notice may not be removed or altered from any source distribution. */ - #define TEST_SERIALIZATION 1 //#undef DESERIALIZE_SOFT_BODIES @@ -57,6 +56,8 @@ subject to the following restrictions: #ifdef USE_AMD_OPENCL #include #include + #include "../SharedOpenCL/btOpenCLUtils.h" + extern cl_context g_cxMainContext; extern cl_device_id g_cdDevice; extern cl_command_queue g_cqCommandQue; @@ -198,7 +199,39 @@ void SerializeDemo::clientMoveAndDisplay() swapBuffers(); } +#ifdef USE_AMD_OPENCL +///the CachingCLFuncs class will try to create/load precompiled binary programs, instead of the slow on-line compilation of programs +class CachingCLFuncs : public CLFunctions +{ + cl_device_id m_device; + + public: + + CachingCLFuncs (cl_command_queue cqCommandQue, cl_context cxMainContext, cl_device_id device) + :CLFunctions(cqCommandQue,cxMainContext), + m_device(device) + { + } + + virtual cl_kernel compileCLKernelFromString( const char* kernelSource, const char* kernelName, const char* additionalMacros, const char* srcFileNameForCaching) + { + + cl_int pErrNum; + cl_program prog; + + prog = btOpenCLUtils::compileCLProgramFromFile( m_cxMainContext,m_device, &pErrNum,additionalMacros ,srcFileNameForCaching); + if (!prog) + { + printf("Using embedded kernel source instead:\n"); + prog = btOpenCLUtils::compileCLProgramFromString( m_cxMainContext,m_device, kernelSource, &pErrNum,additionalMacros); + } + + return btOpenCLUtils::compileCLKernelFromString( m_cxMainContext,m_device, kernelSource, kernelName, &pErrNum, prog,additionalMacros); + } + +}; +#endif void SerializeDemo::displayCallback(void) { @@ -277,18 +310,13 @@ void SerializeDemo::setupEmptyDynamicsWorld() { case kSolverAccelerationOpenCL_GPU: { - fSoftBodySolver - = new btOpenCLSoftBodySolverSIMDAware( g_cqCommandQue, - g_cxMainContext ); - // fSoftBodySolver = new btOpenCLSoftBodySolver( g_cqCommandQue, g_cxMainContext); - - /*if (!fSoftBodySolver->checkInitialized()) - { - btAssert(0); - delete fSoftBodySolver; - fSoftBodySolver = NULL; - } - */ + btOpenCLSoftBodySolverSIMDAware* softSolv= new btOpenCLSoftBodySolverSIMDAware( g_cqCommandQue, g_cxMainContext ); + //btOpenCLSoftBodySolver* softSolv= new btOpenCLSoftBodySolver( g_cqCommandQue, g_cxMainContext); + fSoftBodySolver = softSolv; + + CLFunctions* funcs = new CachingCLFuncs(g_cqCommandQue, g_cxMainContext,g_cdDevice); + softSolv->setCLFunctions(funcs); + break; } @@ -516,14 +544,14 @@ public: psb->m_cfg.diterations=softBodyData->m_config.m_driftIterations; psb->m_cfg.citerations=softBodyData->m_config.m_clusterIterations; psb->m_cfg.viterations=softBodyData->m_config.m_velocityIterations; - + //psb->setTotalMass(0.1); psb->m_cfg.aeromodel = (btSoftBody::eAeroModel::_)softBodyData->m_config.m_aeroModel; psb->m_cfg.kLF = softBodyData->m_config.m_lift; psb->m_cfg.kDG = softBodyData->m_config.m_drag; psb->m_cfg.kMT = softBodyData->m_config.m_poseMatch; psb->m_cfg.collisions = softBodyData->m_config.m_collisionFlags; - psb->m_cfg.kDF = softBodyData->m_config.m_dynamicFriction; + psb->m_cfg.kDF = 1.f;//softBodyData->m_config.m_dynamicFriction; psb->m_cfg.kDP = softBodyData->m_config.m_damping; psb->m_cfg.kPR = softBodyData->m_config.m_pressure; psb->m_cfg.kVC = softBodyData->m_config.m_volume; @@ -715,7 +743,7 @@ void SerializeDemo::initPhysics() setTexturing(true); setShadows(true); - setCameraDistance(btScalar(SCALING*50.)); + setCameraDistance(btScalar(SCALING*30.)); setupEmptyDynamicsWorld(); diff --git a/Demos/SerializeDemo/testFileCloth.bullet b/Demos/SerializeDemo/testFileCloth.bullet new file mode 100644 index 0000000000000000000000000000000000000000..7dd84b68afac15b3ad5218cbd49e11f89f214199 GIT binary patch literal 98740 zcmd?ScVJZ2_Qf9%#fA-36qRDp`X4uR=uu-P zoie@i5uHZ&A29y5Mp;?gWMyUbmLXr7kH+cG_ETkct%v48Tn6EdeV{2_I0XLXKzft|h%M{FfJIiM<9*^!nY|KvS zuN`h{*jn&L_a8nWah+Xf7kSKj$o;6aqUQGzRf!o{l{NBEJ}PT=eQ`;G&phIz=#T6g z^F`=Cl_?kc>#tcy4Lz@GU-uY|DRc9TM)|uHHQ=+FCI4}a^%oB|8m}#}8XNZ(Iwz&NUlS z9P>5@S&dZ-Hm2%(%V_b?w#6 zyW1Fz-Va%ggBG{_$2C4!GiD9P z`uccgRNHJcMqd={Rdz!)j#AghNqa9a8i&re8aJ(4&|nRHeH?RVwbAHwuho!q_aE0d zvF8S(vFa_WA=ltPt}*z+<_V6+vZnSr|GaC9M17BkzRvYtr51EB8YP_rPi(M;KIZSt z9&9uYJK1W8x7S;vuIr;`VV==ga*oxQ-7Hyejk@-#Ag97;EV|8V$ebFiq4#Rp4Yr?8 z-XUpxS2Suk=JZ@?z2>8#{k-G-wx6H%d0q_#T|fUPHGW@h`}q^At%mF84bjlOcAFP$ zKR@bntKs^2eKhnj*Pb}qD>O?zXE^ zc*$;CTMgIG8=^5pd3#wu+t084D%YG>uAkROLmv;lSA&jfYu0hb3ajDz`9G=gYSXc% zpFgt1YIIEZ^ZID$3KVNx?)##V*=k?Lhd!?_Bs#9z~f8*LAB5RHNAcr?Dr_VbR#R>Sr4hG^(@ynU(d z=bv0+HC#V$h{iCrj_c>!e%|*^t0CulgZD~bAN`)O{rq38h8Ry24b#u*hthh@M_v6~ zU+3D}7r$xy`M-Z2Tth+jh`rPfdMT;8x;L>ud;re+)H1zS^TF>~4cE`>qoMaoYy8w?qq%4J@jI*G`gwgc^j>L=b!%-u|LYKY4Z40_ zAC0=!@!^}cpI48v8m^x=L_;6*7uVQ+K5U-VklLmXHbg@|H+X1??dK2OWHnqrZ-|Dz z&JVfS_Vc5wt%mGvgV&+=D!JbF^NaTHXgpCgOs}O^N$WKq_3h{SyxRD??dN-KVQad6 zULTFR&Z{T49%B07Egy8KVbUB2?im`Qp?mG(ojV$hj~}uc)UL0Fz6Nz4e7vRY=Q}O1 z8m^z$XB}EYpSuT-xBdK%F;>I%^M+{XbzHyD_Vb_iu^O(Q*GEGibA3EcxX1SMlee%M zuAetVL!Z05JYxI#=(i6w=auW{4bjlY<3FF;em>)ItKnX6Z-|Co$J0;R=jVIeZ#CTW z^M+{XRV$JJQXy0|mf;PRSJLN(@pZl9k|ZXH~M_1dfD%X&O))^S8j`J8zj z7i`(WyuLfFlhttjJnlLUTe;{VG=96O?IU3tO+UB&{QlWiLu#8I6<4GEPn#Y@W9Vf^ z1nbZ}YWzF4pU=C@YPfzLSA)5uVL#iw($~i)FWP=ywc2X9e%=rbeVxC$#`g1*p0gS< zr+V#`oqz3h=)Jmpn@ZE8KIq({Ce1fJiauf=*Zuu--9Pke8Ee}ljmF?zZB5tD>!YEM zht~M?K-v@Y(GDylhttjJg!FG_G1obuRhr7t6;D6@wnkb+s`lC!D_gE9#><>x#!3pRi>`( zzFV*keO~SKz3u0JeZGr1uUtQmt1+U<)=SY?++|(p@u*sMb&bW0S zb6yp{azm&_@maQ?-@M#vxaa3_H4gmcVsl=t>)HGWcV6k^amN7L&o6#3*elbc;%cm0 zSJ{=l>U8%0frh?kIQuNy&)<8*YPfzLSL28wzZs45Pi-4$=wtqJ=T+H%$ZKR1zpxsv zpU2fGYI*T;*74B)^^(u_8q}}7Zq?I1KR>dWJ$K#n^SBznJTp$Ni^|k1D@Uww8hRZk zZ`UI6rL3d>XshA+d0dSzQ;*6XRi*~-c)w_6UdQa;Z9ji`memkXG(9S=#;x0Tk#n{( zwPn|>dN>XJ+@Qr5wx9QT+-kUf9#>^8}!}jx2*I5nW*K4m%yYfHF(75EXtGYT3y;p~}wEcXynH$sh zFZTI49#0M*QP0mO%($mJ-hS?5_Xb{~uk#mr+kXDvgstiNd0dSH4sXAdb*$g@U%~Ov z8qW-|{ruhSt%mF8aW%epV)q`bBd^Qo(DQ2JsE#$S%Jp&6stsmuTtAPi@zx$kb^lW| zp8Q|yp6n~VRv(Yl*cLS{Mph+`pJp{UpXldtHFoUuNDnlAKQ$GsL-*RA&)a@J>9-%u zI;dSAjkfutmj5mqIfoYq>(H;gHhOwh>KWOqZLY8yuAj$U$F|2l@W>|7`17=^Ktp@` zMPDvR{U{nAbhR3;pU2gB|C;efpwaluXM?@c=T)kEyVOB)JU;*Md$U)rpU2g>cj&OC zXdHgq=YfVk9^>baU-PG|W6u>#>*sMbHg1}+0>4_m z_}5UsI_c#N+1C!QN*q2VxCZOBSJ!vjzvusC9V5#&gs$V{#w`-x%Q~LTvR>l)dE9l( z?)-3%b)xZCkE~wKuk?BK;?ErtYvp+Sd2ir}rborqxOvYB#;?X7xB1-F$Gqa_!HM;< zj;ZHc4cE`(YV5PF(F(l1HlJLLXq@gAMs@z8tK|ixqKaZ<%;C&bL{7N(? z_q;XG&^_w+f)+LV$aVhNdDUkBT|bYj(d45`dfXuz|N3-Ha1HA7s?COtsi9-55*v1~ z8k|q`^SByC4-f2hsA$ZZcTS+8*RfxV)u~3Ks}kouywa?L+V#q0g&l&u*9cMAk9(uN7t; zuAj%%$bb0PBl`SNk=jt$G0@P*BX7eAYu*)&yyvWjdww2QW8&6bd+a9~_db1CprP0C z!l`5WPCBkCG5ijz;hvwz)tLX}={>d*jSiC!4>a`kF?d~@?CVA2+w+3GGCeA;M(YQ6 zSjjqCS8iU1_N$kMFUUShuE9MowHmIU$JIFExkFbr`?Dg|e!Et|I`r{4t<9?J7sX5N z-C#AOcD>ed(T+c@Y%ABtob#>^t`B`Y+BDjj-CnNqvktT8u3Url(#UP}OV8az}RiY*U#hb)sXgg_xMCKj{m=J1HaPixNDn^iPpoa5}mKI8t(ad zT#X&dUt0P2ri#?}QyvS}p})s6yU~zDmR#p!PYAr-+%v@0xaq0KSH2|c=rr zFX{GMUgB5j=kIp38m^zm)j0dVSDSl=-FGSp`nlG4_D5K*m=8t9#^C4kwLw75RJWV z|95a+>FZE{XDM5=&LU5_53dtsWEA|Sq<0E<7!-Y*Y>@a9#fHe z@YyQ^4ZV)v2dz$Bbw^d=lC|^9I=C+A=W#XO7}~V=;-M9(wI!Q1w3j?Ec42DI^JK4P zS`BK~M`QNVcGWY!oRgaL$*N!-`rN(o*LJC`uc=CO-NkCSejaxnzpeYW=ZOO=QfK{j zb*RRen@(EONY=6R^SNfP-1GCe8XtW+ulsH-D^kzxdta!=igu&>_PnktF=DaRaL>=< zYE*spZqMgWs7Q4md`Y08uaA{$+GOu|vHZT;ZB|3}%JitX8cptaedYSYDpG}sTLKMz z&(LPfg6yj=sY=GT|ZQBTv(CX?|)iG8xEd$Dy}suQFIJ@PZrvu((CgT5r;dqhW>qB~7#et@=}~btu3mFt zuT$k7u33k@15eb?4f_8&I5Au9ch}?u-d-<_EkAjpdam5Vef7{w)nT4E{inRdMCnm) z%J=Z#xPBgY9j~48#0u_xrVd&k-23S7vApt??dL=8vKp?R$JO}psNa|KI>WG^+Vd;t z8K2|KO84`&@3%4Q5RDo<5#J(*kEnir$EH`Lk5{IinpW`u+)TZ;?+MA8PvqXGyqT@( z`gvT9e}DGa%6DZQtrk5I?3F&RMxHjL<`TJYKclDBaQ!^4#vvap>Rl|ow)q|T!7rOyp7&q`go=kbAt z_V%BzU7gyebV;J-seQ~kTtAPivHOYhdtWzgR_fDllYxdl9?z~^n3^l+>MPe-4X%s& zoV!cM?$P^`AE&22eSJZoq0g%auH85FL++Bqtq1OH*1`PiqjB`|R;$ilG$S>l;+#;8 zM;1(2)A&D25{V{O!#zKbyH_KxJ*QXO@#U%a_q{Y!qx^NNv*x<7ynY)u`&NHkPGUymon@p}l?I zgBE1B*?mdktoByJ_4BwIe}0)-^~a{^sR28NAM^X?tja$2oyCdIXSQ(}q3l83xJ@W0viEr20WA6HS+;u#;@PJicEuWbxKD;E*(C2PX>E~PL zFHT%?w!O|>KaZ=CyY;NqEoaP5wJI1AXy|pkEYHu+KA2r$ha?^rjZ6O1&KwW-dV5@rJAZs<)mtxDq~2`$NpR1g?|mNpJuh*s z^r%O_w%+dgd0dSfTRqY1b$OlPnOzzM&(HPo81a+s=f50e*WvnkT#bR>-({YkPi~eU zJU`ca_0ECz`T3u7Mz5hCyMB(R*vHZR{Js~zG_SWG{rQ5>bu^lgtT|Qgea@{o&eU}M zoaZ5NG)_A4>gqw#&u^dEBKZD?-mB~rht%BOu_|%x$H$wRuAj$U$Kt)4^||WJiqxtJ z2M5Po_o#<*JJvk$r98*kbb`@v{XDM5`y1N!U48zUsqOO*4bENdiBGlOwPs++gNX&N zO)wg+pU2fWdquOpr4!CbO)YNHm%Yu@&(AsOr_{jhmnOO{&M_LUpT|8O-wwX3&-zEt zOwIXy+dxCFZB?Jhhc)oEGJKttdAY|L4dI&Z?V z#KCvv84cIZwVhxKJ=s6sjTac3Dx+|Iivd)Em@YR zxh!ckTtAPyjx$?bxN`jyGgH&OYf>4 zckJxc180s6G_+sMzGOl6d*?4pJam^`hwJBYH6H)$k(GD+JR>#lhB<+T_V#wUtFtE^ z@^GSP`_s)j-1GCe8ejdqZq@W1W~H|8J~hzL=T*}$e#~z4)^}Y$kE?M|Dz$pi z`*Tva5Byi)?Rp*SzRgQqCOv9a()zCJ=W#WvFTZx>4br1l6dxPBzN_zj-dkt;d7roK zI$S@Gt1)zaH}jhQtIsS8e&$YJAN%ZT`}r@`_PlWYobRv2>9uD+TP{7eGWGP6Tgvz9 zZU4~TUfw-f^Nj4(jpN3fny#PYDRDF|UNv+D-`hUC(e1(aKeU%j9XOPpt)Qa8|0% zHYcxkd!^6a!+&p4v)}a(Ca&J&Sfk87`Wj{Nl`+3}TJXW|{ z@0UBwOm*KdJJ8VgK4Tv{w(sk^Je+v{l|e>>+T1h5T}RL7ZteBNp|ev@?L0ov(AUTO z9olBsY?6D1&yO}5uAj%%IC)%6uW!cBOl^7c%s@kb|Ga16g6xl4Kb&}Bh}|pK&*N&m zc4epPbsJ}=CMAai8v2;8K52FK8DA|+RCl);uAj%%IQfos)uVTvmD*)gL8!*sOMcA0 zedE%^4-*HPy>k6LuExf8kFVZq_L-@zPTV$Bqsv<@67L@KP~yOY>^fXOkE`*{)w`{k z_x|kEpD%9{s`2#t4vCI0FHT&3sJ+fzKaZ=?ZbkPsvrjuC_2Y!8fnVukp8M_K#Kpz( zJgTYHaQ!^4#%C9tyZW{X6{(H+_pWk&rLT`RZ{;P5kF82{I%$+S9+5`5>E~NWAH4o;>+Pw!BZ%KhH>g+p1M?4eEP_ zZ(iJ(dU3*%#9lpG84cIZ<7%9=%bmSHs+yg;>*C>ohTf~j?N_IIS1d`q`C?0>;re-8 zjXgftvAXBt8L2^qXNPKB+ihX0-?mE*sOTG3l?ndwsdM zH1*Kh>jDjZ%ukcQCv(stOA^%`tp>I0vyM4cyY_zT`SR2^TVEV#^i{uqe%j*;vcH#e zcU%j5JX}AIyN>Gy9=>YzyQQi5k6s&S=;QI-qpPwdid45ndj#K0)z`-@ujC~TlwP~b zA=cYnKaZ=??+NQ;e#`Ur`S}Cut%mF8aW&rA@dWez^R4%OJb3P` zuaA9hvi*sMb zuDSFlbI~FY#ZgY)>>*sMbW_)#8@6XFCQb%8XcA%l} zcc(O8mD>K6s>BmtoNY8*KaZ>N>euJ5y0NwVeByn#2O9c#JW{bB)m47pr)*)R(Qy4d zuEyfG->$CtcSY)pt)>MU`n;Mpy=`j0tE&>v)KnM^x%NyykE=1a^~zq4msF&hd^|tU z(8uH8tH-T*TYi>vzw$XogV*2agK;%h-ZQs)_iHLrudT`rH1s--Icr?sF|%c_PO=(2 zC#a7`&!-RQJ$Z6P>XN^c?6)qgO02)sYT#4#(fB%dmzC}2RHTN!Um0j< zPyBK8g6yZK$#a}6dpy{q`e(FSVoq6R_^LyLBTpj$rnqJ4_C)(%d?_6nXx_%y4Bd7D( zN1(Ao^X{Pcq( zhU@2XHP&Qz>)A*&a=w2sxKGp?6_+%xd0Kww=Z+>T%{p8^kE`+0XIu8z+S*<(_iy^ZnVs2J6uGC0RAqspfJ#Zdmn* z(U97v560Cvchn`-pUKap9{%UrKtta%>{m2Db%z{}U)DTqG+aNAtFhhgU7OS$ay)+9w9II@ejZn&<^DhQTp-8e^dTn%>(Ixe^A*Rg883VF-yEyq z`gvT97mBx8xluI!cjcAA_oMVWPG~f)Z)5qLpOP}G;re-8joxjK>2N#v6B4|18IR)tN1VW3KOahg^}DxLU5kPOn=H*U#f>JpMw~O8y?e4aEnA z9`k97Y(M{YwCx|RpU2fW@b^CQ``VSM@{1SA-<7f6uKW4QcDA3d`gx-{9G3YH8~m0}bu%i<1koj}VPc_gf9u z&*N%Lx%Yw|o8)-R?eT1&q5W#!x7FFriZ|Uc8HSWr5z48yy*yr&FLN#8h z-jIDh$9(>`?i#F@#>5woUAa~?uKZ7<4^BKGue~0! zyY)oZ&*N&mH}%GqpNhuS$KDru4Sq2@FTrcbvsPFQ*U#f>EO>9B@x&8zR|fq^1hm5{ub++Q7sZ2GvS-(=NxrF^SBxdTg>l<@9uSa2l;!#neQ2n-$DMyytnKiV<=L(V7DYvXEkZ?}{2-Kig+FMF2x zzGUuz>ePQ_9XEfyvss7h=W#U}eS3@fy9U30x+qwOzGs+mz=8to8D}uFBMkH%lLJ8v1x#v&XUK@4ddi zekZe6uAj%%NDgQ!uUl27cA525uvgkkx_)o}-s{6}TMgIG<7!OGzuX-2<_4BwIlMlUL z&RO}}ZP#rIH1u9oKUjd;n4qk?nS^xC)@r(7|_+>5rj_xjL&zVq}A+5FwXZB|+h zJe_-nxEj|FpJe*S?MFQvT!VTY&)wKOv6X0?+tgm?)UJ=lxA#q5&b`mR2ma8*9doTw zTG=7N-^_e^fz@#RJnlMHoN$cPu1sCuDlgE`$Nbxq1}A&yw0GH`IS!$GWQIh-)c2nKaZ>N_{eFpN0q79zL+QeYyC?9 zyiemV4>cNN`&bRBU9Y`5b;4UR|H{+@`!12!8!~G=zhM`n(dffN6L4HVkGqc75B#DF z8rMJYNjLH{9gj^H*nWOtuGMh;Jg&y_mz$XTtLYuS30=qZ=_^gI?flUp3FhJYd0dT( zd)q$9c~$d&y9dWy|Lovag|?rcQ(-k+KaZ<%aY3i0?A7W64w27xuk?9!S;anP9dmnF z4cE`(YS_Bb*T?sdjWYc_x3Sf5{k%RJtOG4RIOcjCou*DO@2~dzm;Xw@as9j@8v6R! z@ps$Lhkau;TtAPiF<_@X%ys_K6L$yauKwAmmQ(HDZ!iDVYPfzLS7Y*#KbyZp-?;d) z;ND03)$Q}``=5^)5S+WF*T&WOV6VpJ-e>PlRe_i239)Elmc%t5`(wCC8^KCz$Cx3^Rny#P!lN$d!+xGJ_uCW@fpEpE9KOY=< zvF+!rdsq!Q)~25~L_Sr4hG^(@ zd^*JT^FQ7`+N{I%^M+{X?@f#wGtaE!f-9|t>*w{+(D$PH+#Pzm?dRVeVKrPoua8Du z$D{Xk!%RPawzbu8{k$O>dL8Ba9&h^jPHzr0$HVpW`e^9A($~kS>+SooeE5mgaQ(bK z8hWpEkGkc-`PuZSjqh6x*UuZGq4(;&yKO(;w@YyDntom%4ZT-d{3`rif_jXRcD4Qe-3 zL+{o1ALbg3Q#V-+*UuZW4z1Dl)j39E>;ATX$nk6NI`s81td;HOKXgu->s&NUk77Mp z>5saepVxI>?cK`u^Twa&nVPPjH$+3fW_jA-wx7R!zSWT9Q15ldxE0Gkx%y(N@Ft^M+{Xe%@x}3Zrq|>l4j7 zTtELOH8ve;`}x!=tKs^2eKhnj*XPx(JML}P@xc*R!}as}Xw-FmT=U2f)6W~Xu^O(Q zH$+3PqxapXn0`KRO^!JpuAetVLth`u53v3G_>ES>_49^kXm2lRW}ly5ves(2e%=rb zeZTu}`ME3Z8EzV8kB6MQrk^)NLtleSleV8voo+Q;Kd+C5K6mvxuC3X{^p6*|vDdll z=k?L3>)ie6E!)q(oMbgzKd+C5UWY#B=iP1l`M>Y88m^z$N29KFR9D%4e$)o5L2d3C z>Z768p^tf^n{7Wo>k!*NWKQ+g(Cg59b?#GV)o||S?_6oFbK#plh*nnmqrUxI-!nY( z>O`aQ?VA;*rt9bbq{f@;jy4*_XIc%{&+DV1kB5FevDIIP7>!ffS`F9F>!TsuTCa<0 zjV{e>Ki}f1IcBe1KW~VJJ|0aTvHkp`p;p87^M+{Xbu4>*x#@%dx8W?a4%g2cqM=`J z|N76lMq}ljR>Sr4hG^)$8q$3)qp_gEYPf#>Pijo)J=FB`=NelL*U$e+jeajoH2wUt zpU*VM!}arqXy{}9*%S8l_KQEX8m^x=L_@!xxN~3I&$sJg_saG2hG^(x{#45L^JTkQ z4cE`>qoKV-YaDcg?dSVlWHnqruaAasQ(5VcOy5_le8=|l4r{H3>*xQZMwe%8Kkqoo zdWq}j4bjlYeEc=GpYMB_)o}g1AsYI444-TJ`M~$AhU@1I(a^p2&9iJj|7IWCKd8+; zLqjzDz3Mlr-}uYNWo2z8=Y)JSCM%~&R+jv4J{rltM_m4PUmiHDg~H%~{aR%IW(Yj6 z#xnd|uw64?y9ws!lDCD7PM$6pJii_su%A3#F!VKV_Y7D|!F*lvwp2VYcz!)NIv0%n z^XCkzG+sp8C!CK4ExorgVbIEHaBjKG_U~Odhxp2UK_s%P@ z9c1{rU~OgS+;)Qby5#LB!`B6a=jXx!^L4>?lHu!u?JGm)?kAY9OWw{h{9G`2el8rH z3$}|4f6ie0%h0(82B9^RsT1MHv-yn_Yv&o%P2Z-Bw` z=K@FPg5ewfI>6e?(7FE-%wGq2+Bd-9`MGew{5gYZ-vB#AhR!`y<&uX-_;Ut>=jXx! z^XCkveFLn644vCiFn`YEY2N^Y=jX!FxnSBiz&d5fJxnlPmpttoVDNlhI64;$-|+Ve ztg{TAE8i5awO8b|k>Tfp!Si$B=v**-!@st`4ws>Gy9(xCTjb#pzAhL%KNk*|e|*94 z4PO_mn+%=XT`*slJUqhJ1%v13!U6Mj!SD?~7wiaz=fz+>Wcaz{;Sv6v!QlD1aKQXI zgW)BZGpk_>%jr@eHRSh(7eGJupxr^b0!au@O8o9`Ssv{`MSn$GVqScfDIMQ z*Trw}2wwvXo{z!NxnTH)KWDIEGIZ{6!TdRshnM&~FnB%&2h5){*a*eLTks9vo4`h9 zz>XEnp9^_-gr5rr&&S~CTrlk$V52hRjuy;6$H_}59vD1d7Y>*|XE1z2^TuSrjuXt! zB~SYX7(Blo958>*VEBgSjm?0K6U^5o50CJ3!QlBA9Gwe>Z)o1}8L;t!`MTuc5q>Th zJRgIjbHTK4fSr&b_e8<`y&_Nh1{geF7mm&a!#DhE3+yBrI(LF#evcJ}rusp?s1LmKnVA?mpk}`Y^ut_rfT=MV;-#fwJ`MGew z{B?k7-vFB|L+4Hr%-1DP`vw?1KNpV91;aP|IfLcP(7C4x=Fgct?Hge5{9HI-{+z+I zZ-7mep>qpVE_vEFz~K40aC9!1_6@LU8FC8+^XE(+9^vm57(8DW4w%1JVEBfg3sxjU z=N1d*=aPp<_}3N~JUn15}7;Tyg#*y%EKZi!&NE_rx_uL}mx&xHf#>w=vs!`B5X zm7#OX1oL%ImB)C5p9==h&xNCN!SD@#&S2#-bnbM){5g|{m-xD1@cdjjVE&xJW+)!s zqJ0BwrVL*LY?cha9(md~z~K40aC9yhzM*-uGhk;3=HJ(mr+otqo?j0Rm_KJQ?Hgcc z%FwxIsa*2#2%iTA&&S|^`Ev%-z5zBzhR&@J%%3xP+Bd-9`MGd(E|~TWuu2&^_iVv@ zUGlVVfWh-~;pkj2d_(ij$$-rj%-<{W@CZK_44#j{(YavwhJS5=&6A;X=L_b0CwX{; z&jW+!V{pLyIfLOFzAo4T89H~NV7@MSc!bXbgXd##z>|PZT=KMU zfWh0}P&z z!2$E<3|1*a^DfJPT`riPOCBEK=YqlWF*rIG4BybaD>7hL3g+vQr+otqo?j0Rn7>zG z+Bd-dBSYt2rEw;}BUza@X8({GKTsS%x4BznQ40gQ?oqK~|{+!9vUIGTs&xHf#&l&7S#lu^) zZ-CvDfp@cD{<%h;_6;z2{#@YbTrhmY_e8K;Wa!*m1@qTIp7sqecz!M%Fn`Wq+Bd*% zlc95OSGnZj5&oRP;Q6_5!2CIbY2N_5Lx#@1Q!szd=fcssVA?mpR%ggvBbYyD^6&_M zufX8>x^Te!y#m8G{9Le<44qpen4e1?9^qeGU^O!QTsUCX3mlybhHv=m0DDe`&V61m ze;wp$-vEQ>=fVN==M1KO1MCGEI`>7DOCBEK&lwD!p9=@fpEH>D4X~GF=-ihD^XE*S z_6;z2el8rH3#NSo?3E0;uL|btlBazG44$tGN9Tg!8~$E_y(UBFzAl)*SLD^m@N>c7 z`MGd(E*QSyUt3^r$k4fO3g%y1*|XE1!j*9Ch^hR%IkFkhEEJi^xn zgXibM0rPdi@C`o~>>U|8_g%sKT=MV;f6id={9HI-{+z+^5?>c=tqh&}o?yN%d0;*d zztO${_Pz{X0}hy94@~<8*atFn?uROuJUqhZfx+`JIAH$kTwvNaz}CsoxgQDU&zU^! z8({GKTsUBUE*QR{c^_xM{!cJ}&g5y|0E6e(g9GO4f@$9X`$UG${Z!?Whe!1Ae$EIC zo}UXx=Ynb90Q)RM?&pH}b0!au@HN2T`4}89f3Lvs4bA%^1NNoLB@d7AuPrcmemyu~ zzVCwJ8=Ci32JCCW{5g|{NBFv6@cepkz*|XE1!j*9H4UhR*#}FkhEEJi^ZfgXibM(YavwhMx=en+%=%yI{U9d3c1M3kJ{6 zg`;!9v~Pg@H$(3K1oQWbJnb7`@O)i3Iu{Jz@UJbfKV<0KKLzuzE%LN)fWh-~;eh#b z2E#XeU9e3ubnah*`MTt3-vEQ>=fVN=b-}c6fbmOLI=7KvelB@5>QLYv4EdhJUVsH4&aaKR96iy%`w3;pc*FnIU&8l}n!X4KR4VE*vm_ z4~(}6?(2dz&5*mb(#3D^2tOALp05i>=Ynb90NW-*?zV#Yb0!au@O8o9`MPkx{Cx() zH~d_%?K0$UuX4%5Bm7)2c)l(ioePF<__<)sGUV={bjj1c0S3?4g#+gA6`1x7upKkx z?xb?b!y|kRFnB%&2kZ~8&wy#)0NXi3?k-A~JUqhJ0E6daaKL>%!5wVEBgS?V17GO)!7Xxvd5Bb;-jc{9G`2el8rH3x;p_xnON%=-jq~`MTuc z5q>ThJUc7`4}9X3#NSoY`+Y-`wQk@ALQW?z6KaP zAA* zV0ekThJUV z`FjPXeFLn!44r#~$|Vnv@OfbHd<+hl-{ZlwZ-DiXp>ul*=FgctJi_OJ!SgXVV7@Mx z_6@LJGIVZl!F*lvv~PgH^K;?oTrhk?^ZI1K`U>XHnLOnFqSm*n9Qel8e1KNpV91;aNquYU&YNWpwRHF<*hx?u49dUz0+KW8v} z!_Ng9AVcRKC77Q}9v*VEBfw3pP-O&OKT%Uza>Q!p{YR=jX!FxnTH) zp9?lfhRz)X2j-z)O)2tOALo{z!NxnSBiz=meX z9VVE6eUOJo_!?mFd<+hlKW8xQ8(_m_=-d&4`MTuc5k3zLo{zx+^L4@S4b2;w0XtUZ zlBazG44z*P4wye@FuX+bMrFW83+B(6JTN~uA&>D0|6Btbqj+$@{PPqH-|%z6j+5bQ zfQ^;m=aQ#=0}P&@3kS?!2N=HL>w=Awp>vNH%-1CkkMMKB;Q6_5bS{|o4Y2VUa!(M< zpEG%Qgs%$*&)0f=!U2b59n` z*CkK;1{gd)7Y>-eS76#Vz)q2&b5B*deM=fVN=uPrd`8(=vybnZmK{5g|{ zNBFv6@cdjjV7@Mx_6@LH89FylFkhEE?Hge5{9HIX7YyI<=M0vVp>rn*=Fgct?ImFF z{9HI-{+z)kD<0m0Z}{gL*c2K5{J`>M`1juA;Sqi=7(720j?M+sz5#YxhTN%w`FlVf z9^vbP!Si+DfcbLt0c%-1CkkMMKB;Q6_5bS{|o4X~07xut^ndqtl14KR4V zE*zZ;hHq$ISq7|JF#ov*dD=I?;Q967fcbL(#-NBHL&*qJhXUxNeY z_dGCs!_NgfONP#!Bbc8{9vc7`MGd( zE*QR{d1q(9&JoO?GkJJ~uL}mxuLlRrpEDS~;pc+Qm7#Oz3FhaLhe!CiVDS80I64;$ z-|%z6=F8Bz3k36Z$-^Uj9vD0yg9GO86&Svuc?&aOi&QRo+Bd-9`Ssv{`PT;+zM*;N zX28x9%%3xP+Bd-9`Ssv{`MO~EhUT500lPpjUza>Q!p{YR=VNemE|~TWunRNfUL=@5 zXY%k8Ujq!DkHG=+=L~kS;^8g$hJUW{m}}q1c&x8;9&@ewJRWmh`WTP(^~PhaBcI1( zt_2_CG3ULH@tAYh$9SyIQy%N{kH?%>z8CXYpD#S-Jn(ru*2kB}`nd6!z4vo@?C(8) zr*~1({;k-q@;%zu8uQIt-Y0YO`*Jps$1MzMzO9?(-WwnjN6MSOs~N^2j>K*a~;+HXk^|sT0N_PMt05;1J)#5Zts*$Yx0bP29US6Ob%;}En}|Aai0>W2A-=7Q zed^#4_v@thxl&+roQP9rI~i^d^|63MymbVJ`1Uf|)WIR{*U`s|ICYxI7=2#XzVhjB z*I*mLh_{vDIC{>+caYIeM%cRG5T}mg==l(*&W?gx%W#SO0 zP8f$cb#|_ULwtXaqvuR~7a0f02wN8%;?!{*J!j(7X)Ysd&Txny=yCLX@Kt>7U>RZ8 zDIDSlMR15yC(2jh5N|KTY3X$l$LIbfBWzu8h*QUL^ty;sC(2jh5I;nQn~$zT9G^S1 z4i0g@j$RjW>O}b}9O4~hgsqD>KG#u3*qq@Ir%o7$ICY|Y6%O%EGQ#Q*$L9{KgF`&3 z4sq(>d0|>`h&S z;?%+OHos3s?fX?Y#Iqwf#PPX=jIjB@Ax@nz4sq(>c~KnV{USKT@wxso!qx?cICa7} z#HoYlh0O;J@gpNR#PPWSb#RDB)gewDJTGiMa7;XcLmZzQSO>?%|IVotD65GK;?%(x z!}bsk@uMR+#PPX7b#RFLb@cU4oH}@3SigWne6YvS{zDv}8zLiYU2up~$8q$0h*Jm8 z3!4ud;>UO#Js;xu+)x=#D-(w}b;3Brse|W5aflD|IC@>g@wwqL!qx?cICUIH&zU%N z@XfIKz#%@u4)Ku@9OBf$^TO)DAwJ6E=yegt=SIs2yH4Q{r;g+3 zITNQ&LPpqp;1C}Z!6A;%9ajg3xL-%Fi#T=gJU8b|9O7d=j$RjWd~TeKuyw&9PMt6g zaq8fCQ5@pOM{tNAC1ZRY9O6-Rh*Jm83!4ud;wN|295x>~_~xk*9OC#~P8}TLr+OScXX4bs^W2;>afna!IC@>g@ylEpVe5iJoH~x9 zwelo}Uq<;V9O6kCVdp$?d~Q-59O8Z*y)NR^iSkuA#3#!LTNiPBZiO}b}9OC&hTphhG;`rQYGTfXqafnmLarE&bPMs)Ug+qL*3^yNLhd4f0PzQ&2R2|~f ziSkuA#HYy!n-6h(uCNXcalcM_4=V-EY9fO;b)tL~4)G!xZlCqBAdb%!*TEs~*U|e- zoH|jy3WxaVGMp#q`4GqFN@Rr184hvkIF6nVaq2|*DjedaGQ#FV{8SlbGMtv4GaTa7 z3F8nyRYsJr!XaKR!`0F2B96~Zm*M7|i9?(^j-%&HoH|jy3WxX%8E!th4xWk6&8mY# zd}aiPICb#6FfBO5XGd^|<8x=o2)hr5L!3Hc9OBf$^TOr>hxnNh9OC%gS#@xTN7W%t z9Xu~=&Txp&iQo{&=PG1`tqTru>V$EKQwPtB;t;Ql;1I{>&aQ()yfT7AoH}@3*nHp+ zKPQ4i9G{yjBW$1H5T{NUhd6ccyeJOwc^*e!%f#`y`7*+K3LN6paU6Y}5~mKH7d9U_ z#20uRJs;xu+(H>nD-(w}bsR^}hd6ccys-JeA-*VrLmZzww+;?*zmA?Waq8fCZqAuF z#Lx3MdR@fv%kyP~tqTru>Nt+pxjNlN*=Ps&)L)@>U*F~H>9!IZ>I6n7p8DZ;!L!3Hc9OBf$^P)J!FY!2fUBvOZOJ#(u3l4GWIF4Qyaq8fU zVe^4Q{IUoR@k$w&*TEqkRfl+`3_LGvK5&R%5y2sj&s|vuhqzxSz0Z{bXEl*QoI3cT z+ruoRR@Q-Uq>G=;?%+O+&xkz4)Lo!j-E4deC`?+se^BZ%?A$g8$FJm51xt7-7F*Qx`0FcrU(vk>O}b}9OAdga9Vm@#PPXXWrVE@ z4sq%@j$RjW>O}b}9OAdhaP!f1h~sm&*TEs~*U{@DPMs)Ug+u%f8DZ-pj?djGBW%ua zh*KwwL!3HMz6yuO}b}9O4hk2wN9%{PH0gVRME< zoH~x9=d)B|H_3?dRXD_#$#Cbhu0tH3d$Mv*ACVEZF5>vyav5QB zhC`e>VI1Pr!SkXx#8-G6y)NST+)5d4&Y3vGspB|$pNUf^%2(kKua*%uXX5zWsyaBt zqv{Z+PL!|0A--Bh*nEiNb8G705clioJtR(@C|`v`JSD^Jv#vuNpR1{ZL)@>Uj~8+3 zMENQl;*ZJ*yWWZ8bC1agn=>5Z)Nvd=AL7)B@>Mv*AD0m}AL97j6Ed8Z-a|OVsT0N_ zPMs)Ug+u&F8Lp0A7jb;);Ug z>*zTXr%sfw!Xf^WjIebP$LC&_5jJNy#Hka;Ax@nrUxh>b6&Yc5h~smw*1;i8onf_q zHoj7ztR^yuQwPse96>g7h`%Po%}4JuaeVG|8Ey|VafnkVj6<9{HFa=^zu|H8`9mC^ zdsBwHmNRjPQ^#@ioQYE>%2(kKe@lj2m)0VV&%G_fX=UONr;g+3`4FcLp69eOafrVY z!6A;%y;}!|xL-%lnK*UuJU8b|9O7#|j-E4d{PH~+Ve5iJoH}8g_a%m3;(1XV;vYnC zh~sk~*1;hjRfjlr@I3dLuwFYH;_D(f#PPY0WQ45?4sq&)afnj~&x_&^|2Tp}9H09? z8DZ;!L!3Hc9OBf$^TOr>hxjKE9OC%gr*&|MN7W%t9eif8DVoKj?aBw z2Zwl69pco%7sK`t4)HG{IK=U}FYDkC_v`3unK*UuJl6{{afpBAar7P%$LGG55wiYSMj+GGQ!SzIK+R5;1H)yl&``ezEMV)7IA#;M;T%3 zfXaKDT`x9O6-R zh*Jm83)6x_yjg@6aeQuvIyl6m>JX<6o)@--+)7WXZak)A&$@OBAl=}!y!(cFb;9*;CW$MaELdT&tV+m_*@GaVOnsAQzwi= zoH}@36o>e(5gg+9+-@?$<_w28b;3Brse^BZ%?A$g-93)>UOW?@+e1d!JpdfyEh9L@ zse|W*)qz8N&j=22e6E!Y*9-KV;Si@z7>77@@WrtCz#+a@1cx|2w|5;J;!$;oQwPrr zn=>5Z`$TYv<8!TLgslq>aq5I|h*Jm8i{cP(6Tu;l&$X?CLp-Vuaq8fCVe^4Qyj=u` zI6k+pjIe!%L!3Hc9OBf$^P)J!_wzXV9)LJLx4(?A>l6-g>Nt)*=ZRAX&kLIm9O4If z96cZ6_}qaqoK_|daq2jZo)2;A;EQ4NfkXVD2o7<4?%+B&#Qi#Y&cvyM=eap&;t+4| zarC-~0^;)h0Xh~sk|>fjKMszaPQcwX54z#-l- zfm(y=U2up~CyYa!I(S|bhxlO;9OC#~XBlDZfKWLLmZzwq7Dx6s5-={ zgXe|aBf%lwBZ5O5pX(_jY+Z1OQzwi=oH}@3*nHp+?-juzj?eX$5vB!)ICa7}#HoWX zMsbMuiQo`7J|~~U)&-}p4C;h&h*JmO44V%e;@KWY`!k-2&-IfLcAo%;cp`#BoH}@3 zSRFXT`$uqy<8w#KaJ@j!84hvkgmH*d2VV@E4;3ed~R?Z9O6-Rh*Jm83!4ud;zJ@h#PPXf zWQ6TA9OBdo;}E9~o)^U-KGfsr`xfH(+%R|{uLHm_agU?VdE(T;^TOr>hxl-hqvt~$ zpBo_~Y(8*^Q^#@ie27yAUksZM9O5G*IK=U}W9#4$_v`3&5vLBG=jNP=LwuCS(d#0P zUyha$wk|ltsT0N-BQg9k%2(kKKTby2IZqs)8(Rm5cvKzY)QR#{IK;=v2wN9%eC~J| zVRMEIuQ;NXelBRItIxf5iB%^41H>V$EKQzyz-;SfJjM%cQD<8vp~!66=1 zhd6bjd=(Dy2{OXwLmZzwxegBTs5-={6XmOLh@T=OY@dnam#5ajA@0}F*D`VHMENQl z;yG|bex`yrJ~vTD*qq@Ir;g+3`4Fd0l&``eo+~44KE(04JQ+?)dlVev)CuDdr%sfw z!XciN;p*sh5y$5y$#8Sd#34={$I){pPMs)Ug+qL@jIjCOnfP43jIif%aEMQd;1H({ zo)=aJ4)N0>IK=U}sWRN!_5Q#iPMt6gaq8fUVe^4QydZ)@9G{z32Zwl69pco%^TOr~ zhj?KGhd4f0BqMBHaEMbUj6<9{cwQ8TcyR=WI6ilJ9US6Ob%;|3&kLIm9O5Mr9OC#~ zsf@6FhC`e>VI1PrF}@gbp9zO}nT)VH#PPXu8DY;o;Si@z7>77@@VqcBIK-z%aERk` zGh~Em!68naFb;9*;CW&5fkS*|1cx|2H>(Z~@u)h)se|W*%^42y*%2J#_~jWg!qx?c zICUIH`}3I+!!Pl?uzLwO#LtT05Xa}{)WIPhRfjlr@I2QG^qk=kuZZ9f$LA_#gslq> zaq5I|h*Jm8i{cPJJAy+TpF2lJ*t*~lr%o7$ICb#Fu=&6tJ~x6x9G{z42Zwl69pco% z^TOr>hxq&m4sm>LK^+|8QFVw@2hR)JLpa12MsSGZbBpTW5cliodjR6p!SmcTn2AID zT#uvAdE)roc{0M*1&26w97k&rrw*PMHXk^|&-XZbKE(043uHL0OdR6W3F8o_4xSgq zA%3C9(d#0P&s`)VY+Z1OQ^#@ioQYEh-wc}%9O4(JIsN3lS!(3R8n%|v!~S@K{9BnV z|EP~BEGp)gvAh@S`QcZH^Yy6y1IB+O>)TGq2gr~w&BwOs&-PRO=A6z+Y%cNtCVz8w z-Ct4P4e1%%ojk$jI9>PIqO9>R`Pd|*v3$+=|K`KZ$?5CI|KF%mb=$XZ?X~UO`}TI@ z#TSOA_eh?nrH^?t;rYjW`&!d+bN0cyu*9RA*>!OpGZW9nqCS)7hJVbn(${0{y8nMa z=A#Go8`)d7lk2**jK5#!ZZGY=opf>I#PWTf%w}*Da>xs`4}*~-(i=xvG0AuHA{T7fW!L@JL(X5 zFzuAm{MkwQl$%qWlbc^QE35sqQ}S}ka)62o3i3E$_vW%6ZSznqB^hvgUABDK++nHJ5llo2J9 zrHpxyGt)jumdfgUg^@WWOt7pZKc{eVK@cg<$xF|6bHa#{F^V;D`jqr+q=b8r!`$V; zv39~aWwLF8w|#yG*L}$i(ss z^Gv8Sp)To1hjsHG9l?Xm5p%%ZnU+5OW77&SrzK0LaMqUEnalZ^-dQAXj7m=~HTf1B%Kg79`W?=F4GH(XgCiwm2uRi#^B^qx9x&H78GD=_EO( zRLq@{Uyx^_lD7Fl`g}1!`ViJ4+tqLp1oL^Cl<9b{czFKQ34I z_^`2^Yd!8fwv&A<5>s+Yr;N#wwV@!73zIX;(qEcm8JqGRMc5O!`a<=9zvh zTQQcCFMSFelbpp@ljRD@$t{y>(`=~BxxA25AelF^2C#qMW>9nN)I6PGO#GwP3ZcZxB5D$Xd-oD$OV=DaadBl9N6sIJxCgEG{ZA zUN|l?IyAp9nNu=0DHkLyNae~wEib4Q&MC|BUxNFx_C;mBsnVZ zo7}u%MNoVd*|~lpb4&Nqf$1|kzf{hvyfHbG50jbX7Rm7~lv;y1l{QawbXiVW5`x)F zv9;iOO6E>AM<6dbF28JwDUlYitvPMO`GwO57h)qhW5iU8O3YypwwOol4bRE2$7a_D z%4~`_du&SZaI&-xd#uu5$+UBm1<4ZE`Zh}sr58wrvZAt_f`LUPVod1=V~R>nH%2RD z={Ct>Ii;NLR%$@bv|`&m(v$N5X+LazPC*(ArbT125qnOhIiBQ>gMMrl@I`9yQ6P3&hs^_QYW=B_J!4#>Sk z$>yAlpgEfglgYf&f%4D9oZP9*tXP_DGH>v-;_`x0jzr$X$pcDqW~2=$FTWJq3@-1| z{>ic#$z-8y=hS3g|DwEn<4r|!yh`%(lEcfVjV>yX2ZQBeH;0}BI8h95s_aNw#NU|w z(xVCy;6yKzyNS(9NT0IB`O}Na1{d~|c=}70UNW(tT$1LpEt%NgeJOh|Dmh86N-^_CqDf_iVl@(+n4eP`oE9a?f}ELVGXkh|hHS@ZxyZSig9n}ELD>xH8bsvlyx^3# zXV!7GCAnsY`cq0TwH7Un)#AxM*~{VYyCtC!lO~lW%jCk`d@S4scgri9ls+z{#U-+Y z^uar^o4Kh!y*$4le_}~~`83mqENY3P%kn2pVsm^jeZpHPeepP+EI&Q*=!c42cVlvp8MGJVVNnI|}aC*?@^T<yjdU6WTLkH!QPcI4{a1nIkKR7$1$vzX@d}R=uCf(ZIsG74@+$BA+;lfORNl!Ol=jV~YB+{P?gINeD z7gc&n=D}to-0X=to4*tfkmnTUDwN|jI$1C&XjtjPDOE}{z=0v>a-!!K%IzBGtFUvq zp*7Ef-O)0yNXR-T$~9=ZFkhHvB<2=2Y{=Bzb~oz9>t|ZNrpgUP%rI#a@3VkY0!R z+DHUP#1%Ftk#q*BXug!)affZp=Bg%S_bJUwUz6!aL(Fp!sUUkgDSvX-nDHZz8YHhX zp`WGEr8Zc_x}Qy(m8~%QTZ@iFs0%HL+~u zlv$JEkX{iy#Vi~+%B~h*czl@I6susM~lTJ+6^I^Z1kJqXbTn3w; zsZT89X+(bRNM3{SVKMKFub9Y=yVptz^U}|ECYB8=FDT2m&w>0RgA3(mOckPL^=)r(=>Rw+#hJV*b1w&cTrvGh8L} zq}j!E&XMvotf;&&&(txmNVsYG0f)_!OHp1U^F@!$DK0H9sI|#X%R5Pq7K@S|`*(E? z$Dr~8@5QT$o1bT=pVyho#7)6H2e&c*1Tp=wzxY)~P@OXFi3f8$Os*@eh!-8x&qvZ~ zLCN%EbCq^B|8C~`av2G=qzB1C^`8Owi&HVbu*fCHaNE3DH$A7Llk#jgKeyX}B54R} z{$6afxtIMLV>AuGy9+i&xl>>x#W2imXDyIxycU3&)&j89T7WUu05O|ty*0p4levkN z)`H}stX(PX!iFRld3v74jTN?e^TrH4YIr00E=MOsg9BN6FI5T`|_8m{*ccv2q#af>?i-8B=6)?BiNLz+Zv4?GwMjbIkIim z7E+nt{cn`Djm&R5c?`C#AlqhbDYIxS_*OESW;GK|QnU=N;tp z&cfYMvYJViUF7SX*i9ciBxJLk;a0+Y9Te738R^0H4%0b!)hTj*;W4AN?1F| z_t?=LD_3U+{K9w_4=*_VUO+*{Sc z+YA0L$vH%3(^ziPJ z`u?)|BLy8`<|scQBdcjFoPm;ap#0ZFcK>MM50bA23ulO|vWdw*P~L^HiJXJ1?O@p# z*3!hh6XO2z{?xS}v-&3T4v9_W9S)nyx!qLWeQA| zyr%L_cTMFT=$gtq%r%vFerqD{?AAoi+9vXjYE8^Lo$W6>Y4@EqFi*R`r^wE3Y39*P zuB08Za)duo)|@N5-BQlXJQ+!u`y|2km9uTKj48s)mvNemsj{O5GN#EWlu;z3SjOox zN~Bh)j4~PJGN#LzA!DYDS!Q25%5G|noke43(P$oFikdjCryb z^JOfMu~5b$8RyD4PsaH&E|6jCa;|ljbE%2ETSF6hH-;`UtHb4;5)L=-gupd)nT*S2 zTp{C1na_V@TqR#$E#n&b`dayXot)`gNbdELeS?e}<&3wFcaX-C?-*c^lMFG;DMqZAmAY1St^ZMXz+P7VzJGNgtvaMp zhm`2G?~f`Gdb<_bxNd%u-khW}C#lIvDsqvX;jC5LVD35h=e_3wqg=F-*Vl*oP#@}J zO?_QuNYV{Sxgo)}TyI_W-nt~)BtbUW^;33<%Um(1{QtOW-Zg5~4pGi6B$tNd(L!oy zdFqsNSH`a!zhT9dXTZ(Xv0Hu`<2HAy_w?AeTO>v%V@WDxr?kc*yIVc)t@h=6?#DUJ zeID?T2_~6hnnyh52~T;(b6zmREOWf{{#U%_4f7neqc!=C1r~YhTi)^B{3Xwq`M^g$ z85?u0+5Ou5T|b-ug|DkM-(0`*gP;5|_qUoqu37HPWJ4a=Mhf>0Y;qrLrjaddrOC5I zGc9?aZ*+3zPR%??bO=ccAxR-595~mE?Qq@6E_PG18*i)jfx|lZ>pJVZbk-l#(SAtB idbdvU9-ZPxba40S&~{Eye%~JU(#}5i(~ #include +#define BT_MAX_CL_DEVICES 16 //who needs 16 devices? +//#define BT_USE_CACHE_DIR + #ifdef _WIN32 #include #include @@ -30,7 +33,7 @@ subject to the following restrictions: #endif //Set the preferred platform vendor using the OpenCL SDK -static const char* spPlatformVendor = +static char* spPlatformVendor = #if defined(CL_PLATFORM_MINI_CL) "MiniCL, SCEA"; #elif defined(CL_PLATFORM_AMD) @@ -105,7 +108,7 @@ void btOpenCLUtils::getPlatformInfo(cl_platform_id platform, btOpenCLPlatformInf oclCHECKERROR(ciErrNum,CL_SUCCESS); } -cl_context btOpenCLUtils::createContextFromPlatform(cl_platform_id platform, cl_device_type deviceType, cl_int* pErrNum, void* pGLContext, void* pGLDC) +cl_context btOpenCLUtils::createContextFromPlatform(cl_platform_id platform, cl_device_type deviceType, cl_int* pErrNum, void* pGLContext, void* pGLDC, int preferredDeviceIndex, int preferredPlatformIndex) { cl_context retContext = 0; cl_int ciErrNum=0; @@ -117,22 +120,52 @@ cl_context btOpenCLUtils::createContextFromPlatform(cl_platform_id platform, cl_ cl_context_properties cps[7] = {0,0,0,0,0,0,0}; cps[0] = CL_CONTEXT_PLATFORM; cps[1] = (cl_context_properties)platform; -#if defined (_WIN32) && defined(_MSC_VER) && !defined (CL_PLATFORM_MINI_CL) if (pGLContext && pGLDC) { +#if defined(CL_PLATFORM_AMD) || defined(CL_PLATFORM_NVIDIA) cps[2] = CL_GL_CONTEXT_KHR; cps[3] = (cl_context_properties)pGLContext; cps[4] = CL_WGL_HDC_KHR; cps[5] = (cl_context_properties)pGLDC; - } #endif + } + + cl_uint num_entries = BT_MAX_CL_DEVICES; + cl_device_id devices[BT_MAX_CL_DEVICES]; + + cl_uint num_devices=-1; + + ciErrNum = clGetDeviceIDs( + platform, + deviceType, + num_entries, + devices, + &num_devices); cl_context_properties* cprops = (NULL == platform) ? NULL : cps; - retContext = clCreateContextFromType(cprops, - deviceType, - NULL, - NULL, - &ciErrNum); + + if (pGLContext) + { + //search for the GPU that relates to the OpenCL context + for (int i=0;i=0 && preferredDeviceIndex=0 && i==preferredPlatformIndex) { cl_platform_id tmpPlatform = platforms[0]; platforms[0] = platforms[i]; platforms[i] = tmpPlatform; break; + } else + { + if(!strcmp(pbuf, spPlatformVendor)) + { + cl_platform_id tmpPlatform = platforms[0]; + platforms[0] = platforms[i]; + platforms[i] = tmpPlatform; + break; + } } } @@ -191,11 +234,11 @@ cl_context btOpenCLUtils::createContextFromType(cl_device_type deviceType, cl_in cl_platform_id platform = platforms[i]; assert(platform); - retContext = btOpenCLUtils::createContextFromPlatform(platform,deviceType,pErrNum,pGLContext,pGLDC); + retContext = btOpenCLUtils::createContextFromPlatform(platform,deviceType,pErrNum,pGLContext,pGLDC,preferredDeviceIndex); if (retContext) { - printf("OpenCL platform details:\n"); +// printf("OpenCL platform details:\n"); btOpenCLPlatformInfo platformInfo; btOpenCLUtils::getPlatformInfo(platform, platformInfo); @@ -270,10 +313,10 @@ void btOpenCLUtils::printDeviceInfo(cl_device_id device) printf(" CL_DEVICE_TYPE:\t\t\t%s\n", "CL_DEVICE_TYPE_DEFAULT"); printf(" CL_DEVICE_MAX_COMPUTE_UNITS:\t\t%u\n", info.m_computeUnits); - printf(" CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS:\t%d\n", (int)info.m_workitemDims); - printf(" CL_DEVICE_MAX_WORK_ITEM_SIZES:\t%d / %d / %d \n", (int)info.m_workItemSize[0], (int)info.m_workItemSize[1],(int) info.m_workItemSize[2]); - printf(" CL_DEVICE_MAX_WORK_GROUP_SIZE:\t%d\n", (int)info.m_workgroupSize); - printf(" CL_DEVICE_MAX_CLOCK_FREQUENCY:\t%d MHz\n", (int)info.m_clockFrequency); + printf(" CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS:\t%u\n", info.m_workitemDims); + printf(" CL_DEVICE_MAX_WORK_ITEM_SIZES:\t%u / %u / %u \n", info.m_workItemSize[0], info.m_workItemSize[1], info.m_workItemSize[2]); + printf(" CL_DEVICE_MAX_WORK_GROUP_SIZE:\t%u\n", info.m_workgroupSize); + printf(" CL_DEVICE_MAX_CLOCK_FREQUENCY:\t%u MHz\n", info.m_clockFrequency); printf(" CL_DEVICE_ADDRESS_BITS:\t\t%u\n", info.m_addressBits); printf(" CL_DEVICE_MAX_MEM_ALLOC_SIZE:\t\t%u MByte\n", (unsigned int)(info.m_maxMemAllocSize/ (1024 * 1024))); printf(" CL_DEVICE_GLOBAL_MEM_SIZE:\t\t%u MByte\n", (unsigned int)(info.m_globalMemSize/ (1024 * 1024))); @@ -291,11 +334,11 @@ void btOpenCLUtils::printDeviceInfo(cl_device_id device) printf(" CL_DEVICE_MAX_READ_IMAGE_ARGS:\t%u\n", info.m_maxReadImageArgs); printf(" CL_DEVICE_MAX_WRITE_IMAGE_ARGS:\t%u\n", info.m_maxWriteImageArgs); printf("\n CL_DEVICE_IMAGE "); - printf("\t\t\t2D_MAX_WIDTH\t %d\n", (int)info.m_image2dMaxWidth); - printf("\t\t\t\t\t2D_MAX_HEIGHT\t %d\n", (int)info.m_image2dMaxHeight); - printf("\t\t\t\t\t3D_MAX_WIDTH\t %d\n", (int)info.m_image3dMaxWidth); - printf("\t\t\t\t\t3D_MAX_HEIGHT\t %d\n", (int)info.m_image3dMaxHeight); - printf("\t\t\t\t\t3D_MAX_DEPTH\t %d\n", (int)info.m_image3dMaxDepth); + printf("\t\t\t2D_MAX_WIDTH\t %u\n", info.m_image2dMaxWidth); + printf("\t\t\t\t\t2D_MAX_HEIGHT\t %u\n", info.m_image2dMaxHeight); + printf("\t\t\t\t\t3D_MAX_WIDTH\t %u\n", info.m_image3dMaxWidth); + printf("\t\t\t\t\t3D_MAX_HEIGHT\t %u\n", info.m_image3dMaxHeight); + printf("\t\t\t\t\t3D_MAX_DEPTH\t %u\n", info.m_image3dMaxDepth); if (info.m_deviceExtensions != 0) printf("\n CL_DEVICE_EXTENSIONS:%s\n",info.m_deviceExtensions); else @@ -389,21 +432,84 @@ void btOpenCLUtils::getDeviceInfo(cl_device_id device, btOpenCLDeviceInfo& info) clGetDeviceInfo(device, CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE, sizeof(cl_uint), &info.m_vecWidthDouble, NULL); } -static const char* strip2(const char* name, const char* pattern) +static char* strip1(char* name, const char* pattern,int* numOccurences=0) +{ + size_t const patlen = strlen(pattern); + char * oriptr; + char * patloc; + // find how many times the pattern occurs in the original string + for (oriptr = name; patloc = strstr(oriptr, pattern); oriptr = patloc + patlen) + { + if (numOccurences) + (*numOccurences)++; + } + return oriptr; +} +static const char* strip2(const char* name, const char* pattern,int* numOccurences=0) { size_t const patlen = strlen(pattern); - size_t patcnt = 0; const char * oriptr; const char * patloc; // find how many times the pattern occurs in the original string for (oriptr = name; patloc = strstr(oriptr, pattern); oriptr = patloc + patlen) { - patcnt++; + if (numOccurences) + (*numOccurences)++; } return oriptr; } -cl_program btOpenCLUtils::compileCLProgramFromString(cl_context clContext, cl_device_id device, const char* kernelSource, cl_int* pErrNum, const char* additionalMacros , const char* clFileNameForCaching) +cl_program btOpenCLUtils::compileCLProgramFromString(cl_context clContext, cl_device_id device, const char* kernelSource, cl_int* pErrNum, const char* additionalMacros) +{ + + cl_int localErrNum; + size_t program_length = strlen(kernelSource); + + cl_program m_cpProgram = clCreateProgramWithSource(clContext, 1, (const char**)&kernelSource, &program_length, &localErrNum); + if (localErrNum!= CL_SUCCESS) + { + if (pErrNum) + *pErrNum = localErrNum; + return 0; + } + + // Build the program with 'mad' Optimization option + + +#ifdef MAC + char* flags = "-cl-mad-enable -DMAC -DGUID_ARG"; +#else + //const char* flags = "-DGUID_ARG= -fno-alias"; + const char* flags = "-DGUID_ARG= "; +#endif + + char* compileFlags = new char[strlen(additionalMacros) + strlen(flags) + 5]; + sprintf(compileFlags, "%s %s", flags, additionalMacros); + localErrNum = clBuildProgram(m_cpProgram, 1, &device, compileFlags, NULL, NULL); + if (localErrNum!= CL_SUCCESS) + { + char *build_log; + size_t ret_val_size; + clGetProgramBuildInfo(m_cpProgram, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size); + build_log = new char[ret_val_size+1]; + clGetProgramBuildInfo(m_cpProgram, device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL); + + // to be carefully, terminate with \0 + // there's no information in the reference whether the string is 0 terminated or not + build_log[ret_val_size] = '\0'; + + + printf("Error in clBuildProgram, Line %u in file %s, Log: \n%s\n !!!\n\n", __LINE__, __FILE__, build_log); + delete[] build_log; + if (pErrNum) + *pErrNum = localErrNum; + return 0; + } + delete[] compileFlags; + return m_cpProgram; +} + +cl_program btOpenCLUtils::compileCLProgramFromFile(cl_context clContext, cl_device_id device, cl_int* pErrNum, const char* additionalMacros , const char* clFileNameForCaching) { cl_program m_cpProgram=0; @@ -411,7 +517,6 @@ cl_program btOpenCLUtils::compileCLProgramFromString(cl_context clContext, cl_de char binaryFileName[522]; -#if defined (_WIN32) && defined(_MSC_VER) if (clFileNameForCaching) { @@ -423,15 +528,23 @@ cl_program btOpenCLUtils::compileCLProgramFromString(cl_context clContext, cl_de const char* strippedName = strip2(clFileNameForCaching,"\\"); strippedName = strip2(strippedName,"/"); - +#ifdef BT_USE_CACHE_DIR + sprintf_s(binaryFileName,"cache/%s.%s.%s.bin",strippedName, deviceName,driverVersion ); +#else sprintf_s(binaryFileName,"%s.%s.%s.bin",strippedName, deviceName,driverVersion ); - //printf("searching for %s\n", binaryFileName); +#endif + + //printf("searching for %s\n", binaryFileName); bool fileUpToDate = false; bool binaryFileValid=false; FILETIME modtimeBinary; +#ifdef _WIN32 +#ifdef BT_USE_CACHE_DIR + CreateDirectory("cache",0); +#endif //BT_USE_CACHE_DIR { HANDLE binaryFileHandle = CreateFile(binaryFileName,GENERIC_READ,0,0,OPEN_EXISTING,FILE_ATTRIBUTE_NORMAL,0); @@ -494,7 +607,6 @@ cl_program btOpenCLUtils::compileCLProgramFromString(cl_context clContext, cl_de } else { -#ifdef _DEBUG DWORD errorCode; errorCode = GetLastError(); switch (errorCode) @@ -516,11 +628,7 @@ cl_program btOpenCLUtils::compileCLProgramFromString(cl_context clContext, cl_de } //we should make sure the src file exists so we can verify the timestamp with binary - assert(0); -#else - //if we cannot find the source, assume it is OK in release builds - fileUpToDate = true; -#endif + fileUpToDate = false; } } @@ -557,88 +665,88 @@ cl_program btOpenCLUtils::compileCLProgramFromString(cl_context clContext, cl_de btAssert(0); m_cpProgram = 0; } + delete[] binary; } } +#endif //_WIN32 } -#endif if (!m_cpProgram) { -// cl_kernel kernel; - cl_int localErrNum; - size_t program_length = strlen(kernelSource); - - m_cpProgram = clCreateProgramWithSource(clContext, 1, (const char**)&kernelSource, &program_length, &localErrNum); - if (localErrNum!= CL_SUCCESS) + + FILE* file = fopen(clFileNameForCaching, "r"); + if (file) { - if (pErrNum) - *pErrNum = localErrNum; - return 0; - } - - // Build the program with 'mad' Optimization option - - - #ifdef MAC - char* flags = "-cl-mad-enable -DMAC -DGUID_ARG"; - #else - //const char* flags = "-DGUID_ARG= -fno-alias"; - const char* flags = "-DGUID_ARG= "; - #endif - - char* compileFlags = new char[strlen(additionalMacros) + strlen(flags) + 5]; - sprintf(compileFlags, "%s %s", flags, additionalMacros); - localErrNum = clBuildProgram(m_cpProgram, 1, &device, compileFlags, NULL, NULL); - if (localErrNum!= CL_SUCCESS) - { - char *build_log; - size_t ret_val_size; - clGetProgramBuildInfo(m_cpProgram, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size); - build_log = new char[ret_val_size+1]; - clGetProgramBuildInfo(m_cpProgram, device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL); - - // to be carefully, terminate with \0 - // there's no information in the reference whether the string is 0 terminated or not - build_log[ret_val_size] = '\0'; - - - printf("Error in clBuildProgram, Line %u in file %s, Log: \n%s\n !!!\n\n", __LINE__, __FILE__, build_log); - delete[] build_log; - if (pErrNum) - *pErrNum = localErrNum; - return 0; - } - -#if defined (_WIN32) && defined(_MSC_VER) - if( clFileNameForCaching ) - { // write to binary - size_t binarySize; - status = clGetProgramInfo( m_cpProgram, CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &binarySize, 0 ); - btAssert( status == CL_SUCCESS ); - - char* binary = new char[binarySize]; - - status = clGetProgramInfo( m_cpProgram, CL_PROGRAM_BINARIES, sizeof(char*), &binary, 0 ); - btAssert( status == CL_SUCCESS ); - + fseek( file, 0L, SEEK_END ); + size_t fileSize= ftell( file ); + rewind( file ); + char* kernelSource2 = new char[fileSize+1]; + fread( kernelSource2, sizeof(char), fileSize, file ); + fclose( file ); + kernelSource2[fileSize]=0; + int numOccurences = 0; + ///patch/remove the MSTRINGIFY( and ); + char* kernelSource = strip1(kernelSource2,"MSTRINGIFY(",&numOccurences); + int newlen = strlen(kernelSource); + if (numOccurences) { - FILE* file = fopen(binaryFileName, "wb"); - if (file) + int i=newlen-1; + + for (;i>=0;i--) { - fwrite( binary, sizeof(char), binarySize, file ); - fclose( file ); - } else + if (kernelSource[i] == ';') + { + kernelSource[i] = 0;//' '; + break; + } + } + for (;i>=0;i--) { - printf("cannot write file %s\n", binaryFileName); + if (kernelSource[i] == ')') + { + kernelSource[i] = 0;//' '; + break; + } } } + + m_cpProgram = compileCLProgramFromString(clContext,device,kernelSource,pErrNum,additionalMacros); + + if( clFileNameForCaching ) + { // write to binary - delete [] binary; + cl_uint numAssociatedDevices; + status = clGetProgramInfo( m_cpProgram, CL_PROGRAM_NUM_DEVICES, sizeof(cl_uint), &numAssociatedDevices, 0 ); + btAssert( status == CL_SUCCESS ); + if (numAssociatedDevices==1) + { + + size_t binarySize; + status = clGetProgramInfo( m_cpProgram, CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &binarySize, 0 ); + btAssert( status == CL_SUCCESS ); + + char* binary = new char[binarySize]; + + status = clGetProgramInfo( m_cpProgram, CL_PROGRAM_BINARIES, sizeof(char*), &binary, 0 ); + btAssert( status == CL_SUCCESS ); + + { + FILE* file = fopen(binaryFileName, "wb"); + if (file) + { + fwrite( binary, sizeof(char), binarySize, file ); + fclose( file ); + } else + { + printf("cannot write file %s\n", binaryFileName); + } + } + + delete [] binary; + } + } } -#endif//defined (_WIN32) && defined(_MSC_VER) - - delete [] compileFlags; } return m_cpProgram; diff --git a/Demos/SharedOpenCL/btOpenCLUtils.h b/Demos/SharedOpenCL/btOpenCLUtils.h index 5e5d20467..e9f422897 100644 --- a/Demos/SharedOpenCL/btOpenCLUtils.h +++ b/Demos/SharedOpenCL/btOpenCLUtils.h @@ -78,7 +78,7 @@ public: /// CL Context optionally takes a GL context. This is a generic type because we don't really want this code /// to have to understand GL types. It is a HGLRC in _WIN32 or a GLXContext otherwise. - static cl_context createContextFromType(cl_device_type deviceType, cl_int* pErrNum, void* pGLCtx = 0, void* pGLDC = 0); + static cl_context createContextFromType(cl_device_type deviceType, cl_int* pErrNum, void* pGLCtx = 0, void* pGLDC = 0, int preferredDeviceIndex = -1, int preferredPlatformIndex= - 1); static int getNumDevices(cl_context cxMainContext); static cl_device_id getDevice(cl_context cxMainContext, int nr); @@ -88,7 +88,10 @@ public: static cl_kernel compileCLKernelFromString( cl_context clContext,cl_device_id device, const char* kernelSource, const char* kernelName, cl_int* pErrNum=0, cl_program prog=0,const char* additionalMacros = "" ); //optional - static cl_program compileCLProgramFromString( cl_context clContext,cl_device_id device, const char* kernelSource, cl_int* pErrNum=0,const char* additionalMacros = "" , const char* srcFileNameForCaching=0); + static cl_program compileCLProgramFromString( cl_context clContext,cl_device_id device, const char* kernelSource, cl_int* pErrNum=0,const char* additionalMacros = ""); + ///compileCLProgramFromFile will attempt to save/load the binary precompiled program + static cl_program compileCLProgramFromFile( cl_context clContext,cl_device_id device, cl_int* pErrNum=0,const char* additionalMacros = "" , const char* srcFileNameForCaching=0); + //the following optional APIs provide access using specific platform information static int getNumPlatforms(cl_int* pErrNum=0); @@ -96,7 +99,7 @@ public: static cl_platform_id getPlatform(int nr, cl_int* pErrNum=0); static void getPlatformInfo(cl_platform_id platform, btOpenCLPlatformInfo& platformInfo); static const char* getSdkVendorName(); - static cl_context createContextFromPlatform(cl_platform_id platform, cl_device_type deviceType, cl_int* pErrNum, void* pGLCtx = 0, void* pGLDC = 0); + static cl_context createContextFromPlatform(cl_platform_id platform, cl_device_type deviceType, cl_int* pErrNum, void* pGLCtx = 0, void* pGLDC = 0,int preferredDeviceIndex = -1, int preferredPlatformIndex= -1); }; diff --git a/Extras/Serialize/BulletWorldImporter/btBulletWorldImporter.cpp b/Extras/Serialize/BulletWorldImporter/btBulletWorldImporter.cpp index 96b8b2ee6..7e58f3c53 100644 --- a/Extras/Serialize/BulletWorldImporter/btBulletWorldImporter.cpp +++ b/Extras/Serialize/BulletWorldImporter/btBulletWorldImporter.cpp @@ -1,3 +1,18 @@ +/* +Bullet Continuous Collision Detection and Physics Library +Copyright (c) 2003-2010 Erwin Coumans http://continuousphysics.com/Bullet/ + +This software is provided 'as-is', without any express or implied warranty. +In no event will the authors be held liable for any damages arising from the use of this software. +Permission is granted to anyone to use this software for any purpose, +including commercial applications, and to alter it and redistribute it freely, +subject to the following restrictions: + +1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required. +2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software. +3. This notice may not be removed or altered from any source distribution. +*/ + #include "btBulletWorldImporter.h" #include "../BulletFileLoader/btBulletFile.h" @@ -454,23 +469,24 @@ btCollisionShape* btBulletWorldImporter::convertCollisionShape( btCollisionShap { case 0: { - shape = createCapsuleShapeX(implicitShapeDimensions.getY(),2*implicitShapeDimensions.getX()); + shape = createCapsuleShapeX(implicitShapeDimensions.getY()+bsd->m_collisionMargin*2,2*implicitShapeDimensions.getX()); break; } case 1: { - shape = createCapsuleShapeY(implicitShapeDimensions.getX(),2*implicitShapeDimensions.getY()); + shape = createCapsuleShapeY(implicitShapeDimensions.getX()+bsd->m_collisionMargin*2,2*implicitShapeDimensions.getY()); break; } case 2: { - shape = createCapsuleShapeZ(implicitShapeDimensions.getX(),2*implicitShapeDimensions.getZ()); + shape = createCapsuleShapeZ(implicitShapeDimensions.getX()+bsd->m_collisionMargin*2,2*implicitShapeDimensions.getZ()); break; } default: { printf("error: wrong up axis for btCapsuleShape\n"); } + bsd->m_collisionMargin = 0.f; }; @@ -567,7 +583,7 @@ btCollisionShape* btBulletWorldImporter::convertCollisionShape( btCollisionShap if (shape) { - shape->setMargin(bsd->m_collisionMargin); + //shape->setMargin(bsd->m_collisionMargin); btVector3 localScaling; localScaling.deSerializeFloat(bsd->m_localScaling); shape->setLocalScaling(localScaling); @@ -654,7 +670,9 @@ btCollisionShape* btBulletWorldImporter::convertCollisionShape( btCollisionShap compoundShape->addChildShape(localTransform,childShape); } else { +#ifdef _DEBUG printf("error: couldn't create childShape for compoundShape\n"); +#endif } } @@ -668,7 +686,9 @@ btCollisionShape* btBulletWorldImporter::convertCollisionShape( btCollisionShap } default: { +#ifdef _DEBUG printf("unsupported shape type (%d)\n",shapeData->m_shapeType); +#endif } } @@ -827,6 +847,10 @@ bool btBulletWorldImporter::convertAllObjects( bParse::btBulletFile* bulletFile } bool isDynamic = mass!=0.f; btRigidBody* body = createRigidBody(isDynamic,mass,startTransform,shape,colObjData->m_collisionObjectData.m_name); + body->setFriction(colObjData->m_collisionObjectData.m_friction); + body->setRestitution(colObjData->m_collisionObjectData.m_restitution); + + #ifdef USE_INTERNAL_EDGE_UTILITY if (shape->getShapeType() == TRIANGLE_MESH_SHAPE_PROXYTYPE) { @@ -857,7 +881,9 @@ bool btBulletWorldImporter::convertAllObjects( bParse::btBulletFile* bulletFile startTransform.deSerializeDouble(colObjData->m_worldTransform); btCollisionShape* shape = (btCollisionShape*)*shapePtr; btCollisionObject* body = createCollisionObject(startTransform,shape,colObjData->m_name); - + body->setFriction(colObjData->m_friction); + body->setRestitution(colObjData->m_restitution); + #ifdef USE_INTERNAL_EDGE_UTILITY if (shape->getShapeType() == TRIANGLE_MESH_SHAPE_PROXYTYPE) { diff --git a/Extras/Serialize/BulletWorldImporter/btBulletWorldImporter.h b/Extras/Serialize/BulletWorldImporter/btBulletWorldImporter.h index ee57c2e2e..05c3ccbfc 100644 --- a/Extras/Serialize/BulletWorldImporter/btBulletWorldImporter.h +++ b/Extras/Serialize/BulletWorldImporter/btBulletWorldImporter.h @@ -55,6 +55,10 @@ namespace bParse }; + +///The btBulletWorldImporter is a starting point to import .bullet files. +///note that not all data is converted yet. You are expected to override or modify this class. +///See Bullet/Demos/SerializeDemo for a derived class that extract btSoftBody objects too. class btBulletWorldImporter { protected: diff --git a/src/BulletMultiThreaded/GpuSoftBodySolvers/DX11/btSoftBodySolver_DX11.cpp b/src/BulletMultiThreaded/GpuSoftBodySolvers/DX11/btSoftBodySolver_DX11.cpp index 5efd3576e..b66722bdc 100644 --- a/src/BulletMultiThreaded/GpuSoftBodySolvers/DX11/btSoftBodySolver_DX11.cpp +++ b/src/BulletMultiThreaded/GpuSoftBodySolvers/DX11/btSoftBodySolver_DX11.cpp @@ -20,7 +20,7 @@ subject to the following restrictions: #include "btSoftBodySolverVertexBuffer_DX11.h" #include "BulletSoftBody/btSoftBody.h" #include "BulletCollision/CollisionShapes/btCapsuleShape.h" - +#include //printf #define MSTRINGIFY(A) #A static char* PrepareLinksHLSLString = #include "HLSL/PrepareLinks.hlsl" @@ -2190,7 +2190,9 @@ void btDX11SoftBodySolver::processCollision( btSoftBody *softBody, btCollisionOb m_collisionObjectDetails.push_back( newCollisionShapeDescription ); } else { - btAssert("Unsupported collision shape type\n"); +#ifdef _DEBUG + printf("Unsupported collision shape type\n"); +#endif } } else { btAssert("Unknown soft body"); diff --git a/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/OpenCLC10/SolveCollisionsAndUpdateVelocitiesSIMDBatched.cl b/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/OpenCLC10/SolveCollisionsAndUpdateVelocitiesSIMDBatched.cl index ba57c8869..870902258 100644 --- a/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/OpenCLC10/SolveCollisionsAndUpdateVelocitiesSIMDBatched.cl +++ b/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/OpenCLC10/SolveCollisionsAndUpdateVelocitiesSIMDBatched.cl @@ -1,9 +1,23 @@ MSTRINGIFY( -float mydot3(float4 a, float4 b) + +//#pragma OPENCL EXTENSION cl_amd_printf:enable\n + +float mydot3a(float4 a, float4 b) { return a.x*b.x + a.y*b.y + a.z*b.z; } +float mylength3(float4 a) +{ + a.w = 0; + return length(a); +} + +float4 mynormalize3(float4 a) +{ + a.w = 0; + return normalize(a); +} typedef struct { @@ -37,8 +51,7 @@ typedef struct // From btBroadphaseProxy.h __constant int CAPSULE_SHAPE_PROXYTYPE = 10; - -/* Multiply column-major matrix against vector */ +// Multiply column-major matrix against vector float4 matrixVectorMul( float4 matrix[4], float4 vector ) { float4 returnVector; @@ -66,7 +79,8 @@ SolveCollisionsAndUpdateVelocitiesKernel( __global float4 * g_vertexForces, __global float4 *g_vertexVelocities, __global float4 *g_vertexPositions, - __local CollisionShapeDescription *localCollisionShapes) + __local CollisionShapeDescription *localCollisionShapes, + __global float * g_vertexInverseMasses) { int nodeID = get_global_id(0); float4 forceOnVertex = (float4)(0.f, 0.f, 0.f, 0.f); @@ -78,18 +92,20 @@ SolveCollisionsAndUpdateVelocitiesKernel( return; - float4 position = (float4)(g_vertexPositions[nodeID].xyz, 1.f); - float4 previousPosition = (float4)(g_vertexPreviousPositions[nodeID].xyz, 1.f); + float4 position = (float4)(g_vertexPositions[nodeID].xyz, 0.f); + float4 previousPosition = (float4)(g_vertexPreviousPositions[nodeID].xyz, 0.f); + float clothFriction = g_perClothFriction[clothIdentifier]; float dampingFactor = g_clothDampingFactor[clothIdentifier]; float velocityCoefficient = (1.f - dampingFactor); - - // Update velocity float4 difference = position - previousPosition; float4 velocity = difference*velocityCoefficient*isolverdt; + float inverseMass = g_vertexInverseMasses[nodeID]; + CollisionObjectIndices collisionObjectIndices = g_perClothCollisionObjectIndices[clothIdentifier]; int numObjects = collisionObjectIndices.endObject - collisionObjectIndices.firstObject; + if( numObjects > 0 ) { // We have some possible collisions to deal with @@ -113,7 +129,7 @@ SolveCollisionsAndUpdateVelocitiesKernel( // We have some possible collisions to deal with for( int collision = 0; collision < numObjects; ++collision ) { - //CollisionShapeDescription shapeDescription = localCollisionShapes[collision]; + CollisionShapeDescription shapeDescription = localCollisionShapes[collision]; float colliderFriction = localCollisionShapes[collision].friction; if( localCollisionShapes[collision].collisionShapeType == CAPSULE_SHAPE_PROXYTYPE ) @@ -125,14 +141,14 @@ SolveCollisionsAndUpdateVelocitiesKernel( float capsuleMargin = localCollisionShapes[collision].margin; int capsuleupAxis = localCollisionShapes[collision].upAxis; + if ( capsuleHalfHeight <= 0 ) + capsuleHalfHeight = 0.0001f; float4 worldTransform[4]; worldTransform[0] = localCollisionShapes[collision].shapeTransform[0]; worldTransform[1] = localCollisionShapes[collision].shapeTransform[1]; worldTransform[2] = localCollisionShapes[collision].shapeTransform[2]; worldTransform[3] = localCollisionShapes[collision].shapeTransform[3]; - //float4 c1 = (float4)(0.f, -capsuleHalfHeight, 0.f, 1.f); - //float4 c2 = (float4)(0.f, +capsuleHalfHeight, 0.f, 1.f); // Correctly define capsule centerline vector float4 c1 = (float4)(0.f, 0.f, 0.f, 1.f); float4 c2 = (float4)(0.f, 0.f, 0.f, 1.f); @@ -145,65 +161,72 @@ SolveCollisionsAndUpdateVelocitiesKernel( float4 worldC1 = matrixVectorMul(worldTransform, c1); float4 worldC2 = matrixVectorMul(worldTransform, c2); - float4 segment = (worldC2 - worldC1); + float4 segment = (float4)((worldC2 - worldC1).xyz, 0.f); + float4 segmentNormalized = mynormalize3(segment); + float distanceAlongSegment =mydot3a( (position - worldC1), segmentNormalized ); - // compute distance of tangent to vertex along line segment in capsule - float distanceAlongSegment = -( mydot3( (worldC1 - position), segment ) / mydot3(segment, segment) ); - - float4 closestPoint = (worldC1 + (float4)(segment * distanceAlongSegment)); - float distanceFromLine = length(position - closestPoint); - float distanceFromC1 = length(worldC1 - position); - float distanceFromC2 = length(worldC2 - position); - + float4 closestPointOnSegment = (worldC1 + (float4)(segmentNormalized * distanceAlongSegment)); + float distanceFromLine = mylength3(position - closestPointOnSegment); + float distanceFromC1 = mylength3(worldC1 - position); + float distanceFromC2 = mylength3(worldC2 - position); + // Final distance from collision, point to push from, direction to push in // for impulse force float dist; float4 normalVector; + if( distanceAlongSegment < 0 ) { dist = distanceFromC1; - normalVector = normalize(position - worldC1); - } else if( distanceAlongSegment > 1.f ) { + normalVector = (float4)(normalize(position - worldC1).xyz, 0.f); + } else if( distanceAlongSegment > length(segment) ) { dist = distanceFromC2; - normalVector = normalize(position - worldC2); + normalVector = (float4)(normalize(position - worldC2).xyz, 0.f); } else { dist = distanceFromLine; - normalVector = normalize(position - closestPoint); + normalVector = (float4)(normalize(position - closestPointOnSegment).xyz, 0.f); } - float4 colliderLinearVelocity = localCollisionShapes[collision].linearVelocity; - float4 colliderAngularVelocity = localCollisionShapes[collision].angularVelocity; - float4 velocityOfSurfacePoint = colliderLinearVelocity + cross(colliderAngularVelocity, position - (float4)(worldTransform[0].w, worldTransform[1].w, worldTransform[2].w, 0.f)); - float minDistance = capsuleRadius + capsuleMargin; - - // In case of no collision, this is the value of velocity - velocity = (position - previousPosition) * velocityCoefficient * isolverdt; + float4 closestPointOnSurface = (float4)((position + (minDistance - dist) * normalVector).xyz, 0.f); + + float4 colliderLinearVelocity = shapeDescription.linearVelocity; + float4 colliderAngularVelocity = shapeDescription.angularVelocity; + float4 velocityOfSurfacePoint = colliderLinearVelocity + cross(colliderAngularVelocity, closestPointOnSurface - (float4)(worldTransform[0].w, worldTransform[1].w, worldTransform[2].w, 0.f)); // Check for a collision if( dist < minDistance ) { // Project back to surface along normal - position = position + (float4)((minDistance - dist)*normalVector*0.9f); + position = closestPointOnSurface; velocity = (position - previousPosition) * velocityCoefficient * isolverdt; float4 relativeVelocity = velocity - velocityOfSurfacePoint; - float4 p1 = (float4)(normalize(cross(normalVector, segment)).xyz, 0.f); - float4 p2 = (float4)(normalize(cross(p1, normalVector)).xyz, 0.f); - // Full friction is sum of velocities in each direction of plane - float4 frictionVector = p1*mydot3(relativeVelocity, p1) + p2*mydot3(relativeVelocity, p2); - - // Real friction is peak friction corrected by friction coefficients - frictionVector = frictionVector * (colliderFriction*clothFriction); - - float approachSpeed = dot(relativeVelocity, normalVector); - - if( approachSpeed <= 0.0f ) - forceOnVertex -= frictionVector; - } + float4 p1 = mynormalize3(cross(normalVector, segment)); + float4 p2 = mynormalize3(cross(p1, normalVector)); + float4 tangentialVel = p1*mydot3a(relativeVelocity, p1) + p2*mydot3a(relativeVelocity, p2); + float frictionCoef = (colliderFriction * clothFriction); + if (frictionCoef>1.f) + frictionCoef = 1.f; + + //only apply friction if objects are not moving apart + float projVel = mydot3a(relativeVelocity,normalVector); + if ( projVel >= -0.001f) + { + if ( inverseMass > 0 ) + { + //float4 myforceOnVertex = -tangentialVel * frictionCoef * isolverdt * (1.0f / inverseMass); + position += (-tangentialVel * frictionCoef) / (isolverdt); + } + } + + // In case of no collision, this is the value of velocity + velocity = (position - previousPosition) * velocityCoefficient * isolverdt; + + } } } } diff --git a/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/btSoftBodySolver_OpenCL.cpp b/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/btSoftBodySolver_OpenCL.cpp index a9b5d01c3..0a5ca008b 100644 --- a/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/btSoftBodySolver_OpenCL.cpp +++ b/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/btSoftBodySolver_OpenCL.cpp @@ -739,7 +739,7 @@ void btOpenCLSoftBodySolver::optimize( btAlignedObjectArray< btSoftBody * > &sof m_perClothDragFactor.push_back( softBody->m_cfg.kDG ); m_perClothMediumDensity.push_back(softBody->getWorldInfo()->air_density); // Simple init values. Actually we'll put 0 and -1 into them at the appropriate time - m_perClothFriction.push_back( softBody->getFriction() ); + m_perClothFriction.push_back(softBody->m_cfg.kDF); m_perClothCollisionObjects.push_back( CollisionObjectIndices(-1, -1) ); // Add space for new vertices and triangles in the default solver for now @@ -1737,7 +1737,9 @@ void btOpenCLSoftBodySolver::processCollision( btSoftBody *softBody, btCollision } else { +#ifdef _DEBUG printf("Unsupported collision shape type\n"); +#endif //btAssert(0 && "Unsupported collision shape type\n"); } } else { diff --git a/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/btSoftBodySolver_OpenCLSIMDAware.cpp b/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/btSoftBodySolver_OpenCLSIMDAware.cpp index e9c6f7de4..2216768a9 100644 --- a/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/btSoftBodySolver_OpenCLSIMDAware.cpp +++ b/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/btSoftBodySolver_OpenCLSIMDAware.cpp @@ -217,7 +217,7 @@ void btOpenCLSoftBodySolverSIMDAware::optimize( btAlignedObjectArray< btSoftBody m_perClothDragFactor.push_back( softBody->m_cfg.kDG ); m_perClothMediumDensity.push_back(softBody->getWorldInfo()->air_density); // Simple init values. Actually we'll put 0 and -1 into them at the appropriate time - m_perClothFriction.push_back( softBody->getFriction() ); + m_perClothFriction.push_back(softBody->m_cfg.kDF); m_perClothCollisionObjects.push_back( CollisionObjectIndices(-1, -1) ); // Add space for new vertices and triangles in the default solver for now @@ -253,6 +253,10 @@ void btOpenCLSoftBodySolverSIMDAware::optimize( btAlignedObjectArray< btSoftBody m_anchorIndex.push_back(-1.0); } + for( int vertex = numVertices; vertex < maxVertices; ++vertex ) + { + m_anchorIndex.push_back(-1.0); + } // Copy triangles similarly // We're assuming here that vertex indices are based on the firstVertex rather than the entire scene @@ -524,6 +528,7 @@ void btOpenCLSoftBodySolverSIMDAware::solveCollisionsAndUpdateVelocities( float ciErrNum = clSetKernelArg(m_solveCollisionsAndUpdateVelocitiesKernel, 9, sizeof(cl_mem),&m_vertexData.m_clVertexVelocity.m_buffer); ciErrNum = clSetKernelArg(m_solveCollisionsAndUpdateVelocitiesKernel, 10, sizeof(cl_mem),&m_vertexData.m_clVertexPosition.m_buffer); ciErrNum = clSetKernelArg(m_solveCollisionsAndUpdateVelocitiesKernel, 11, sizeof(CollisionShapeDescription)*16,0); + ciErrNum = clSetKernelArg(m_solveCollisionsAndUpdateVelocitiesKernel, 12, sizeof(cl_mem),&m_vertexData.m_clVertexInverseMass.m_buffer); size_t numWorkItems = workGroupSize*((m_vertexData.getNumVertices() + (workGroupSize-1)) / workGroupSize); if (numWorkItems) diff --git a/src/MiniCL/MiniCL.cpp b/src/MiniCL/MiniCL.cpp index e998f1dfe..24f6751fc 100644 --- a/src/MiniCL/MiniCL.cpp +++ b/src/MiniCL/MiniCL.cpp @@ -646,7 +646,9 @@ extern CL_API_ENTRY cl_int CL_API_CALL clGetContextInfo(cl_context /* co return 0; } -CL_API_ENTRY cl_context CL_API_CALL clCreateContextFromType(cl_context_properties * /* properties */, + + +CL_API_ENTRY cl_context CL_API_CALL clCreateContextFromType(const cl_context_properties * /* properties */, cl_device_type device_type , void (*pfn_notify)(const char *, const void *, size_t, void *) /* pfn_notify */, void * /* user_data */, @@ -706,6 +708,28 @@ CL_API_ENTRY cl_context CL_API_CALL clCreateContextFromType(cl_context_propertie return (cl_context)scheduler; } +CL_API_ENTRY cl_int CL_API_CALL +clGetDeviceIDs(cl_platform_id /* platform */, + cl_device_type /* device_type */, + cl_uint /* num_entries */, + cl_device_id * /* devices */, + cl_uint * /* num_devices */) CL_API_SUFFIX__VERSION_1_0 +{ + return 0; +} + +CL_API_ENTRY cl_context CL_API_CALL +clCreateContext(const cl_context_properties * properties , + cl_uint num_devices , + const cl_device_id * devices , + void (*pfn_notify)(const char *, const void *, size_t, void *), + void * user_data , + cl_int * errcode_ret ) CL_API_SUFFIX__VERSION_1_0 +{ + + return clCreateContextFromType(properties,CL_DEVICE_TYPE_ALL,pfn_notify,user_data,errcode_ret); +} + CL_API_ENTRY cl_int CL_API_CALL clReleaseContext(cl_context context ) CL_API_SUFFIX__VERSION_1_0 { diff --git a/src/MiniCL/cl.h b/src/MiniCL/cl.h index 053491ee2..352829883 100644 --- a/src/MiniCL/cl.h +++ b/src/MiniCL/cl.h @@ -437,7 +437,7 @@ clGetDeviceInfo(cl_device_id /* device */, // Context APIs extern CL_API_ENTRY cl_context CL_API_CALL -clCreateContext(cl_context_properties * /* properties */, +clCreateContext(const cl_context_properties * /* properties */, cl_uint /* num_devices */, const cl_device_id * /* devices */, void (*pfn_notify)(const char *, const void *, size_t, void *) /* pfn_notify */, @@ -445,7 +445,7 @@ clCreateContext(cl_context_properties * /* properties */, cl_int * /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0; extern CL_API_ENTRY cl_context CL_API_CALL -clCreateContextFromType(cl_context_properties * /* properties */, +clCreateContextFromType(const cl_context_properties * /* properties */, cl_device_type /* device_type */, void (*pfn_notify)(const char *, const void *, size_t, void *) /* pfn_notify */, void * /* user_data */,