From 274aba1a89a85ec749fe08a40bbf4ffec381baf1 Mon Sep 17 00:00:00 2001 From: Andrey Pavlenko Date: Mon, 27 Jul 2015 02:57:01 +0300 Subject: [PATCH] adding OpenCL processing (Image2D-to-Image2D only, others will be added later) --- .../android/tutorial-4-opencl/jni/Android.mk | 11 +- .../tutorial-4-opencl/jni/Application.mk | 4 +- .../tutorial-4-opencl/jni/CLprocessor.cpp | 168 ++++++++++++++++++ .../tutorial-4-opencl/jni/GLrender.cpp | 32 +++- 4 files changed, 201 insertions(+), 14 deletions(-) create mode 100644 samples/android/tutorial-4-opencl/jni/CLprocessor.cpp diff --git a/samples/android/tutorial-4-opencl/jni/Android.mk b/samples/android/tutorial-4-opencl/jni/Android.mk index 1981bbdc0..a641a931b 100644 --- a/samples/android/tutorial-4-opencl/jni/Android.mk +++ b/samples/android/tutorial-4-opencl/jni/Android.mk @@ -1,9 +1,14 @@ LOCAL_PATH := $(call my-dir) include $(CLEAR_VARS) +LOCAL_MODULE := OpenCL +LOCAL_SRC_FILES := $(OPENCL_SDK)/lib/$(TARGET_ARCH_ABI)/libOpenCL.so +LOCAL_EXPORT_C_INCLUDES := $(OPENCL_SDK)/include +include $(PREBUILT_SHARED_LIBRARY) +include $(CLEAR_VARS) LOCAL_MODULE := JNIrender -LOCAL_SRC_FILES := jni.c GLrender.cpp -LOCAL_LDLIBS += -llog -lGLESv2 -lEGL - +LOCAL_SRC_FILES := jni.c GLrender.cpp CLprocessor.cpp +LOCAL_LDLIBS := -llog -lGLESv2 -lEGL +LOCAL_SHARED_LIBRARIES := OpenCL include $(BUILD_SHARED_LIBRARY) \ No newline at end of file diff --git a/samples/android/tutorial-4-opencl/jni/Application.mk b/samples/android/tutorial-4-opencl/jni/Application.mk index 5c7ca0e47..06db65762 100644 --- a/samples/android/tutorial-4-opencl/jni/Application.mk +++ b/samples/android/tutorial-4-opencl/jni/Application.mk @@ -1,4 +1,4 @@ -#APP_STL := gnustl_shared -#APP_GNUSTL_FORCE_CPP_FEATURES := exceptions rtti +APP_STL := gnustl_static +APP_GNUSTL_FORCE_CPP_FEATURES := exceptions rtti APP_ABI := armeabi-v7a APP_PLATFORM := android-14 diff --git a/samples/android/tutorial-4-opencl/jni/CLprocessor.cpp b/samples/android/tutorial-4-opencl/jni/CLprocessor.cpp new file mode 100644 index 000000000..03fec9c30 --- /dev/null +++ b/samples/android/tutorial-4-opencl/jni/CLprocessor.cpp @@ -0,0 +1,168 @@ +#define __CL_ENABLE_EXCEPTIONS +#include + +#include + +#include "common.hpp" + +const char oclProgB2B[] = "// clBuffer to clBuffer"; +const char oclProgI2B[] = "// clImage to clBuffer"; +const char oclProgI2I[] = \ + "__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST; \n" \ + "\n" \ + "__kernel void blur( \n" \ + " __read_only image2d_t imgIn, \n" \ + " __write_only image2d_t imgOut, \n" \ + " __private int size \n" \ + " ) { \n" \ + " \n" \ + " const int2 pos = {get_global_id(0), get_global_id(1)}; \n" \ + " \n" \ + " float4 sum = (float4) 0.0f; \n" \ + " for(int x = -size/2; x <= size/2; x++) { \n" \ + " for(int y = -size/2; y <= size/2; y++) { \n" \ + " sum += read_imagef(imgIn, sampler, pos + (int2)(x,y)); \n" \ + " } \n" \ + " } \n" \ + " \n" \ + " write_imagef(imgOut, pos, sum/size/size); \n" \ + "} \n"; + +void dumpCLinfo() +{ + LOGD("*** OpenCL info ***"); + try + { + std::vector platforms; + cl::Platform::get(&platforms); + LOGD("OpenCL info: Found %d OpenCL platforms", platforms.size()); + for (int i = 0; i < platforms.size(); ++i) + { + std::string name = platforms[i].getInfo(); + std::string version = platforms[i].getInfo(); + std::string profile = platforms[i].getInfo(); + std::string extensions = platforms[i].getInfo(); + LOGD( "OpenCL info: Platform[%d] = %s, ver = %s, prof = %s, ext = %s", + i, name.c_str(), version.c_str(), profile.c_str(), extensions.c_str() ); + } + + std::vector devices; + platforms[0].getDevices(CL_DEVICE_TYPE_ALL, &devices); + + for (int i = 0; i < devices.size(); ++i) + { + std::string name = devices[i].getInfo(); + std::string extensions = devices[i].getInfo(); + cl_ulong type = devices[i].getInfo(); + LOGD( "OpenCL info: Device[%d] = %s (%s), ext = %s", + i, name.c_str(), (type==CL_DEVICE_TYPE_GPU ? "GPU" : "CPU"), extensions.c_str() ); + } + } + catch(cl::Error& e) + { + LOGE( "OpenCL info: error while gathering OpenCL info: %s (%d)", e.what(), e.err() ); + } + catch(std::exception& e) + { + LOGE( "OpenCL info: error while gathering OpenCL info: %s", e.what() ); + } + catch(...) + { + LOGE( "OpenCL info: unknown error while gathering OpenCL info" ); + } + LOGD("*******************"); +} + +cl::Context theContext; +cl::CommandQueue theQueue; +cl::Program theProgB2B, theProgI2B, theProgI2I; + +void initCL() +{ + EGLDisplay mEglDisplay = eglGetCurrentDisplay(); + if (mEglDisplay == EGL_NO_DISPLAY) + LOGE("initCL: eglGetCurrentDisplay() returned 'EGL_NO_DISPLAY', error = %x", eglGetError()); + + EGLContext mEglContext = eglGetCurrentContext(); + if (mEglContext == EGL_NO_CONTEXT) + LOGE("initCL: eglGetCurrentContext() returned 'EGL_NO_CONTEXT', error = %x", eglGetError()); + + cl_context_properties props[] = + { CL_GL_CONTEXT_KHR, (cl_context_properties) mEglContext, + CL_EGL_DISPLAY_KHR, (cl_context_properties) mEglDisplay, + CL_CONTEXT_PLATFORM, 0, + 0 }; + + try + { + cl::Platform p = cl::Platform::getDefault(); + std::string ext = p.getInfo(); + if(ext.find("cl_khr_gl_sharing") == std::string::npos) + LOGE("Warning: CL-GL sharing isn't supported by PLATFORM"); + props[5] = (cl_context_properties) p(); + + theContext = cl::Context(CL_DEVICE_TYPE_GPU, props); + std::vector devs = theContext.getInfo(); + LOGD("Context returned %d devices, taking the 1st one", devs.size()); + ext = devs[0].getInfo(); + if(ext.find("cl_khr_gl_sharing") == std::string::npos) + LOGE("Warning: CL-GL sharing isn't supported by DEVICE"); + + theQueue = cl::CommandQueue(theContext, devs[0]); + + cl::Program::Sources src(1, std::make_pair(oclProgI2I, sizeof(oclProgI2I))); + theProgI2I = cl::Program(theContext, src); + theProgI2I.build(devs); + } + catch(cl::Error& e) + { + LOGE("cl::Error: %s (%d)", e.what(), e.err()); + } + catch(std::exception& e) + { + LOGE("std::exception: %s", e.what()); + } + catch(...) + { + LOGE( "OpenCL info: unknown error while initializing OpenCL stuff" ); + } + LOGD("initCL completed"); +} + +void closeCL() +{ +} + +#define GL_TEXTURE_2D 0x0DE1 +void procOCL_I2I(int texIn, int texOut, int w, int h) +{ + LOGD("procOCL_I2I(%d, %d, %d, %d)", texIn, texOut, w, h); + cl::ImageGL imgIn (theContext, CL_MEM_READ_ONLY, GL_TEXTURE_2D, 0, texIn); + cl::ImageGL imgOut(theContext, CL_MEM_WRITE_ONLY, GL_TEXTURE_2D, 0, texOut); + std::vector < cl::Memory > images; + images.push_back(imgIn); + images.push_back(imgOut); + + int64_t t = getTimeMs(); + theQueue.enqueueAcquireGLObjects(&images); + theQueue.finish(); + LOGD("enqueueAcquireGLObjects() costs %d ms", getTimeInterval(t)); + + t = getTimeMs(); + cl::Kernel blur(theProgI2I, "blur"); //TODO: may be done once + blur.setArg(0, imgIn); + blur.setArg(1, imgOut); + blur.setArg(2, 5); //5x5 + theQueue.finish(); + LOGD("Kernel() costs %d ms", getTimeInterval(t)); + + t = getTimeMs(); + theQueue.enqueueNDRangeKernel(blur, cl::NullRange, cl::NDRange(w, h), cl::NullRange); + theQueue.finish(); + LOGD("enqueueNDRangeKernel() costs %d ms", getTimeInterval(t)); + + t = getTimeMs(); + theQueue.enqueueReleaseGLObjects(&images); + theQueue.finish(); + LOGD("enqueueReleaseGLObjects() costs %d ms", getTimeInterval(t)); +} diff --git a/samples/android/tutorial-4-opencl/jni/GLrender.cpp b/samples/android/tutorial-4-opencl/jni/GLrender.cpp index a94f33f73..ad239dd4f 100644 --- a/samples/android/tutorial-4-opencl/jni/GLrender.cpp +++ b/samples/android/tutorial-4-opencl/jni/GLrender.cpp @@ -51,7 +51,7 @@ const char fss2D[] = \ int progOES = 0; int prog2D = 0; -GLuint FBOtex = 0; +GLuint FBOtex = 0, FBOtex2 = 0; GLuint FBO = 0; GLuint texOES = 0; @@ -75,6 +75,7 @@ static void releaseFBO() FBO = 0; } deleteTex(&FBOtex); + deleteTex(&FBOtex2); glDeleteProgram(prog2D); prog2D = 0; } @@ -127,10 +128,20 @@ static int makeShaderProg(const char* vss, const char* fss) return program; } + static void initFBO(int width, int height) { + LOGD("initFBO(%d, %d)", width, height); releaseFBO(); + glGenTextures(1, &FBOtex2); + glBindTexture(GL_TEXTURE_2D, FBOtex2); + glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA, width, height, 0, GL_RGBA, GL_UNSIGNED_BYTE, 0); + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_EDGE); + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_EDGE); + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST); + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST); + glGenTextures(1, &FBOtex); glBindTexture(GL_TEXTURE_2D, FBOtex); glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA, width, height, 0, GL_RGBA, GL_UNSIGNED_BYTE, 0); @@ -227,20 +238,17 @@ void drawFrameProcCPU() drawTex(FBOtex, GL_TEXTURE_2D, 0); } -void procOCL(int tex, int w, int h) -{ - //TODO: not yet implemented -} - +void procOCL(int tex, int w, int h); +void procOCL_I2I(int texIn, int texOut, int w, int h); void drawFrameProcOCL() { drawTex(texOES, GL_TEXTURE_EXTERNAL_OES, FBO); // modify pixels in FBO texture using OpenCL and CL-GL interop - procOCL(FBOtex, texWidth, texHeight); + procOCL_I2I(FBOtex, FBOtex2, texWidth, texHeight); // render to screen - drawTex(FBOtex, GL_TEXTURE_2D, 0); + drawTex(FBOtex2, GL_TEXTURE_2D, 0); } @@ -249,13 +257,16 @@ extern "C" void drawFrame() LOGD("*** drawFrame() ***"); int64_t t = getTimeMs(); //drawFrameOrig(); - drawFrameProcCPU(); + //drawFrameProcCPU(); + drawFrameProcOCL(); glFinish(); LOGD("*** drawFrame() costs %d ms ***", getTimeInterval(t)); } +void closeCL(); extern "C" void closeGL() { + closeCL(); LOGD("closeGL"); deleteTex(&texOES); @@ -266,6 +277,7 @@ extern "C" void closeGL() releaseFBO(); } +void initCL(); extern "C" int initGL() { LOGD("initGL"); @@ -287,6 +299,8 @@ extern "C" int initGL() glTexParameteri(GL_TEXTURE_EXTERNAL_OES, GL_TEXTURE_MIN_FILTER, GL_NEAREST); glTexParameteri(GL_TEXTURE_EXTERNAL_OES, GL_TEXTURE_MAG_FILTER, GL_NEAREST); + initCL(); + return texOES; }