From c8f863dfc55bbd95b6c3265eb12aa29e31663766 Mon Sep 17 00:00:00 2001 From: Andrey Pavlenko Date: Wed, 29 Jul 2015 00:13:35 +0300 Subject: [PATCH] minor fixes --- .../tutorial-4-opencl/jni/CLprocessor.cpp | 26 ++++--- .../tutorial-4-opencl/jni/GLrender.cpp | 71 +++++++++++++------ 2 files changed, 62 insertions(+), 35 deletions(-) diff --git a/samples/android/tutorial-4-opencl/jni/CLprocessor.cpp b/samples/android/tutorial-4-opencl/jni/CLprocessor.cpp index 66699aa52b..6d843acdd3 100644 --- a/samples/android/tutorial-4-opencl/jni/CLprocessor.cpp +++ b/samples/android/tutorial-4-opencl/jni/CLprocessor.cpp @@ -13,22 +13,21 @@ 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" \ + "__kernel void Laplacian( \n" \ " __read_only image2d_t imgIn, \n" \ - " __write_only image2d_t imgOut, \n" \ - " __private int size \n" \ + " __write_only image2d_t imgOut \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" \ + " sum += read_imagef(imgIn, sampler, pos + (int2)(-1,0)); \n" \ + " sum += read_imagef(imgIn, sampler, pos + (int2)(+1,0)); \n" \ + " sum += read_imagef(imgIn, sampler, pos + (int2)(0,-1)); \n" \ + " sum += read_imagef(imgIn, sampler, pos + (int2)(0,+1)); \n" \ + " sum -= read_imagef(imgIn, sampler, pos) * 4; \n" \ " \n" \ - " write_imagef(imgOut, pos, sum/size/size); \n" \ + " write_imagef(imgOut, pos, sum*10); \n" \ "} \n"; void dumpCLinfo() @@ -160,15 +159,14 @@ void procOCL_I2I(int texIn, int texOut, int w, int h) 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 + cl::Kernel Laplacian(theProgI2I, "Laplacian"); //TODO: may be done once + Laplacian.setArg(0, imgIn); + Laplacian.setArg(1, imgOut); theQueue.finish(); LOGD("Kernel() costs %d ms", getTimeInterval(t)); t = getTimeMs(); - theQueue.enqueueNDRangeKernel(blur, cl::NullRange, cl::NDRange(w, h), cl::NullRange); + theQueue.enqueueNDRangeKernel(Laplacian, cl::NullRange, cl::NDRange(w, h), cl::NullRange); theQueue.finish(); LOGD("enqueueNDRangeKernel() 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 da671fd59f..9e9342af3b 100644 --- a/samples/android/tutorial-4-opencl/jni/GLrender.cpp +++ b/samples/android/tutorial-4-opencl/jni/GLrender.cpp @@ -3,7 +3,7 @@ #include "common.hpp" -float vertexes[] = { +float vertices[] = { -1.0f, -1.0f, -1.0f, 1.0f, 1.0f, -1.0f, @@ -28,7 +28,7 @@ const char vss[] = \ "varying vec2 texCoord;\n" \ "void main() {\n" \ " texCoord = vTexCoord;\n" \ - " gl_Position = vec4 ( vPosition.x, vPosition.y, 0.0, 1.0 );\n" \ + " gl_Position = vec4 ( vPosition, 0.0f, 1.0f );\n" \ "}"; const char fssOES[] = \ @@ -48,8 +48,11 @@ const char fss2D[] = \ " gl_FragColor = texture2D(sTexture,texCoord);\n" \ "}"; -int progOES = 0; -int prog2D = 0; +GLuint progOES = 0; +GLuint prog2D = 0; + +GLint vPosOES, vTCOES; +GLint vPos2D, vTC2D; GLuint FBOtex = 0, FBOtex2 = 0; GLuint FBO = 0; @@ -80,13 +83,16 @@ static void releaseFBO() prog2D = 0; } -static inline void logShaderCompileError(GLuint shader) +static inline void logShaderCompileError(GLuint shader, bool isProgram = false) { GLchar msg[512]; msg[0] = 0; GLsizei len; - glGetShaderInfoLog(shader, sizeof(msg) - 1, &len, msg); - LOGE("Could not compile shader: %s", msg); + if(isProgram) + glGetProgramInfoLog(shader, sizeof(msg)-1, &len, msg); + else + glGetShaderInfoLog(shader, sizeof(msg)-1, &len, msg); + LOGE("Could not compile shader/program: %s", msg); } static int makeShaderProg(const char* vss, const char* fss) @@ -96,9 +102,9 @@ static int makeShaderProg(const char* vss, const char* fss) const GLchar* text = vss; glShaderSource(vshader, 1, &text, 0); glCompileShader(vshader); - int compiled; + GLint compiled; glGetShaderiv(vshader, GL_COMPILE_STATUS, &compiled); - if (compiled == 0) { + if (!compiled) { logShaderCompileError(vshader); glDeleteShader(vshader); vshader = 0; @@ -110,7 +116,7 @@ static int makeShaderProg(const char* vss, const char* fss) glShaderSource(fshader, 1, &text, 0); glCompileShader(fshader); glGetShaderiv(fshader, GL_COMPILE_STATUS, &compiled); - if (compiled == 0) { + if (!compiled) { logShaderCompileError(fshader); glDeleteShader(fshader); fshader = 0; @@ -121,6 +127,23 @@ static int makeShaderProg(const char* vss, const char* fss) glAttachShader(program, vshader); glAttachShader(program, fshader); glLinkProgram(program); + GLint linked; + glGetProgramiv(program, GL_LINK_STATUS, &linked); + if (!linked) + { + logShaderCompileError(program, true); + glDeleteProgram(program); + program = 0; + } + glValidateProgram(program); + GLint validated; + glGetProgramiv(program, GL_VALIDATE_STATUS, &validated); + if (!validated) + { + logShaderCompileError(program, true); + glDeleteProgram(program); + program = 0; + } if(vshader) glDeleteShader(vshader); if(fshader) glDeleteShader(fshader); @@ -160,6 +183,10 @@ static void initFBO(int width, int height) LOGE("initFBO failed: %d", glCheckFramebufferStatus(GL_FRAMEBUFFER)); prog2D = makeShaderProg(vss, fss2D); + vPos2D = glGetAttribLocation(prog2D, "vPosition"); + vTC2D = glGetAttribLocation(prog2D, "vTexCoord"); + glEnableVertexAttribArray(vPos2D); + glEnableVertexAttribArray(vTC2D); } void drawTex(int tex, GLenum texType, GLuint fbo) @@ -171,20 +198,18 @@ void drawTex(int tex, GLenum texType, GLuint fbo) glClear(GL_COLOR_BUFFER_BIT); - int prog = texType == GL_TEXTURE_EXTERNAL_OES ? progOES : prog2D; + GLuint prog = texType == GL_TEXTURE_EXTERNAL_OES ? progOES : prog2D; + GLint vPos = texType == GL_TEXTURE_EXTERNAL_OES ? vPosOES : vPos2D; + GLint vTC = texType == GL_TEXTURE_EXTERNAL_OES ? vTCOES : vTC2D; + float* texCoord = texType == GL_TEXTURE_EXTERNAL_OES ? texCoordOES : texCoord2D; glUseProgram(prog); - int vPos = glGetAttribLocation(prog, "vPosition"); - int vTC = glGetAttribLocation(prog, "vTexCoord"); + glVertexAttribPointer(vPos, 2, GL_FLOAT, false, 4*2, vertices); + glVertexAttribPointer(vTC, 2, GL_FLOAT, false, 4*2, texCoord); glActiveTexture(GL_TEXTURE0); glBindTexture(texType, tex); glUniform1i(glGetUniformLocation(prog, "sTexture"), 0); - glVertexAttribPointer(vPos, 2, GL_FLOAT, false, 4*2, vertexes); - glVertexAttribPointer(vTC, 2, GL_FLOAT, false, 4*2, texType == GL_TEXTURE_EXTERNAL_OES ? texCoordOES : texCoord2D); - glEnableVertexAttribArray(vPos); - glEnableVertexAttribArray(vTC); - glDrawArrays(GL_TRIANGLE_STRIP, 0, 4); glFlush(); LOGD("drawTex(%u) costs %d ms", tex, getTimeInterval(t)); @@ -245,11 +270,11 @@ void drawFrameProcOCL() drawTex(texOES, GL_TEXTURE_EXTERNAL_OES, FBO); // modify pixels in FBO texture using OpenCL and CL-GL interop - //procOCL_I2I(FBOtex, FBOtex2, texWidth, texHeight); - procOCL_OCV(FBOtex, texWidth, texHeight); + procOCL_I2I(FBOtex, FBOtex2, texWidth, texHeight); + //procOCL_OCV(FBOtex, texWidth, texHeight); // render to screen - drawTex(/*FBOtex2*/FBOtex, GL_TEXTURE_2D, 0); + drawTex(FBOtex2, GL_TEXTURE_2D, 0); } @@ -289,6 +314,10 @@ extern "C" int initGL() LOGD("GL_VERSION = %s", vs); progOES = makeShaderProg(vss, fssOES); + vPosOES = glGetAttribLocation(progOES, "vPosition"); + vTCOES = glGetAttribLocation(progOES, "vTexCoord"); + glEnableVertexAttribArray(vPosOES); + glEnableVertexAttribArray(vTCOES); glClearColor(1.0f, 1.0f, 1.0f, 1.0f);