
ãããè¡ãã«ã¯ãOpenCLããã³RenderScriptãã¯ãããžãŒã«æ³šæãã䟡å€ããããŸãã
ãã®èšäºã§ã¯ãããã°ã©ãã³ã°èšèªOpenCLããã³RenderScriptã䜿çšããé«æ§èœç»ååŠçã®æ¹æ³ã瀺ãAndroidã¢ããªã±ãŒã·ã§ã³ã®äŸã説æããŸãã ãããã®ãã¯ãããžãŒã¯ãããŒã¿ïŒã·ã§ãŒããŒãŠãããïŒã®äžŠååŠççšã«èšèšãããã°ã©ãã£ãã¯ããŒããŠã§ã¢ã®æ©èœã«æ³šç®ããŠèšèšãããŠããŸãã ããã«ããã倧éã®ããŒã¿ã®åŠçãé«éåããå€æ°ã®ã³ãã³ãã®ç¹°ãè¿ãã«äŒŽãåé¡ã解決ã§ããŸãã ä»ã®ãã¯ãããžãŒã䜿çšããŠAndroidã¢ããªã±ãŒã·ã§ã³ã®ã°ã©ãã£ãã¯åŠçãé«éåã§ããŸããããã®èšäºã§ã¯ãã¢ããªã±ãŒã·ã§ã³ã€ã³ãã©ã¹ãã©ã¯ãã£ãæ§ç¯ããOpenCLããã³RenderScriptã§ã°ã©ãã£ã«ã«ã¢ã«ãŽãªãºã ãå®è£ ããäŸã説æããŸãã ãŸããOpenCL APIã®ã©ãããŒã¯ã©ã¹ã«ã€ããŠã説æããŸããããã«ãããã°ã©ãã£ãã¯ã¹ã§åäœããOpenCLã䜿çšããã¢ããªã±ãŒã·ã§ã³ã®äœæãšå®è¡ã容æã«ãªããŸãã ãããžã§ã¯ãã§ãã®ã¯ã©ã¹ã®ãœãŒã¹ã³ãŒãã䜿çšããå Žåãã©ã€ã»ã³ã¹ã¯å¿ èŠãããŸããã
ãã®è³æãæºåããã«ããããèªè ã¯OpenCLããã³RenderScriptãã¯ãããžãŒã«ç²ŸéããŠãããAndroidãã©ãããã©ãŒã åãã®ããã°ã©ãã³ã°ãã¯ããã¯ãç¿åŸããŠããããšãåæãšãªããŸããã ãããã£ãŠãç»åã®åŠçãŸãã¯ããã°ã©ã ã«ããäœæãå éããã¡ã«ããºã ã®æ€èšã«äž»ã«æ³šæãæããŸãã
åäœäŸã確èªããã«ã¯ãOpenCLã³ãŒããå®è¡ã§ããããã«Androidããã€ã¹ãæ§æããå¿ èŠããããŸãã 以äžã§ã¯ãIntel INDEãšAndroid Studioã䜿çšããŠOpenCLéçºçšã®äœæ¥ç°å¢ãç·šæããæ¹æ³ã«ã€ããŠèª¬æããŸãã
ãã®èšäºã®ç®çã¯ãOpenCLããã³RenderScriptã³ãŒãã®æ©èœã瀺ãããšã§ãããããã§ã¯ä»ã®æè¡ã«ã€ããŠã¯èª¬æããŠããŸããã ããã«ãOpenCLã³ãŒããšãããªãããïŒGPUïŒã§å®è¡ãããRenderScriptã䜿çšããã¢ããªã±ãŒã·ã§ã³ã®ããã©ãŒãã³ã¹ã®åæã«é¢ããè³æãèšç»ãããŠããŸãã
1.1ã¢ããªã±ãŒã·ã§ã³ã€ã³ã¿ãŒãã§ã€ã¹
åé¡ã®ã¢ããªã±ãŒã·ã§ã³ã®ç»é¢ã«ã¯3ã€ã®ã¹ã€ããããããRenderScriptãOpenCLããŸãã¯Androidãã€ãã£ãã³ãŒãã䜿çšããŠç»åãæäœããããã®ãµãã·ã¹ãã ãéžæã§ããŸãã ãã®ã¡ãã¥ãŒã䜿çšãããšãCPUïŒäžå€®åŠçè£ çœ®ïŒãŸãã¯GPUïŒã°ã©ãã£ãã¯ã³ã¢ïŒã§ã®OpenCLã³ãŒãã®å®è¡ãåãæ¿ããããšãã§ããŸãã ããã«ãã¡ãã¥ãŒããã°ã©ãã£ãã¯å¹æãéžæã§ããŸãã ã¿ãŒã²ããããã€ã¹ã®éžæã¯ãOpenCLã³ãŒãã§ã®ã¿äœ¿çšã§ããŸãã Intel x86ãã©ãããã©ãŒã ã¯ãCPUãšGPUã®äž¡æ¹ã§OpenCLããµããŒãããŠããŸãã
以äžã«ãOpenCLã«ãã£ãŠçæããããã©ãºãå¹æã衚瀺ããã¢ããªã±ãŒã·ã§ã³ã®ã¡ã€ã³ç»é¢ã瀺ããŸãã

ããã°ã©ã ã®ã¡ã€ã³ãŠã£ã³ããŠãOpenCLã³ãŒãã®å®è¡ã®ããã®ã¿ãŒã²ããããã€ã¹ã®éžæ
ãŠã£ã³ããŠã®å³äžé ã«ãããã©ãŒãã³ã¹ã€ã³ãžã±ãŒã¿ã衚瀺ãããŸãã ãããã¯ãããã°ã©ã ã§ãµããŒããããŠããã°ã©ãã£ãã¯ã¹ãæäœãã3ã€ã®æ¹æ³ãã¹ãŠã§è¡šç€ºãããŸãã
ããã©ãŒãã³ã¹ã¡ããªãã¯ã«ã¯ã1ç§ãããã®ãã¬ãŒã æ°ïŒFPSïŒããã¬ãŒã ã¬ã³ããªã³ã°æéãå¹æã®èšç®ã«å¿ èŠãªæéïŒå¹æã®èšç®çµéæéïŒãå«ãŸããŸãã

ããã©ãŒãã³ã¹ææš
ããã¯ããã©ãŒãã³ã¹ã®äžäŸã«ãããªãããšã«æ³šæããŠãã ããã ããã©ãŒãã³ã¹ã¯ãã³ãŒããå®è¡ãããããã€ã¹ã«ãã£ãŠç°ãªããŸãã
1.2ã 䜿çšãããŠããAPIãšSDK
ADTïŒAndroidéçºããŒã«ãAndroid SDKãå«ãAndroidéçºããŒã«ïŒã«å ããŠããµã³ãã«éçºã§ã¯ãAndroidç°å¢ã§åäœããããã«èšèšãããRenderScript SDKããã³Intel OpenCL SDKã䜿çšããŸããã
Intel OpenCL SDKã¯OpenCLä»æ§ã«åºã¥ããŠããããã®èŠå®ãé å®ããŠããŸãã ãã®ä»æ§ã¯ããªãŒãã³ã§ã¯ãã¹ãã©ãããã©ãŒã ã®éçºæšæºã§ãããç¡æã§äœ¿çšã§ããŸãã 詳现ã¯ã¯ããã¹ã®ãŠã§ããµã€ãã«ãããŸãã
RenderScriptã¯ADT 2.2ã§ç»å ŽããŸããã ïŒAPIã¬ãã«8ïŒã Androidç°å¢ã§ã®é«æ§èœã³ã³ãã¥ãŒãã£ã³ã°ãã©ãããã©ãŒã ã§ãã RenderScriptã¯ãäž»ã«èšç®ã®äžŠåå®è¡ãå¯èœã«ããã¿ã¹ã¯ãå®è¡ããããã«èšèšãããŠããŸãããé£ç¶ããŠå®è¡ãããèšç®ã®æ©æµãåããããšãã§ããŸãã ããã§ã¯ãRenderScriptã®è©³çŽ°ã確èªã§ããŸãã
Google Open Repositoryããå ¥æã§ããææ°ããŒãžã§ã³ã®ADTã«ã¯ãRenderScriptãJNIïŒJavaãã€ãã£ãã€ã³ã¿ãŒãã§ã€ã¹ããã€ãã£ãã³ãŒãä»ãJavaïŒãããã³äžé£ã®ã©ã³ã¿ã€ã APIã䜿çšããããã«ã€ã³ããŒãããå¿ èŠãããããã±ãŒãžãå«ãŸããŠããŸãã
RenderScriptã䜿çšããéçºã®è©³çŽ°ã«ã€ããŠã¯ãOpenCLã®è³æãåç §ããŠãã ãã ã
1.3ã¢ããªã±ãŒã·ã§ã³ãµããŒãã€ã³ãã©ã¹ãã©ã¯ãã£ã³ãŒã
åé¡ã®ã¢ããªã±ãŒã·ã§ã³ã®ã€ã³ãã©ã¹ãã©ã¯ãã£ã¯ãäž»èŠãªã¢ã¯ãã£ããã£ãšè£å©æ©èœã§æ§æãããŠããŸãã ããã§ã¯ããããã®é¢æ°ãšããŠãŒã¶ãŒã€ã³ã¿ãŒãã§ã€ã¹ã®ã«ã¹ã¿ãã€ãºãã°ã©ãã£ãã¯å¹æã®éžæãOpenCLã®äœ¿çšãããã³äœ¿çšããã³ã³ãã¥ãŒãã£ã³ã°ããã€ã¹ã®éžæã«äœ¿çšãããã³ãŒããèŠãŠãããŸãã
ããã§ã¯ã2ã€ã®äž»èŠãªè£å©æ©èœã«ã€ããŠæ€èšããŸãã
æåã®backgroundThreadïŒïŒã¯ãç¬ç«ããå®è¡ã¹ã¬ãããèµ·åããŸãããã®ã¹ã¬ãããããã°ã©ãã£ãã¯å¹æã®æ®µéçãªã¢ããªã±ãŒã·ã§ã³ãå®è¡ããé¢æ°ãå®æçã«åŒã³åºãããŸãã ãã®é¢æ°ã¯ãRenderScriptå ¥éèšäºã§èª¬æãããŠããã¢ããªã±ãŒã·ã§ã³ããååŸãããŸã;ãã®äŸã®è©³çŽ°ã«ã€ããŠã¯ããã¡ããåç §ããŠãã ãã ã
2çªç®ã®é¢æ°processStepïŒïŒã¯ã backgroundThreadïŒïŒããåŒã³åºãããŸãã 圌女㯠ãé çªã«ãç»ååŠçã®ã³ãã³ããåŒã³åºããŸãã ãã®é¢æ°ã¯ãã¹ã€ããã®ç¶æ ã«åºã¥ããŠã䜿çšããã¢ã«ãŽãªãºã ã®å®è£ ã決å®ããŸãã é¢æ°processStepïŒïŒã§ã¯ãèšå®ã®ç¶æ ã«å¿ããŠãOpenCLãRenderScriptããŸãã¯C / C ++ã§èšè¿°ãããéåžžã®ãã·ã³ã³ãŒãã䜿çšããŠç»åãåŠçããããã®ã¡ãœãããåŒã³åºãããŸãã ãã®ã³ãŒãã¯ããã¯ã°ã©ãŠã³ãã¹ã¬ããã§å®è¡ãããããããŠãŒã¶ãŒã€ã³ã¿ãŒãã§ã€ã¹ã¯ãããã¯ãããªãããããã€ã§ãã°ã©ãã£ã«ã«ã¢ã«ãŽãªãºã ã®å®è£ ãåãæ¿ããããšãã§ããŸãã ãŠãŒã¶ãŒããšãã§ã¯ããéžæãããšãã¢ããªã±ãŒã·ã§ã³ã¯ããã«ãã®ãšãã§ã¯ãã«åãæ¿ãããŸãã
// processStep() () . private void processStep() { try { switch (this.type.getCheckedRadioButtonId()) { case R.id.type_renderN: oclFlag = 0; // OpenCL stepRenderNative(); break; case R.id.type_renderOCL: oclFlag = 1; // OpenCL stepRenderOpenCL(); break; case R.id.type_renderRS: oclFlag = 0; // OpenCL stepRenderScript(); break; default: return; } } catch (RuntimeException ex) { // Log.wtf("Android Image Processing", "render failed", ex); } }
1.4 Javaã®ãã·ã³ã³ãŒãã«å®è£ ãããé¢æ°ã®å®çŸ©
åé¡ã®ã¢ããªã±ãŒã·ã§ã³ã¯ãã°ã©ãã£ã«ã«å¹æãå®è£ ããJNIã䜿ââçšããŠãã·ã³ã³ãŒãã¬ãã«ã®ã³ãã³ããåŒã³åºãããã«äœ¿çšãããé¢æ°ãå®çŸ©ããNativeLibã¯ã©ã¹ãå®è£ ããŸãã ã¢ããªã±ãŒã·ã§ã³ã«ã¯ããã©ãºãå¹æïŒãã©ãºãïŒãã»ãã¢ïŒã»ãã¢ïŒã®ç»åã®èª¿è²ãããã³å€è²ïŒã¢ãã¯ãïŒã®3ã€ã®å¹æããããŸãã ãããã£ãŠãé¢æ°ã¯renderPlasmaïŒ...ïŒ ã renderSepiaïŒ...ïŒããã³renderMonoChromeïŒ...ïŒãå®çŸ©ããŸãã ã ãããã®Javaé¢æ°ã¯ãJNIãšã³ããªãã€ã³ãã®åœ¹å²ãæãããŸããJNIãšã³ããªãã€ã³ããä»ããŠããã·ã³ã³ãŒãã«å®è£ ãããæ©èœãåŒã³åºãããããã°ã©ãã£ã«ã«ã¢ã«ãŽãªãºã ã®OpenCLããŒãžã§ã³ãåŒã³åºãããŸãã
察å¿ããJNIé¢æ°ã¯ãéžæããã°ã©ãã£ãã¯å¹æãéå§ãããšãã«ãC / C ++ã§èšè¿°ãããã³ãŒããå®è¡ããããOpenCLã§ããã°ã©ã ãæ§æããŠå®è¡ããŸãã èšè¿°ãããã¯ã©ã¹ã¯ãããã±ãŒãžandroid.graphics.Bitmapããã³android.content.res.AssetManagerã䜿çšããŸãã BitMapãªããžã§ã¯ã㯠ãåŠçã®ããã«ã°ã©ãã£ãã¯ããŒã¿ããµãã·ã¹ãã ã«éä¿¡ããçµæãååŸããããã«äœ¿çšãããŸãã ã¢ããªã±ãŒã·ã§ã³ã¯ã AssetManagerã¯ã©ã¹ã®ãªããžã§ã¯ãã䜿çšããŠãOpenCLãã¡ã€ã«ïŒããšãã°ãsepia.clïŒã«ã¢ã¯ã»ã¹ããŸãã ãããã®ãã¡ã€ã«ã¯ãã°ã©ãã£ã«ã«ã¢ã«ãŽãªãºã ãå®è£ ããé¢æ°ã§ããOpenCLã«ãŒãã«ãèšè¿°ããŠããŸãã
以äžã¯ã NativeLibã¯ã©ã¹ã®ã³ãŒãã§ãã // TODOã§ã³ã¡ã³ããããŠããããã«ãç°¡åã«æ¡åŒµããŠã°ã©ãã£ãã¯å¹æãè¿œå ã§ããŸãã
package com.example.imageprocessingoffload; import android.content.res.AssetManager; import android.graphics.Bitmap; public class NativeLib { // libimageeffects.so public static native void renderPlasma(Bitmap bitmapIn, int renderocl, long time_ms, String eName, int devtype, AssetManager mgr); public static native void renderMonoChrome(Bitmap bitmapIn, Bitmap bitmapOut, int renderocl, long time_ms, String eName, int simXtouch, int simYtouch, int radHi, int radLo, int devtype, AssetManager mgr); public static native void renderSepia(Bitmap bitmapIn, Bitmap bitmapOut, int renderocl, long time_ms, String eName, int simXtouch, int simYtouch, int radHi, int radLo, int devtype, AssetManager mgr); //TODO public static native <return type> render<Effectname>(âŠ); // static { System.loadLibrary("imageeffectsoffloading"); } }
Android AssetManagerããã³BitMapãªããžã§ã¯ãã¯ãå ¥åããã³åºåãã©ã¡ãŒã¿ãŒãšããŠãã·ã³ã³ãŒãã«æž¡ãããããšã«æ³šæããŠãã ããã AssetManagerãªããžã§ã¯ã㯠ãOpenCLã«ãŒãã«ãèšè¿°ããCLãã¡ã€ã«ã«ã¢ã¯ã»ã¹ããããã«ãã·ã³ã³ãŒãã«ãã£ãŠäœ¿çšãããŸãã BitMapãªããžã§ã¯ãã«ã¯ããã·ã³ã³ãŒãã§åŠçããããã¯ã»ã«ããŒã¿ãå«ãŸããŠããŸãã åŠçã®çµæãè¿ãããã«åãããŒã¿åã䜿çšãããŸãã
deviceTypeãŠãŒã¶ãŒã€ã³ã¿ãŒãã§ã€ã¹ãã©ã¡ãŒã¿ãŒã¯ãOpenCLã³ãŒããå®è¡ãããã¿ãŒã²ããããã€ã¹ïŒCPUãŸãã¯GPUïŒã瀺ãããã«äœ¿çšãããŸãã ãã®æè»æ§ãå®çŸããã«ã¯ãããã«å¿ããŠAndroid OSãæ§æããå¿ èŠããããŸãã ææ°ã®Intel Atomããã³Intel Coreããã»ããµã¯ãOpenCLåœä»€ãç¬ç«ããŠå®è¡ã§ããã·ã¹ãã ã®çµ±åã°ã©ãã£ãã¯ãããã䜿çšããŸãã
eNameãã©ã¡ãŒã¿ãŒã¯ãã©ã®OpenCLã«ãŒãã«ãã³ã³ãã€ã«ããŠå®è¡ããããæå®ããŸãã ã¢ããªã±ãŒã·ã§ã³ã§ã¯ãåã°ã©ãã£ãã¯å¹æã«ç¬èªã®JNIé¢æ°ããããŸãããã®çµæãã«ãŒãã«åã®è»¢éã¯äžèŠã«æãããããããŸããã ãã ãã1ã€ã®CLãã¡ã€ã«ïŒåãããšãJNIæ©èœã«ãåœãŠã¯ãŸããŸãïŒã§ã¯ãããã€ãã®åæ§ã®ã°ã©ãã£ã«ã«ã¢ã«ãŽãªãºã ããšã³ã³ãŒãã§ããŸãã ãã®ãããªç¶æ³ã§ã¯ã eNameãã©ã¡ãŒã¿ãŒã䜿çšããŠãã©ã®ç¹å®ã®CLããã°ã©ã ïŒãŸãã¯ã«ãŒãã«ïŒãã³ã³ãã€ã«ããã³ããŒãããå¿ èŠããããã瀺ãããšãã§ããŸãã
renderoclãã©ã¡ãŒã¿ãŒã¯ãOpenCLã³ãŒããŸãã¯C / C ++ã§èšè¿°ããããã·ã³ã³ãŒããå®è¡ãããã©ããã瀺ããã©ã°ãšããŠæ©èœããŸãã ãŠãŒã¶ãŒãOpenCLã¹ã€ãããã¢ã¯ãã£ãã«ããå Žåããã®å€ã¯trueãšè§£éãããŸãããã以å€ã®å Žåããã©ã°ã¯æªèšå®ã®ãŸãŸã§ãã
time_msãã©ã¡ãŒã¿ãŒã¯ ãããã©ãŒãã³ã¹ã€ã³ãžã±ãŒã¿ãŒã®èšç®ã«äœ¿çšãããã¿ã€ã ã¹ã¿ã³ãïŒããªç§åäœïŒãæž¡ãããã«äœ¿çšãããŸãã ããã«ããã©ãºãå¹æã¯ãç»åã®æ®µéçãªèšç®äžã«ãã®å€ã«çŠç¹ãåãããŸãã
ä»ã®åŒæ°ã¯ãã°ã©ãã£ãã¯å¹æã®å®è£ ã®æ©èœãç¹ã«ãæ²»çé åã®æŸå°ç¶ã®æ¡å€§ãåæ ããŠããŸãã ããšãã°ããã©ã¡ãŒã¿ãŒsimXtouch ã simYTouch ã radLo ãããã³radHi㯠ãå¹ ãšé«ãã®å€ãšãšãã«ãç»åãã»ãã¢ãšå€è²ã«èª¿æŽããå¹æã§äœ¿çšãããŸãã ãããã䜿çšããŠãç¹å®ã®ãã€ã³ãã§ç»ååŠçãéå§ãããæ¹æ³ã®èšç®ãšè¡šç€ºãè¡ããããã®åŸãé åå šäœãç»åå šäœã«åºãããŸã§æŸå°ç¶ã«æ¡å€§ããŸãã
1.5ãã·ã³ã³ãŒãã®å®è¡ã«å¿ èŠãªå®çŸ©ãšãªãœãŒã¹ïŒCãŸãã¯OpenCLïŒ
ããã§ã¯ãäŸã«ç€ºãããŠããå¹æãå®è£ ãããã·ã³JNIé¢æ°ã®å®çŸ©ãèŠãŠãããŸãã æ¢ã«è¿°ã¹ãããã«ãåé¢æ°ã«ã¯1ã€ã®é¢æ°ããããŸãã ããã¯ã話ãè€éã«ãããOpenCLã䜿çšããŠç»ååŠçãå éããããã«äœ¿çšãããæ©èœèŠçŽ ãããæ確ã«åŒ·èª¿ããããã«è¡ãããŸãã
Cã§èšè¿°ãããã³ãŒããžã®ãªã³ã¯ãããããã®ã³ãŒãã®ãã©ã°ã¡ã³ããããã«å«ãŸããŠããŸãã ããã¯ãäŸã®å°æ¥ã®ããŒãžã§ã³ã§æ€èšäžã®ãã¯ãããžãŒã®ããã©ãŒãã³ã¹ã®æ¯èŒã«åºã¥ããŠè¡ãããŸãã
1ã€ã®JNIé¢æ°ã¯ã1ã€ã®ãã·ã³Javaé¢æ°ã«å¯Ÿå¿ããŸãã ãããã£ãŠãJNIé¢æ°ãæ£ãã宣èšããã³å®çŸ©ããããšãéåžžã«éèŠã§ãã Java SDKã«ã¯ãæ£ç¢ºã§æ£ç¢ºãªJNIé¢æ°å®£èšã®çæã«åœ¹ç«ã€javahããŒã«ããããŸã ã ãã®ããŒã«ã¯ãã³ãŒããæ£ããã³ã³ãã€ã«ãããŠãå®è¡æãšã©ãŒãçºçããå°é£ãªç¶æ³ã«é¥ããªãããã«ããããã«äœ¿çšããããšããå§ãããŸãã
以äžã¯ãç»ååŠçãé«éåããããã®äœã¬ãã«é¢æ°ã«å¯Ÿå¿ããJNIé¢æ°ã§ãã é¢æ°ã·ã°ããã£ã¯ãjavahããŒã«ã䜿çšããŠçæãããŸãã
// JNI-, // #ifndef _Included_com_example_imageprocessingoffload_NativeLib #define _Included_com_example_imageprocessingoffload_NativeLib #ifdef __cplusplus extern "C" { #endif /* * Class: com_example_imageprocessingoffload_NativeLib * Method: renderPlasma * Signature: (Landroid/graphics/Bitmap;IJLjava/lang/String;)Ljava/lang/String; */ JNIEXPORT void JNICALL Java_com_example_imageprocessingoffload_NativeLib_renderPlasma (JNIEnv *, jclass, jobject, jint, jlong, jstring, jint, jobject); /* * Class: com_example_imageprocessingoffload_NativeLib * Method: renderMonoChrome * Signature: (Landroid/graphics/Bitmap;Landroid/graphics/Bitmap;IJLjava/lang/String;)Ljava/lang/String; */ JNIEXPORT void JNICALL Java_com_example_imageprocessingoffload_NativeLib_renderMonoChrome (JNIEnv *, jclass, jobject, jobject, jint, jlong, jstring, jint, jint, jint, jint, jint, jobject); /* * Class: com_example_imageprocessingoffload_NativeLib * Method: renderSepia * Signature: (Landroid/graphics/Bitmap;Landroid/graphics/Bitmap;IJLjava/lang/String;)Ljava/lang/String; */ JNIEXPORT void JNICALL Java_com_example_imageprocessingoffload_NativeLib_renderSepia (JNIEnv *, jclass, jobject, jobject, jint, jlong, jstring, jint, jint, jint, jint, jint, jobject); } #endif
javahããŒã«ã¯ãJNIé¢æ°ã®æ£ãã眲åãçæã§ããŸãã ãã ããJavaãã·ã³é¢æ°ãå®çŸ©ãã1ã€ãŸãã¯è€æ°ã®ã¯ã©ã¹ã¯ãAndroidãããžã§ã¯ãã«æ¢ã«ã³ã³ãã€ã«ãããŠããå¿ èŠããããŸãã ããããŒãã¡ã€ã«ãçæããå¿ èŠãããå Žåã¯ã次ã®ããã«javahã³ãã³ãã䜿çšã§ããŸãã
{javahLocation} -o {outputFile} -classpath {classpath} {importName}
ãã®äŸã§ã¯ãé¢æ°ã·ã°ããã£ã¯æ¬¡ã®ã³ãã³ãã§äœæãããŸããã
javah -o junk.h -classpath bin \ classes com.example.imageprocessingoffloading.NativeLib
次ã«ã junk.hãã¡ã€ã«ã®JNIé¢æ°ã®çœ²åãimageeffects.cppãã¡ã€ã«ã«è¿œå ãããOpenCLãŸãã¯Cã³ãŒãã®æºåãšå®è¡ãå®è£ ãããŸãã 次ã«ãOpenCLã®å®è¡ã«å¿ èŠãªãªãœãŒã¹ããŸãã¯ãã©ãºãå¹æãå€è²ãã»ãã¢è²åãã®ãã·ã³ã³ãŒããå²ãåœãŠãŸãã
1.5.1ãã©ãºãå¹æ
Java_com_example_imageprocessingoffload_NativeLib_renderPlasmaïŒ...ïŒé¢æ°ã¯ããã©ãºãå¹æãå®è£ ããOpenCLãŸãã¯ãã·ã³ã³ãŒããå®è¡ããããã®ãšã³ããªãã€ã³ãã§ãã é¢æ°startPlasmaOpenCLïŒ...ïŒ ã runPlasmaOpenCLïŒ...ïŒ ãããã³runPlasmaNativeïŒ...ïŒã¯ã imageeffects.cppãã¡ã€ã«ã®ã³ãŒãã®å€éšã«ãããå¥ã®plasmaEffect.cppãã¡ã€ã«ã§å®£èšãããŠããŸãã plasmaEffect.cppã®ãœãŒã¹ã³ãŒãã¯ã ããã«ãããŸã ã
ãšã³ããªãã€ã³ãã§ããrenderPlasmaïŒ...ïŒé¢æ°ã¯ãOpenCLã©ãããŒã䜿çšããŠãOpenCLã®AndroidãµããŒããèŠæ±ããŸãã ã©ãããŒã¯ã©ã¹é¢æ°:: initOpenCLïŒ...ïŒãåŒã³åºããŠãOpenCLç°å¢ãåæåããŸãã ããã€ã¹ã¿ã€ããšããŠãOpenCLã³ã³ããã¹ããäœæãããšãã«ãCPUãŸãã¯GPUã転éãããŸãã AndroidãªãœãŒã¹ãããŒãžã£ãŒã¯ceNameãã©ã¡ãŒã¿ãŒã䜿çšããŠãå¿ èŠãªã«ãŒãã«ãå«ãCLãã¡ã€ã«ãèå¥ãããŠã³ããŒããã³ã³ãã€ã«ããŸãã
OpenCLç°å¢ãæ£åžžã«æ§æã§ããå Žåã renderPlasmaïŒ...ïŒé¢æ°ã®æ¬¡ã®ã¹ãããã¯ã startPlasmaOpenCLïŒïŒé¢æ°ãåŒã³åºãããšã§ã ããã®é¢æ°ã¯ãOpenCLãªãœãŒã¹ãå²ãåœãŠããã©ãºãå¹æãå®è£ ããã«ãŒãã«ã®å®è¡ãéå§ããŸãã gOCLã¯ãOpenCLã©ãããŒã¯ã©ã¹ã®ã€ã³ã¹ã¿ã³ã¹ãæ ŒçŽããã°ããŒãã«å€æ°ã§ããããšã«æ³šæããŠãã ããã ãã®å€æ°ã¯ããšã³ããªãã€ã³ãã§ãããã¹ãŠã®JNIé¢æ°ã«è¡šç€ºãããŸãã ãã®ã¢ãããŒãã®ãããã§ããµããŒããããŠããã°ã©ãã£ãã¯å¹æã«ã¢ã¯ã»ã¹ãããšãã«OpenCLç°å¢ãåæåã§ããŸãã
ãã©ãºãå¹æãå®èšŒããå Žåãæ¢è£œã®ç»åã¯äœ¿çšãããŸããã ç»é¢ã«è¡šç€ºããããã®ã¯ãã¹ãŠãããã°ã©ã ã«ãã£ãŠçæãããŸãã bitmapInãã©ã¡ãŒã¿ãŒã¯ãã¢ã«ãŽãªãºã ã®æäœäžã«çæãããã°ã©ãã£ãã¯ããŒã¿ãæ ŒçŽããBitMapã¯ã©ã¹ã®ãªããžã§ã¯ãã§ãã startPlasmaïŒ...ïŒé¢æ°ã«æž¡ããããã¯ã»ã«ãã©ã¡ãŒã¿ãŒã¯ãã©ã¹ã¿ãŒãã¯ã¹ãã£ã«è¡šç€ºããããã·ã³ã³ãŒããŸãã¯OpenCLã«ãŒãã«ã³ãŒãã«ãã£ãŠäœ¿çšãããç»é¢ã«è¡šç€ºããããã¯ã»ã«ããŒã¿ãèªã¿æžãããŸãã ããäžåºŠã assetManagerãªããžã§ã¯ãã䜿çšããŠããã©ãºãå¹æãå®è£ ããOpenCLã«ãŒãã«ãå«ãCLãã¡ã€ã«ã«ã¢ã¯ã»ã¹ããŸãã
JNIEXPORT void Java_com_example_imageprocessingoffload_NativeLib_renderPlasma(JNIEnv * env, jclass, jobject bitmapIn, jint renderocl, jlong time_ms, jstring ename, jint devtype, jobject assetManager) { ⊠// // BitMapIn âpixelsâ, OpenCL- . ret = AndroidBitmap_lockPixels(env, bitmapIn, &pixels); ⊠// If OCL not initialized AAssetManager *amgr = AAssetManager_fromJava(env, assetManager); gOCL.initOpenCL(clDeviceType, ceName, amgr); startPlasmaOpenCL((cl_ushort *) pixels, infoIn.height, infoIn.width, (float) time_ms, ceName, cpinit); else runPlasmaOpenCL(infoIn.width, infoIn.height, (float) time_ms, (cl_ushort *) pixels); ⊠// }
å€éšã®startPlasmaOpenCLïŒ...ïŒé¢æ°ã¯ã Paletteãããã¡ãŒãšAnglesãããã¡ãŒãäœæããŠå¡ãã€ã¶ããŸãããããã®ãããã¡ãŒã«ã¯ããã©ãºãå¹æã®äœæã«å¿ èŠãªããŒã¿ãå«ãŸããŠããŸãã ãã®å¹æã®åå ãšãªãOpenCLã«ãŒãã«ãå®è¡ããããã«ãé¢æ°ã¯ãã©ãããŒã¯ã©ã¹ã®ã¡ã³ããŒããŒã¿ãšããŠå®çŸ©ãããŠããOpenCLã³ãã³ããã¥ãŒãã³ã³ããã¹ããããã³ã«ãŒãã«ã«äŸåããŠããŸãã
runPlasmaOpenCLïŒ...ïŒé¢æ°ã¯ããã©ãºãç»åãçæããOpenCLã«ãŒãã«ãç¶ç¶çã«åŒã³åºããŸãã OpenCLã«ãŒãã«ã®èµ·åæã«å¥ã®é¢æ°ã1å䜿çšãããŸãããã®åŸã®ã«ãŒãã«åŒã³åºãã§ã¯ãå ¥åã«æ°ããã¿ã€ã ã¹ã¿ã³ãå€ã®ã¿ãå¿ èŠã§ãã åŸç¶ã®ã«ãŒãã«åŒã³åºãã§ã¯ãã¿ã€ã ã¹ã¿ã³ãã®åœ¢åŒã§åŒæ°ãæž¡ãã ãã§ãããããè¿œå ã®é¢æ°ãå¿ èŠã§ãã
extern int startPlasmaOpenCL(cl_ushort* pixels, cl_int height, cl_int width, cl_float ts, const char* eName, int inittbl); extern int runPlasmaOpenCL(int width, int height, cl_float ts, cl_ushort *pixels); extern void runPlasmaNative( AndroidBitmapInfo* info, void* pixels, double t, int inittbl );
runPlasmaNativeïŒ...ïŒé¢æ°ã«ã¯ãCã«ãã©ãºãå¹æãäœæããããã®ã¢ã«ãŽãªãºã ã®å®è£ ãå«ãŸããŠããŸããinittblåŒæ°ã¯è«çåŒæ°ãšããŠäœ¿çšãããŸãããã®å€ã¯ãã¢ã«ãŽãªãºã ãæ©èœããããã«å¿ èŠãªPaletteããã³AnglesããŒã¿ã»ãããå¿ èŠãã©ããã瀺ããŸãã ãã©ãºãå¹æãå®è£ ããOpenCLã«ãŒãã«ã³ãŒãã¯ã plasmaEffect.cppãã¡ã€ã«ã«ãããŸãã
#define FBITS 16 #define FONE (1 << FBITS) #define FFRAC(x) ((x) & ((1 << FBITS)-1)) #define FIXED_FROM_FLOAT(x) ((int)((x)*FONE)) /* , */ #define PBITS 8 #define ABITS 9 #define PSIZE (1 << PBITS) #define ANGLE_2PI (1 << ABITS) #define ANGLE_MSK (ANGLE_2PI - 1) #define YT1_INCR FIXED_FROM_FLOAT(1/100.0f) #define YT2_INCR FIXED_FROM_FLOAT(1/163.0f) #define XT1_INCR FIXED_FROM_FLOAT(1/173.0f) #define XT2_INCR FIXED_FROM_FLOAT(1/242.0f) #define ANGLE_FROM_FIXED(x) ((x) >> (FBITS - ABITS)) & ANGLE_MSK ushort pfrom_fixed(int x, __global ushort *palette) { if (x < 0) x = -x; if (x >= FONE) x = FONE-1; int idx = FFRAC(x) >> (FBITS - PBITS); return palette[idx & (PSIZE-1)]; } __kernel void plasma(__global ushort *pixels, int height, int width, float t, __global ushort *palette, __global int *angleLut) { int yt1 = FIXED_FROM_FLOAT(t/1230.0f); int yt2 = yt1; int xt10 = FIXED_FROM_FLOAT(t/3000.0f); int xt20 = xt10; int x = get_global_id(0); int y = get_global_id(1); int tid = x+y*width; yt1 += y*YT1_INCR; yt2 += y*YT2_INCR; int base = angleLut[ANGLE_FROM_FIXED(yt1)] + angleLut[ANGLE_FROM_FIXED(yt2)]; int xt1 = xt10; int xt2 = xt20; xt1 += x*XT1_INCR; xt2 += x*XT2_INCR; int ii = base + angleLut[ANGLE_FROM_FIXED(xt1)] + angleLut[ANGLE_FROM_FIXED(xt2)]; pixels[tid] = pfrom_fixed(ii/4, palette); }
1.5.2ç»åã®å€è²
Java_com_example_imageprocessingoffload_NativeLib_renderMonochromeïŒ...ïŒé¢æ°ã¯ããã·ã³ã³ãŒããŸãã¯OpenCLããŒã«ã䜿çšããŠå®è£ ãããç»åæŒçœæ©èœãåŒã³åºãããã®ãšã³ããªãã€ã³ãã§ãã é¢æ°executeMonochromeOpenCLïŒ...ïŒããã³executeMonochromeNativeïŒ...ïŒã¯ã imageeffects.cppã®ã³ãŒãã®å€éšã«ãããå¥ã®ãã¡ã€ã«ã§å®£èšãããŠããŸãã ãã©ãºãå¹æã®å Žåãšåæ§ã«ããšã³ããªãã€ã³ããšããŠæ©èœããé¢æ°ã¯ãOpenCLã©ãããŒã䜿çšããŠãAndroidããã€ã¹ç®¡çãµãã·ã¹ãã ãžã®OpenCLãµããŒãã«é¢é£ããèŠæ±ãå®è¡ããŸãã ãŸããOpenCLç°å¢ãåæåãã:: initOpenCLïŒ...ïŒé¢æ°ãåŒã³åºããŸãã
次ã®ã³ãŒãè¡ã¯ã executeMonochromeOpenCLïŒ...ïŒããã³executeMonochromeNativeïŒ...ïŒé¢æ°ãexternããŒã¯ãŒãã§å®£èšãããŠããããšã瀺ããŠããŸãã ããã«ããããããã¯NDKã³ã³ãã€ã©ããèŠããããã«ãªããŸãã ãããã®é¢æ°ã¯å¥ã®ãã¡ã€ã«ã§å®£èšãããŠããããããããå¿ èŠã§ãã
extern int executeMonochromeOpenCL(cl_uchar4 *srcImage, cl_uchar4 *dstImage, int radiHi, int radiLo, int xt, int yt, int nWidth, int nHeight); extern int executeMonochromeNative(cl_uchar4 *srcImage, cl_uchar4 *dstImage, int radiHi, int radiLo, int xt, int yt, int nWidth, int nHeight);
ãã©ãºãå¹æãšã¯ç°ãªããããã§ã¯å ¥åç»åãšåºåç»åã䜿çšãããŸãã bitmapInãšbitmapOutã¯äž¡æ¹ãšãARGB_888圢åŒã®ããããããç»åã§ãã ã©ã¡ããcl_uchar4ãªã©ã®ãã¯ã¿ãŒã®CLãããã¡ãŒã«ããããããŸãã ããã§ã¯ pixelsInãšpixelsOutã®åå€æ ãå®è¡ãããããšã«æ³šæããŠãã ãããããã¯OpenCLãBitMapãªããžã§ã¯ããcl_uchar4ãã¯ãã«ãããã¡ã«ãããã§ããããã«ããããã«å¿ èŠã§ãã
JNIEXPORT void JNICALL Java_com_example_imageprocessingoffload_NativeLib_renderMonochrome(JNIEnv * env, jclass obj, jobject bitmapIn, jobject bitmapOut, jint renderocl, jlong time_ms, jstring ename, jint xto, jint yto, jint radHi, jint radLo, jint devtype, jobject assetManager) { ⊠// // BitMapIn âpixelsInâ, OpenCL- . ret = AndroidBitmap_lockPixels(env, bitmapIn, &pixelsIn); // BitMapOut âpixelsOutâ, OpenCL- ret = AndroidBitmap_lockPixels(env, bitmapOut, &pixelsOut); ⊠// If OpenCL If OCL not initialized AAssetManager *amgr = AAssetManager_fromJava(env, assetManager); gOCL.initOpenCL(clDeviceType, ceName, amgr); else executeMonochromeOpenCL((cl_uchar4*) pixelsIn,(cl_uchar4*) pixelsOut, radiHi, radiLo, xt, yt, infoIn.width, infoIn.height); // , , OCL else executeMonochromeNative((cl_uchar4*) pixelsIn,(cl_uchar4*) pixelsOut, radiHi, radiLo, xt, yt, infoIn.width, infoIn.height); // OpenCL ⊠// }
executeMonochromeOpenCLïŒ...ïŒé¢æ°ãåŒã³åºããããšã pixelsInãšpixelsOut㯠cl_uchar4ãããã¡ãŒã®ã¿ã€ãã«å€æãããŠè»¢éãããŸãã ãã®é¢æ°ã¯ãOpenCL APIã䜿çšããŠãäœæ¥ã«å¿ èŠãªãããã¡ãŒããã³ãã®ä»ã®ãªãœãŒã¹ãäœæããŸãã ã«ãŒãã«åŒæ°ãèšå®ããOpenCLã«ãŒãã«ã®å®è¡ã«å¿ èŠãªã³ãã³ãããã¥ãŒã«å ¥ããŸãã å ¥åã€ã¡ãŒãžãããã¡ãŒã¯èªã¿åãå°çšãããã¡ãŒïŒread_onlyïŒãšããŠåœ¢æããã pixelsInãã€ã³ã¿ãŒã䜿çšããŠã¢ã¯ã»ã¹ããŸãã ã«ãŒãã«ã³ãŒãã¯ãã®ãã€ã³ã¿ãŒã䜿çšããŠãç»åã®å ¥åãã¯ã»ã«ããŒã¿ãååŸããŸãã ãã®ããŒã¿ã¯ãã«ãŒãã«ã«ãã£ãŠåŠçãããå ¥åç»åã¯å€è²ããŸãã åºåãããã¡ãŒã¯ãèªã¿åããšæžã蟌ã¿ïŒread_writeïŒã®äž¡æ¹ã®ããã«èšèšããããããã¡ãŒã§ãç»ååŠçã®çµæãä¿åãã pixelsOutããããæããŸãã OpenCLã®è©³çŽ°ã«ã€ããŠã¯ã Intelã®ããã°ã©ãã³ã°ããã³æé©åã¬ã€ããåç §ããŠãã ããã
executeMonochromeNativeïŒ...ïŒé¢æ°ã§ã¯ãç»åæŒçœã¢ã«ãŽãªãºã ã¯Cã§å®è£ ãããŸããããã¯ããã¹ãïŒã«ãŒãïŒã䜿çšããéåžžã«åçŽãªã¢ã«ãŽãªãºã ã§ããå€éšïŒyïŒããã³å éšïŒxïŒã pixelInãæãsrcImageå€æ°ã¯ãã¢ã«ãŽãªãºã åŒã§å ¥åãã¯ã»ã«ããŒã¿ãååŸããããã«äœ¿çšãããã«ã©ãŒã€ã¡ãŒãžãã¢ãã¯ãã«å€æãããŸãã
æŒçœå¹æãå®è£ ããOpenCLã«ãŒãã«ã³ãŒãã¯æ¬¡ã®ãšããã§ãã
constant uchar4 cWhite = {1.0f, 1.0f, 1.0f, 1.0f}; constant float3 channelWeights = {0.299f, 0.587f, 0.114f}; constant float saturationValue = 0.0f; __kernel void mono (__global uchar4 *in, __global uchar4 *out, int4 intArgs, int width) { int x = get_global_id(0); int y = get_global_id(1); int xToApply = intArgs.x; int yToApply = intArgs.y; int radiusHi = intArgs.z; int radiusLo = intArgs.w; int tid = x + y * width; uchar4 c4 = in[tid]; float4 f4 = convert_float4 (c4); int xRel = x - xToApply; int yRel = y - yToApply; int polar = xRel*xRel + yRel*yRel; if (polar > radiusHi || polar < radiusLo) { if (polar < radiusLo) { float4 outPixel = dot (f4.xyz, channelWeights); outPixel = mix ( outPixel, f4, saturationValue); outPixel.w = f4.w; out[tid] = convert_uchar4_sat_rte (outPixel); } else { out[tid] = convert_uchar4_sat_rte (f4); } } else { out[tid] = convert_uchar4_sat_rte (cWhite); } }
1.5.3ã»ãã¢èª¿
ã»ãã¢è²ã®ç»åã«è²ãä»ããã³ãŒãã¯ãããªãŒãã¢ã«ãŽãªãºã ã®å®è£ ã«éåžžã«äŒŒãŠããŸãã äž»ãªéãã¯ããã¯ã»ã«ã®è²æ å ±ã®åŠçæ¹æ³ã§ãã ããã§ã¯ãä»ã®åŒãšå®æ°ã䜿çšãããŸãã 以äžã¯ãOpenCLããã³Cã«ãã£ãŠå®è£ ãããã¢ã«ãŽãªãºã ã®å®è£ ãåŒã³åºãããã®é¢æ°ã®å®£èšã§ããã芧ã®ãšãããååãé€ãé¢æ°ã¯ãæŒçœã¢ã«ãŽãªãºã ã®å®è£ ãåŒã³åºãããã®é¢æ°ãšåãã«èŠããŸãã
extern int executeSepiaOpenCL(cl_uchar4 *srcImage, cl_uchar4 *dstImage, it int radiHi, int radiLo, int xt, int yt, int nWidth, int nHeight); extern int executeSepiaNative(cl_uchar4 *srcImage, cl_uchar4 *dstImage, int radiHi, int radiLo, int xt, int yt, int nWidth, int nHeight); JNIEXPORT jstring JNICALL Java_com_example_imageprocessingoffload_NativeLib_renderSepia(JNIEnv * env, jclass obj, jobject bitmapIn, jobject bitmapOut, jint renderocl, jlong time_ms, jstring ename, jint xto, jint yto, jint radHi, jint radLo, jint devtype, jobject assetManager) { ⊠}
Java_com_example_imageprocessingoffload_NativeLib_renderSepiaïŒ...ïŒã®ã³ãŒãããããªãŒãã¢ã«ãŽãªãºã ã§èŠããã®ãšéåžžã«äŒŒãŠãããããããã§ã¯ç€ºããŸããã
executeSepiaOpenCLïŒ...ïŒé¢æ°ãåŒã³åºããããš ãæž¡ãããå€ãç®çã®ã¿ã€ãã«å€æãã cl_uchar4ãããã¡ãŒã®åœ¢åŒã§ã pixelsInãšpixelsOutãæž¡ããŸãã OpenCL APIã䜿çšããŠããããã¡ãŒããã®ä»ã®å¿ èŠãªãªãœãŒã¹ãäœæããŸãã ãŸããOpenCLã«ãŒãã«ã®åŒæ°ãèšå®ããã³ãã³ããå®è¡åŸ ã¡ã«ããŸãã å ¥åã€ã¡ãŒãžãããã¡ãŒã¯èªã¿åãå°çšãããã¡ãŒïŒread_onlyïŒãšããŠåœ¢æããã pixelsInãã€ã³ã¿ãŒã䜿çšããŠã¢ã¯ã»ã¹ããŸãã ã«ãŒãã«ã³ãŒãã¯ããã€ã³ã¿ãŒã䜿çšããŠãã¯ã»ã«ç»åããŒã¿ãååŸããŸãã 次ã«ããã®ããŒã¿ã¯ã«ãŒãã«ã«ãã£ãŠåŠçãããå ¥åç»åã¯ã»ãã¢è²ã«ãªããŸãã åºåãããã¡ãŒã¯ãèªã¿åããšæžã蟌ã¿ïŒread_writeïŒã®äž¡æ¹ã®ããã«èšèšããããããã¡ãŒã§ãç»ååŠçã®çµæãä¿åãã pixelsOutããããæããŸãã
executeSepiaNativeïŒ...ïŒé¢æ°ã«ã¯ã Cã®ã»ãã¢èª¿è²ã¢ã«ãŽãªãºã ã®å®è£ ãå«ãŸããŠããŸããããã¯ãå€éšïŒyïŒããã³å éšïŒxïŒã®ãã¹ããããã«ãŒãã®ãã¢ã§æ§æãããåçŽãªã¢ã«ãŽãªãºã ã§ãã ããŒã¿åŠçã¯ãµã€ã¯ã«ã§å®è¡ãããçµæã¯ã pixelsOutã§ç€ºãããdstImageå€æ°ã«ä¿åãããŸãã pixelInãæãsrcImageå€æ°ã¯ãã¢ã«ãŽãªãºã åŒã§å ¥åãã¯ã»ã«ããŒã¿ãååŸããããã«äœ¿çšãããŸããããã§ãã«ã©ãŒã€ã¡ãŒãžã¯ã»ãã¢ããŒã³ã§ãã€ã³ããããŸãã
以äžã¯ãç»åãã»ãã¢è²ã«çè²ããããã®OpenCLã«ãŒãã«ã³ãŒãã§ãã
constant uchar4 cWhite = {1, 1, 1, 1}; constant float3 sepiaRed = {0.393f, 0.769f, 0.189f}; constant float3 sepiaGreen = {0.349f, 0.686f, 0.168f}; constant float3 sepiaBlue = {0.272f, 0.534f, 0.131f}; __kernel void sepia(__global uchar4 *in, __global uchar4 *out, int4 intArgs, int2 wh) { int x = get_global_id(0); int y = get_global_id(1); int width = wh.x; int height = wh.y; if(width <= x || height <= y) return; int xTouchApply = intArgs.x; int yTouchApply = intArgs.y; int radiusHi = intArgs.z; int radiusLo = intArgs.w; int tid = x + y * width; uchar4 c4 = in[tid]; float4 f4 = convert_float4(c4); int xRel = x - xTouchApply; int yRel = y - yTouchApply; int polar = xRel*xRel + yRel*yRel; uchar4 pixOut; if(polar > radiusHi || polar < radiusLo) { if(polar < radiusLo) { float4 outPixel; float tmpR = dot(f4.xyz, sepiaRed); float tmpG = dot(f4.xyz, sepiaGreen); float tmpB = dot(f4.xyz, sepiaBlue); outPixel = (float4)(tmpR, tmpG, tmpB, f4.w); pixOut = convert_uchar4_sat_rte(outPixel); } else { pixOut= c4; } } else { pixOut = cWhite; } out[tid] = pixOut; }
1.6 , RenderScript
, RenderScript- ? , â , , , . Android- , .
MainActivity.java .
private RenderScript rsContext;
rsContext RenderScript, RS-. RenderScript. RenderScript.
private ScriptC_plasma plasmaScript; private ScriptC_mono monoScript; private ScriptC_sepia sepiaScript;
plasmaScript , monoScript , sepiaScript â -, RS-. Eclipse IDE Android Studio Java- rs-. , plasma.rs ScriptC_plasma , mono.rs â ScriptC_mono . sepia.rs ScriptC_sepia . RenderScript- , gen . , sepia.rs ScriptC_sepia.java . Java-, rs- , RenderScript-, . - - MainActivity.java.
private Allocation allocationIn; private Allocation allocationOut; private Allocation allocationPalette; private Allocation allocationAngles;
Allocation RenderScript-. , allocationIn allocationOut . , allocationIn , allocationOut , RS- , , .
RenderScript-, , Activity . , , allocationPalette allocationAngle .
, RS-, RS-. initRS(âŠ) .
protected void initRS() { ⊠};
RenderScript, create RenderScript . , RenderScript, . RenderScript RS-. , RenderScript MainActivity . RenderScript.create(âŠ) « this ».
rsContext = RenderScript.create(this);
, RS-, RenderScript-, . , , initRS() , RenderScript- , .
if (effectName.equals("plasma")) { plasmaScript = new ScriptC_plasma(rsContext); } else if (effectName.equals("mono")) { monoScript = new ScriptC_mono(rsContext); } else if (effectName.equals("sepia")) { sepiaScript = new ScriptC_sepia(rsContext); } //
stepRenderScript(âŠ) , RenderScript- . RenderScript- RS-. stepRenderScript(âŠ) , .
private void stepRenderScript(âŠ) { ⊠// if(effectName.equals("plasma")) { plasmaScript.bind_gPalette(allocationPalette); plasmaScript.bind_gAngles(allocationAngles); plasmaScript.set_gx(inX - stepCount); plasmaScript.set_gy(inY - stepCount); plasmaScript.set_ts(System.currentTimeMillis() - mStartTime); plasmaScript.set_gScript(plasmaScript); plasmaScript.invoke_filter(plasmaScript, allocationIn, allocationOut); } else if(effectName.equals("mono")) { // , int radius = (stepApply == -1 ? -1 : 10*(stepCount - stepApply)); int radiusHi = (radius + 2)*(radius + 2); int radiusLo = (radius - 2)*(radius - 2); // . monoScript.set_radiusHi(radiusHi); monoScript.set_radiusLo(radiusLo); monoScript.set_xInput(xToApply); monoScript.set_yInput(yToApply); // . monoScript.forEach_root(allocationIn, allocationOut); if(stepCount > FX_COUNT) { stepCount = 0; stepApply = -1; } } else if(effectName.equals("sepia")) { ⊠// , } ⊠// };
RenderScript-, , gPalette , gAngles , gx , gy gScript . RS , . plasma.rs . , rs_allocation , bind_<var> . , bind_<gvars> , allocationPalette allocationAngles RenderScript-. , , gx , gy , ts gScript, set_<var> , . , , RenderScript- , x, y , . invoke_filter(âŠ) RenderScript. , , filter() , , , RenderScript.
, radius radiusHi radiusLo . , xInput yInput , . , , , , forEach_root() . forEach_root(âŠ) â, , RenderScript. , radiusHi , radiusLo , xInput yInput . set_<var> .
RenderScript , .
RenderScript :
#pragma version(1) #pragma rs java_package_name(com.example.imageprocessingoffload) rs_allocation *gPalette; rs_allocation *gAngles; rs_script gScript; float ts; int gx; int gy; static int32_t intFromFloat(float xfl) { return (int32_t)((xfl)*(1 << 16)); } const float YT1_INCR = (1/100.0f); const float YT2_INCR = (1/163.0f); const float XT1_INCR = (1/173.0f); const float XT2_INCR = (1/242.0f); static uint16_t pfrom_fixed(int32_t dx) { unsigned short *palette = (unsigned short *)gPalette; uint16_t ret; if (dx < 0) dx = -dx; if (dx >= (1 << 16)) dx = (1 << 16)-1; int idx = ((dx & ((1 << 16)-1)) >> 8); ret = palette[idx & ((1<<8)-1)]; return ret; } uint16_t __attribute__((kernel)) root(uint16_t in, uint32_t x, uint32_t y) { unsigned int *angles = (unsigned int *)gAngles; uint32_t out = in; int yt1 = intFromFloat(ts/1230.0f); int yt2 = yt1; int xt10 = intFromFloat(ts/3000.0f); int xt20 = xt10; int y1 = y*intFromFloat(YT1_INCR); int y2 = y*intFromFloat(YT2_INCR); yt1 = yt1 + y1; yt2 = yt2 + y2; int a1 = (yt1 >> 7) & ((1<<9)-1); int a2 = (yt2 >> 7) & ((1<<9)-1); int base = angles[a1] + angles[a2]; int xt1 = xt10; int xt2 = xt20; xt1 += x*intFromFloat(XT1_INCR); xt2 += x*intFromFloat(XT2_INCR); a1 = (xt1 >> (16-9)) & ((1<<9)-1); a2 = (xt2 >> (16-9)) & ((1<<9)-1); int ii = base + angles[a1] + angles[a2]; out = pfrom_fixed(ii/4); return out; } void filter(rs_script gScript, rs_allocation alloc_in, rs_allocation alloc_out) { //rsDebug("Inputs TS, X, Y:", ts, gx, gy); rsForEach(gScript, alloc_in, alloc_out); }
RenderScript :
#pragma version(1) #pragma rs java_package_name(com.example.imageprocessingoffload) int radiusHi; int radiusLo; int xToApply; int yToApply; const float4 gWhite = {1.f, 1.f, 1.f, 1.f}; const float3 channelWeights = {0.299f, 0.587f, 0.114f}; float saturationValue = 0.0f; uchar4 __attribute__((kernel)) root(const uchar4 in, uint32_t x, uint32_t y) { float4 f4 = rsUnpackColor8888(in); int xRel = x - xToApply; int yRel = y - yToApply; int polar = xRel*xRel + yRel*yRel; uchar4 out; if(polar > radiusHi || polar < radiusLo) { if(polar < radiusLo) { float3 outPixel = dot(f4.rgb, channelWeights); outPixel = mix( outPixel, f4.rgb, saturationValue); out = rsPackColorTo8888(outPixel); } else { out = rsPackColorTo8888(f4); } } else { out = rsPackColorTo8888(gWhite); } return out; }
RenderScript- .
#pragma version(1) #pragma rs java_package_name(com.example.imageprocessingoffload) #pragma rs_fp_relaxed int radiusHi; int radiusLo; int xTouchApply; int yTouchApply; rs_script gScript; const float4 gWhite = {1.f, 1.f, 1.f, 1.f}; const static float3 sepiaRed = {0.393f, 0.769f, 0.189f}; const static float3 sepiaGreen = {0.349f, 0.686, 0.168f}; const static float3 sepiaBlue = {0.272f, 0.534f, 0.131f}; uchar4 __attribute__((kernel)) sepia(uchar4 in, uint32_t x, uint32_t y) { uchar4 result; float4 f4 = rsUnpackColor8888(in); int xRel = x - xTouchApply; int yRel = y - yTouchApply; int polar = xRel*xRel + yRel*yRel; if(polar > radiusHi || polar < radiusLo) { if(polar < radiusLo) { float3 out; float tmpR = dot(f4.rgb, sepiaRed); float tmpG = dot(f4.rgb, sepiaGreen); float tmpB = dot(f4.rgb, sepiaBlue); out.r = tmpR; out.g = tmpG; out.b = tmpB; result = rsPackColorTo8888(out); } else { result = rsPackColorTo8888(f4); } } else { result = rsPackColorTo8888(gWhite); } return result; }
1.7
OpenCL, .
Intel INDE . IDE, . â Android Studio. , , -, IDE, ( â Android SDK NDK, ), - â OpenCL- OpenCL-. Android-, . , Root-.
, , OpenCL Android . Eclipse IDE, , Android Studio.
Android Studio . , Android Studio, . , , Android SDK, NDK , Intel OpenCL.
, Android.mk , OpenCL-. :
INTELOCLSDKROOT="C:\Intel\INDE\code_builder_5.1.0.25"
local.properties Android SDK NDK.
sdk.dir=C\:\\Intel\\INDE\\IDEintegration\\android-sdk-windows ndk.dir=C\:\\Intel\\INDE\\IDEintegration\\android-ndk-r10d
Android-. Intel Nexus 7 x86. Android Virtual Device Manager.
, , , OpenCL. , Run Android Studio . , Serial Number. emulator-5554 .
Windows :
C:\Intel\INDE\code_builder_5.1.0.25\android-preinstall>opencl_android_install âd emulator-5554
OpenCL- , . , , Android Studio , , OK. .
OpenCL-
, OpenCL- .
, OpenCL, RenderScript-. Android Studio Eclipse- RenderScript Android . , , RenderScript. .
OpenCL-. OpenCL .
2. - OpenCL
- OpenCL API OpenCL-. , - API, OpenCL. . . .
class openclWrapper { private: cl_device_id* mDeviceIds; // OpenCL- (CPU, GPU, ) cl_kernel mKernel; // cl_command_queue mCmdQue; // CL- cl_context mContext; // OpenCL cl_program mProgram; // OpenCL- public: openclWrapper() { mDeviceIds = NULL; mKernel = NULL; mCmdQue = NULL; mContext = NULL; mProgram = NULL; }; ~openclWrapper() { }; cl_context getContext() { return mContext; }; cl_kernel getKernel() { return mKernel; }; cl_command_queue getCmdQue() { return mCmdQue; }; int createContext(cl_device_type deviceType); bool LoadInlineSource(char* &sourceCode, const char* eName); bool LoadFileSource(char* &sourceCode, const char* eName, AAssetManager *mgr); int buildProgram(const char* eName, AAssetManager *mgr); int createCmdQueue(); int createKernel(const char *kname); // int initOpenCL(cl_device_type clDeviceType, const char* eName, AAssetManager *mgr=NULL); };
::createContext(cl device) . , (CPU GPU), , OpenCL . , OpenCL. OpenCL. , , SUCCESS ( mContext ). , , OpenCL, FAIL.
::createCmdQue() , OpenCL-. mContext . SUCCESS ( mCmdQue ). , , createContext(âŠ) , FAIL.
::buildProgram(effectName, AssetManager) . ( effectName ) Android JNI. OpenCL-, . - (inline) OpenCL-. NULL . , effectName , , . , OpenCL-, , . , , â OpenCL-. , OpenCL- , â , API OpenCL- .
- buildProgram(âŠ) OpenCL API clCreateProgramWithSource(âŠ) . API , OpenCL- . . OpenCL . , clCreateProgramWithSource(âŠ) .
- clBuildProgram(âŠ) , clCreateProgramWithSource(âŠ) clCreateProgramWithBinary(âŠ) . , , OpenCL-. - , clGetProgramBuildInfo(âŠ) . -.
::createKernel(âŠ) . SUCCESS. mKernel , , , .
::getContext() , ::getCmdQue() ::getKernel() , , , . JNI , OpenCL-.
ãŸãšã
OpenCL-, Android-. OpenCL RenderScript. , , . OpenCL , , , . , .