ShuriZma
/
dolphin
mirror of https://github.com/dolphin-emu/dolphin.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

More OpenCL work, got XFB converting which needs more optimization. I haven't checked for FPS changes. My desktop isn't the best to test on anyway (Phenom 1, 32 stream processors). The package check doesn't work for me, so I just checked true if you compile with opencl=true. Requires a bit of cleanup still 

git-svn-id: https://dolphin-emu.googlecode.com/svn/trunk@4369 8ced0084-cf51-0410-be5f-012b33b47a6e
This commit is contained in:
Sonicadvance1 2009-10-07 02:48:21 +00:00
parent f474788e01
commit 0146f3f58f
8 changed files with 406 additions and 51 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
SConstruct
View File

@ -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'])

65
Source/Core/Common/Src/OpenCL.cpp
View File

@ -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,16 +77,52 @@ cl_command_queue GetCommandQueue() {
return g_cmdq;
}
cl_program CompileProgram(const char *program, unsigned int size) {
// TODO
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

10
Source/Core/Common/Src/OpenCL.h
View File

@ -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);
};

1
Source/Core/Common/Src/SConscript
View File

@ -23,6 +23,7 @@ files = [
"MemoryUtil.cpp",
"Misc.cpp",
"MsgHandler.cpp",
"OpenCL.cpp",
"Plugin.cpp",
"PluginDSP.cpp",
"PluginWiimote.cpp",

167
Source/Core/VideoCommon/Src/OpenCL/TextureDecoder.cpp
View File

@ -18,6 +18,7 @@
#include "TextureDecoder.h"
#include "OpenCL.h"
#include <CL/cl.h>
#include <fcntl.h>
#include <stdio.h>
@ -28,9 +29,6 @@
#include <sys/types.h>
#include <sys/stat.h>
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
}
#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

2
Source/Core/VideoCommon/Src/TextureDecoder.cpp
View File

@ -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

200
Source/Core/VideoCommon/Src/XFBConvert.cpp
View File

@ -19,6 +19,8 @@
#include <intrin.h>
#endif
#include "OpenCL.h"
#include <xmmintrin.h>
#include "XFBConvert.h"
@ -62,6 +64,39 @@ 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) {
@ -69,6 +104,79 @@ void ConvertFromXFB(u32 *dst, const u8* _pXFB, int width, int height)
}
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
}

3
Source/Plugins/Plugin_VideoOGL/Src/SConscript
View File

@ -40,7 +40,8 @@ files += [
'main.cpp',
'GLUtil.cpp',
]
if gfxenv['HAVE_OPENCL']:
libs += [ 'OpenCL']
if gfxenv['HAVE_WX']:
files += [
'GUI/ConfigDlg.cpp',