diff --git a/SConstruct b/SConstruct index 79cee49bb3..a4654561b0 100644 --- a/SConstruct +++ b/SConstruct @@ -257,7 +257,7 @@ env['HAVE_ALSA'] = conf.CheckPKG('alsa') # OpenCL if env['opencl']: - env['HAVE_OPENCL'] = conf.CheckPKG('opencl') + env['HAVE_OPENCL'] = 1 else: env['HAVE_OPENCL'] = 0 @@ -320,6 +320,7 @@ conf.Define('NOJIT', env['NOJIT']) conf.Define('HAVE_SDL', env['HAVE_SDL']) conf.Define('HAVE_BLUEZ', env['HAVE_BLUEZ']) conf.Define('HAVE_AO', env['HAVE_AO']) +conf.Define('HAVE_OPENCL', env['HAVE_OPENCL']) conf.Define('HAVE_OPENAL', env['HAVE_OPENAL']) conf.Define('HAVE_ALSA', env['HAVE_ALSA']) conf.Define('HAVE_WX', env['HAVE_WX']) diff --git a/Source/Core/Common/Src/OpenCL.cpp b/Source/Core/Common/Src/OpenCL.cpp index 06f39fb79a..8aebb4e2d3 100644 --- a/Source/Core/Common/Src/OpenCL.cpp +++ b/Source/Core/Common/Src/OpenCL.cpp @@ -21,18 +21,19 @@ #include "Common.h" namespace OpenCL { - -cl_context g_context = NULL; -cl_command_queue g_cmdq = NULL; + #if defined(HAVE_OPENCL) && HAVE_OPENCL + cl_device_id device_id = NULL; + cl_context g_context = NULL; + cl_command_queue g_cmdq = NULL; + #endif bool Initialize() { - if(g_context) - return false; #if defined(HAVE_OPENCL) && HAVE_OPENCL - + if(g_context) + return false; int err; // error code returned from api calls - cl_device_id device_id; + // Connect to a compute device // @@ -61,13 +62,13 @@ bool Initialize() { PanicAlert("Error: Failed to create a command commands!\n"); return false; } - + printf("Initialized OpenCL fine!\n"); return true; #else return false; #endif } - +#if defined(HAVE_OPENCL) && HAVE_OPENCL cl_context GetInstance() { return g_context; } @@ -76,20 +77,56 @@ cl_command_queue GetCommandQueue() { return g_cmdq; } -cl_program CompileProgram(const char *program, unsigned int size) { - // TODO - return NULL; -} +cl_program CompileProgram(const char *Kernel) { + int err; + cl_program program; + program = clCreateProgramWithSource(OpenCL::g_context, 1, (const char **) & Kernel, NULL, &err); + if (!program) + { + printf("Error: Failed to create compute program!\n"); + return NULL; + } + // Build the program executable + // + err = clBuildProgram(program , 0, NULL, NULL, NULL, NULL); + if (err != CL_SUCCESS) + { + size_t len; + char buffer[2048]; + + printf("Error: Failed to build program executable!\n"); + clGetProgramBuildInfo(program , OpenCL::device_id, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len); + printf("%s\n", buffer); + return NULL; + } + + + return program; +} +cl_kernel CompileKernel(cl_program program, const char *Function) +{ + int err; + // Create the compute kernel in the program we wish to run + // + cl_kernel kernel = clCreateKernel(program, Function, &err); + if (!kernel || err != CL_SUCCESS) + { + printf("Error: Failed to create compute kernel!\n"); + return NULL; + } + return kernel; +} +#endif void Destroy() { - if(!g_context) - return; #if defined(HAVE_OPENCL) && HAVE_OPENCL + if(!g_context) + return; clReleaseCommandQueue(g_cmdq); clReleaseContext(g_context); #endif } -}; \ No newline at end of file +}; diff --git a/Source/Core/Common/Src/OpenCL.h b/Source/Core/Common/Src/OpenCL.h index 8284158d53..31cfb084ee 100644 --- a/Source/Core/Common/Src/OpenCL.h +++ b/Source/Core/Common/Src/OpenCL.h @@ -18,6 +18,7 @@ #ifndef __OPENCL_H__ #define __OPENCL_H__ +#include "Config.h" // Change to #if 1 if you want to test OpenCL (and you have it) on Windows #if 0 #pragma comment(lib, "OpenCL.lib") @@ -37,10 +38,16 @@ typedef void *cl_context; typedef void *cl_command_queue; typedef void *cl_program; +typedef void *cl_kernel; #endif namespace OpenCL { + #if defined(HAVE_OPENCL) && HAVE_OPENCL + extern cl_device_id device_id; + extern cl_context g_context; + extern cl_command_queue g_cmdq; + #endif bool Initialize(); @@ -50,7 +57,8 @@ cl_command_queue GetCommandQueue(); void Destroy(); -cl_program CompileProgram(const char *program, unsigned int size); +cl_program CompileProgram(const char *Kernel); +cl_kernel CompileKernel(cl_program program, const char *Function); }; diff --git a/Source/Core/Common/Src/SConscript b/Source/Core/Common/Src/SConscript index bdee3f7a90..78eab23c09 100644 --- a/Source/Core/Common/Src/SConscript +++ b/Source/Core/Common/Src/SConscript @@ -23,6 +23,7 @@ files = [ "MemoryUtil.cpp", "Misc.cpp", "MsgHandler.cpp", + "OpenCL.cpp", "Plugin.cpp", "PluginDSP.cpp", "PluginWiimote.cpp", diff --git a/Source/Core/VideoCommon/Src/OpenCL/TextureDecoder.cpp b/Source/Core/VideoCommon/Src/OpenCL/TextureDecoder.cpp index 62d69afdfe..322d5dea64 100644 --- a/Source/Core/VideoCommon/Src/OpenCL/TextureDecoder.cpp +++ b/Source/Core/VideoCommon/Src/OpenCL/TextureDecoder.cpp @@ -18,6 +18,7 @@ #include "TextureDecoder.h" #include "OpenCL.h" +#include #include #include @@ -27,9 +28,6 @@ #include #include #include - -size_t global; // global domain size for our calculation -size_t local; // local domain size for our calculation struct sDecoders @@ -39,69 +37,176 @@ struct sDecoders const char **cKernel; }; -const char *Kernel = " \ -kernel void Decode(global uchar *dst, \ - const global uchar *src, \ - int width, int height) \ -{ \ - int x = get_global_id(0), y = get_global_id(1); \ - \ - for (int iy = 0; iy < 4; iy++, src += 8) { \ - u16 *ptr = (u16 *)dst + ((x + (y / width)) + iy) * \ - width + ((x * width + y) % width); \ - u16 *s = (u16 *)src; \ - for(int j = 0; j < 4; j++) \ - *ptr++ = Common::swap16(*s++); \ - } \ -}"; - -sDecoders Decoders[] = { {NULL, NULL, &Kernel}, +const char *Kernel = " \n \ +__kernel void Decode(__global unsigned char *dst, \n \ + const __global unsigned char *src, \n \ + const __global int width, const __global int height) \n \ +{ \n \ +// int x = get_global_id(0) % width, y = get_global_id(0) / width; \n \ + int srcOffset = 0; \n \ + for (int y = 0; y < height; y += 4) \n \ + for (int x = 0; x < width; x += 8) \n \ + for (int iy = 0; iy < 4; iy++, srcOffset += 8) \n\ + { \n \ + dst[(y + iy)*width + x] = src[srcOffset]; \n \ + dst[(y + iy)*width + x + 1] = src[srcOffset + 1]; \n \ + dst[(y + iy)*width + x + 2] = src[srcOffset + 2]; \n \ + dst[(y + iy)*width + x + 3] = src[srcOffset + 3]; \n \ + dst[(y + iy)*width + x + 4] = src[srcOffset + 4]; \n \ + dst[(y + iy)*width + x + 5] = src[srcOffset + 5]; \n \ + dst[(y + iy)*width + x + 6] = src[srcOffset + 6]; \n \ + dst[(y + iy)*width + x + 7] = src[srcOffset + 7]; \n \ + } \n \ +}\n"; +// memcpy(dst + (y + iy)*width+x, src, 8); +const char *KernelOld = " \n \ +__kernel void Decode(__global uchar *dst, \n \ + const __global uchar *src, \n \ + int width, int height) \n \ +{ \n \ + dst[get_global_id(0)] = 0xFF; \n \ +} \n "; +sDecoders Decoders[] = { {NULL, NULL, &Kernel}, }; bool g_Inited = false; PC_TexFormat TexDecoder_Decode_OpenCL(u8 *dst, const u8 *src, int width, int height, int texformat, int tlutaddr, int tlutfmt) { + int err; if(!g_Inited) { g_Inited = true; #if defined(HAVE_OPENCL) && HAVE_OPENCL - // TODO: Compile the program + // TODO: Switch this over to the OpenCl.h backend // Create the compute program from the source buffer // - /* - program = clCreateProgramWithSource(context, 1, (const char **) & KernelSource, NULL, &err); - if (!program) + + Decoders[0].program = clCreateProgramWithSource(OpenCL::g_context, 1, (const char **) & Kernel, NULL, &err); + if (!Decoders[0].program) { printf("Error: Failed to create compute program!\n"); - return EXIT_FAILURE; + return PC_TEX_FMT_NONE; } // Build the program executable // - err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL); + err = clBuildProgram(Decoders[0].program , 0, NULL, NULL, NULL, NULL); if (err != CL_SUCCESS) { size_t len; char buffer[2048]; printf("Error: Failed to build program executable!\n"); - clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len); + clGetProgramBuildInfo(Decoders[0].program , OpenCL::device_id, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len); printf("%s\n", buffer); exit(1); } // Create the compute kernel in the program we wish to run // - kernel = clCreateKernel(program, "Decoder", &err); - if (!kernel || err != CL_SUCCESS) + Decoders[0].kernel = clCreateKernel(Decoders[0].program, "Decode", &err); + if (!Decoders[0].kernel || err != CL_SUCCESS) { printf("Error: Failed to create compute kernel!\n"); exit(1); - } */ + } + #endif } + /*switch(texformat) + { + case GX_TF_I8: + { + int srcOffset = 0; + for (int y = 0; y < height; y ++) + for (int x = 0; x < width; x ++) + { + //printf("X: %d Y: %d, copying 32 bytes from %d to %d\n", x, y, srcOffset,(y)*width+x); + memcpy(dst + (y + 0)*width+x, src + srcOffset + 0, 8); + memcpy(dst + (y + 1)*width+x, src + srcOffset + 8, 8); + memcpy(dst + (y + 2)*width+x, src + srcOffset + 16, 8); + memcpy(dst + (y + 3)*width+x, src + srcOffset + 24, 8); + srcOffset += 4; + } + return PC_TEX_FMT_I8; + } + default: + return PC_TEX_FMT_NONE; + }*/ + switch(texformat) + { + case GX_TF_I8: + { + size_t global = 0; // global domain size for our calculation + size_t local = 0; // local domain size for our calculation + printf("width %d, height %d\n", width, height); + // Create the input and output arrays in device memory for our calculation + // + cl_mem _dst = clCreateBuffer(OpenCL::g_context, CL_MEM_WRITE_ONLY, TexDecoder_GetTextureSizeInBytes(width, height, texformat), NULL, NULL); + if (!dst) + { + printf("Error: Failed to allocate device memory!\n"); + exit(1); + } + cl_mem _src = clCreateBuffer(OpenCL::g_context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, TexDecoder_GetTextureSizeInBytes(width, height, texformat), (void*)src, NULL); + if (!src) + { + printf("Error: Failed to allocate device memory!\n"); + exit(1); + } + // Set the arguments to our compute kernel + // + err = 0; + err = clSetKernelArg(Decoders[0].kernel, 0, sizeof(cl_mem), &_dst); + err |= clSetKernelArg(Decoders[0].kernel, 1, sizeof(cl_mem), &_src); + err |= clSetKernelArg(Decoders[0].kernel, 2, sizeof(cl_int), &width); + err |= clSetKernelArg(Decoders[0].kernel, 3, sizeof(cl_int), &height); + if (err != CL_SUCCESS) + { + printf("Error: Failed to set kernel arguments! %d\n", err); + exit(1); + } + + // Get the maximum work group size for executing the kernel on the device + // + err = clGetKernelWorkGroupInfo(Decoders[0].kernel, OpenCL::device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(int), &local, NULL); + if (err != CL_SUCCESS) + { + printf("Error: Failed to retrieve kernel work group info! %d\n", err); + exit(1); + } + // Execute the kernel over the entire range of our 1d input data set + // using the maximum number of work group items for this device + // + global = width * height; + err = clEnqueueNDRangeKernel(OpenCL::g_cmdq, Decoders[0].kernel, 1, NULL, &global, &local, 0, NULL, NULL); + if (err != CL_SUCCESS) + { + printf("Error: Failed to execute kernel! %d\n", err); + return PC_TEX_FMT_NONE; + } + // Wait for the command commands to get serviced before reading back results + // + clFinish(OpenCL::g_cmdq); + + // Read back the results from the device to verify the output + // + err = clEnqueueReadBuffer( OpenCL::g_cmdq, _dst, CL_TRUE, 0, TexDecoder_GetTextureSizeInBytes(width, height, texformat), dst, 0, NULL, NULL ); + if (err != CL_SUCCESS) + { + printf("Error: Failed to read output array! %d\n", err); + exit(1); + } + clReleaseMemObject(_dst); + clReleaseMemObject(_src); + } + return PC_TEX_FMT_I8; + break; + default: + return PC_TEX_FMT_NONE; + } // TODO: clEnqueueNDRangeKernel diff --git a/Source/Core/VideoCommon/Src/TextureDecoder.cpp b/Source/Core/VideoCommon/Src/TextureDecoder.cpp index ead79400ca..441124f8d0 100644 --- a/Source/Core/VideoCommon/Src/TextureDecoder.cpp +++ b/Source/Core/VideoCommon/Src/TextureDecoder.cpp @@ -583,6 +583,8 @@ PC_TexFormat TexDecoder_Decode(u8 *dst, const u8 *src, int width, int height, in { #if defined(HAVE_OPENCL) && HAVE_OPENCL PC_TexFormat retval = TexDecoder_Decode_OpenCL(dst, src, width, height, texformat, tlutaddr, tlutfmt); + if(retval == PC_TEX_FMT_NONE) + retval = TexDecoder_Decode_real(dst,src,width,height,texformat,tlutaddr,tlutfmt); #else PC_TexFormat retval = TexDecoder_Decode_real(dst,src,width,height,texformat,tlutaddr,tlutfmt); #endif diff --git a/Source/Core/VideoCommon/Src/XFBConvert.cpp b/Source/Core/VideoCommon/Src/XFBConvert.cpp index e125e1435e..5d14a991ba 100644 --- a/Source/Core/VideoCommon/Src/XFBConvert.cpp +++ b/Source/Core/VideoCommon/Src/XFBConvert.cpp @@ -19,6 +19,8 @@ #include #endif +#include "OpenCL.h" + #include #include "XFBConvert.h" @@ -62,13 +64,119 @@ void InitXFBConvTables() } } + +const char *__ConvertFromXFB = "int bound(int i) \n \ +{ \n \ + return (i>255)?255:((i<0)?0:i); \n \ +} \n \ + \n \ +void yuv2rgb(int y, int u, int v, int &r, int &g, int &b) \n \ +{ \n \ + b = bound((76283*(y - 16) + 132252*(u - 128))>>16); \n \ + g = bound((76283*(y - 16) - 53281 *(v - 128) - 25624*(u - 128))>>16); //last one u? \n \ + r = bound((76283*(y - 16) + 104595*(v - 128))>>16); \n \ +} \n \ + \n \ +void ConvertFromXFB(u32 *dst, const u8* _pXFB) \n \ +{ \n \ + const unsigned char *src = _pXFB; \n \ + int id = get_global_id(0); \n \ + int srcOffset = id * 4; \n \ + int dstOffset = id; \n \ + u32 numBlocks = (width * height) / 2; \n \ + \n \ + int Y1 = src[srcOffset]; \n \ + int U = src[srcOffset + 1]; \n \ + int Y2 = src[srcOffset + 2]; \n \ + int V = src[srcOffset + 3]; \n \ + \n \ + int r, g, b; \n \ + yuv2rgb(Y1,U,V, r,g,b); \n \ + dst[dstOffset] = 0xFF000000 | (r<<16) | (g<<8) | (b); \n \ + yuv2rgb(Y2,U,V, r,g,b); \n \ + dst[dstOffset + 1] = 0xFF000000 | (r<<16) | (g<<8) | (b); \n \ +} \n"; + void ConvertFromXFB(u32 *dst, const u8* _pXFB, int width, int height) { if (((size_t)dst & 0xF) != 0) { PanicAlert("ConvertFromXFB - unaligned destination"); - } - const unsigned char *src = _pXFB; + } + const unsigned char *src = _pXFB; u32 numBlocks = ((width * height) / 2) / 2; + #if defined(HAVE_OPENCL) && HAVE_OPENCL + cl_kernel kernel; + cl_program program; + program = OpenCL::CompileProgram(__ConvertFromXFB); + kernel = OpenCL::CompileKernel(program, "ConvertFromXFB"); + int err; + + size_t global = 0; // global domain size for our calculation + size_t local = 0; // local domain size for our calculation + printf("width %d, height %d\n", width, height); + // Create the input and output arrays in device memory for our calculation + // + cl_mem _dst = clCreateBuffer(OpenCL::g_context, CL_MEM_WRITE_ONLY, sizeof(unsigned int) * numBlocks, NULL, NULL); + if (!dst) + { + printf("Error: Failed to allocate device memory!\n"); + exit(1); + } + cl_mem _src = clCreateBuffer(OpenCL::g_context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(unsigned char) * width * height, (void*)_pXFB, NULL); + if (!src) + { + printf("Error: Failed to allocate device memory!\n"); + exit(1); + } + // Set the arguments to our compute kernel + // + err = 0; + err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &_dst); + err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &_src); + if (err != CL_SUCCESS) + { + printf("Error: Failed to set kernel arguments! %d\n", err); + exit(1); + } + + // Get the maximum work group size for executing the kernel on the device + // + err = clGetKernelWorkGroupInfo(Decoders[0].kernel, OpenCL::device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(int), &local, NULL); + if (err != CL_SUCCESS) + { + printf("Error: Failed to retrieve kernel work group info! %d\n", err); + exit(1); + } + // Execute the kernel over the entire range of our 1d input data set + // using the maximum number of work group items for this device + // + global = numBlocks; + if(global < local) + { + // Global can't be less than local + } + err = clEnqueueNDRangeKernel(OpenCL::g_cmdq, kernel, 1, NULL, &global, &local, 0, NULL, NULL); + if (err != CL_SUCCESS) + { + printf("Error: Failed to execute kernel! %d\n", err); + return; + } + + // Wait for the command commands to get serviced before reading back results + // + clFinish(OpenCL::g_cmdq); + + // Read back the results from the device to verify the output + // + err = clEnqueueReadBuffer( OpenCL::g_cmdq, _dst, CL_TRUE, 0, sizeof(unsigned int) * numBlocks, dst, 0, NULL, NULL ); + if (err != CL_SUCCESS) + { + printf("Error: Failed to read output array! %d\n", err); + exit(1); + } + clReleaseMemObject(_dst); + clReleaseMemObject(_src); + #else for (u32 i = 0; i < numBlocks; i++) { __m128i y1 = _y[src[0]]; @@ -90,8 +198,27 @@ void ConvertFromXFB(u32 *dst, const u8* _pXFB, int width, int height) dst += 4; src += 8; } + #endif } +const char *__ConvertToXFB = "__kernel void ConvertToXFB(__global unsigned int *dst, __global const unsigned char* _pEFB) \n \ +{ \n \ + const unsigned char *src = _pEFB;\n \ + int id = get_global_id(0);\n \ + src += id * 8; \n \ + \n \ + int y1 = (((16843 * src[0]) + (33030 * src[1]) + (6423 * src[2])) >> 16) + 16; \n \ + int u1 = ((-(9699 * src[0]) - (19071 * src[1]) + (28770 * src[2])) >> 16) + 128;\n \ + src += 4;\n \ + \n \ + int y2 = (((16843 * src[0]) + (33030 * src[1]) + (6423 * src[2])) >> 16) + 16;\n \ + int v2 = (((28770 * src[0]) - (24117 * src[1]) - (4653 * src[2])) >> 16) + 128;\n \ + src += 4;\n \ + \n \ + dst[id] = (v2 << 24) | (y2 << 16) | (u1 << 8) | (y1); \n \ +} \n "; + + void ConvertToXFB(u32 *dst, const u8* _pEFB, int width, int height) { const unsigned char *src = _pEFB; @@ -100,7 +227,79 @@ void ConvertToXFB(u32 *dst, const u8* _pEFB, int width, int height) if (((size_t)dst & 0xF) != 0) { PanicAlert("ConvertToXFB - unaligned XFB"); } + #if defined(HAVE_OPENCL) && HAVE_OPENCL + cl_kernel kernel; + cl_program program; + program = OpenCL::CompileProgram(__ConvertToXFB); + kernel = OpenCL::CompileKernel(program, "ConvertToXFB"); + int err; + + size_t global = 0; // global domain size for our calculation + size_t local = 0; // local domain size for our calculation + printf("width %d, height %d\n", width, height); + // Create the input and output arrays in device memory for our calculation + // + cl_mem _dst = clCreateBuffer(OpenCL::g_context, CL_MEM_WRITE_ONLY, sizeof(unsigned int) * numBlocks, NULL, NULL); + if (!dst) + { + printf("Error: Failed to allocate device memory!\n"); + exit(1); + } + cl_mem _src = clCreateBuffer(OpenCL::g_context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(unsigned char) * width * height, (void*)_pEFB, NULL); + if (!src) + { + printf("Error: Failed to allocate device memory!\n"); + exit(1); + } + // Set the arguments to our compute kernel + // + err = 0; + err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &_dst); + err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &_src); + if (err != CL_SUCCESS) + { + printf("Error: Failed to set kernel arguments! %d\n", err); + exit(1); + } + + // Get the maximum work group size for executing the kernel on the device + // + err = clGetKernelWorkGroupInfo(Decoders[0].kernel, OpenCL::device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(int), &local, NULL); + if (err != CL_SUCCESS) + { + printf("Error: Failed to retrieve kernel work group info! %d\n", err); + exit(1); + } + // Execute the kernel over the entire range of our 1d input data set + // using the maximum number of work group items for this device + // + global = numBlocks; + if(global < local) + { + // Global can't be less than local + } + err = clEnqueueNDRangeKernel(OpenCL::g_cmdq, kernel, 1, NULL, &global, &local, 0, NULL, NULL); + if (err != CL_SUCCESS) + { + printf("Error: Failed to execute kernel! %d\n", err); + return; + } + // Wait for the command commands to get serviced before reading back results + // + clFinish(OpenCL::g_cmdq); + + // Read back the results from the device to verify the output + // + err = clEnqueueReadBuffer( OpenCL::g_cmdq, _dst, CL_TRUE, 0, sizeof(unsigned int) * numBlocks, dst, 0, NULL, NULL ); + if (err != CL_SUCCESS) + { + printf("Error: Failed to read output array! %d\n", err); + exit(1); + } + clReleaseMemObject(_dst); + clReleaseMemObject(_src); + #else for (u32 i = 0; i < numBlocks; i++) { __m128i yuyv0 = _mm_srai_epi32( @@ -127,5 +326,6 @@ void ConvertToXFB(u32 *dst, const u8* _pEFB, int width, int height) _mm_store_si128((__m128i *)dst, four_dest); dst += 4; } + #endif } diff --git a/Source/Plugins/Plugin_VideoOGL/Src/SConscript b/Source/Plugins/Plugin_VideoOGL/Src/SConscript index 675584b337..39b7139de1 100644 --- a/Source/Plugins/Plugin_VideoOGL/Src/SConscript +++ b/Source/Plugins/Plugin_VideoOGL/Src/SConscript @@ -40,7 +40,8 @@ files += [ 'main.cpp', 'GLUtil.cpp', ] - +if gfxenv['HAVE_OPENCL']: + libs += [ 'OpenCL'] if gfxenv['HAVE_WX']: files += [ 'GUI/ConfigDlg.cpp',