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

Better TF_I8 decoding via OCL. Can still get way better though. Seems XFB speed boost was a fluke 

git-svn-id: https://dolphin-emu.googlecode.com/svn/trunk@4379 8ced0084-cf51-0410-be5f-012b33b47a6e
This commit is contained in:
Sonicadvance1 2009-10-07 21:10:39 +00:00
parent 9c45ac7e35
commit 5523f28b51
1 changed files with 15 additions and 8 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -42,10 +42,12 @@ __kernel void Decode(__global unsigned char *dst, \n \
const __global unsigned char *src, \n \ const __global unsigned char *src, \n \
const __global int width, const __global int height) \n \ const __global int width, const __global int height) \n \
{ \n \ { \n \
// int x = get_global_id(0) % width, y = get_global_id(0) / width; \n \ int x = get_global_id(0) % width, y = get_global_id(0) / width; \n \
int srcOffset = 0; \n \ if((y % 4) == 0 && (x % 8) == 0) \n \
for (int y = 0; y < height; y += 4) \n \ { \n \
for (int x = 0; x < width; x += 8) \n \ int srcOffset = (x * 4) + (y * width); \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\ for (int iy = 0; iy < 4; iy++, srcOffset += 8) \n\
{ \n \ { \n \
dst[(y + iy)*width + x] = src[srcOffset]; \n \ dst[(y + iy)*width + x] = src[srcOffset]; \n \
@ -57,6 +59,7 @@ __kernel void Decode(__global unsigned char *dst, \n \
dst[(y + iy)*width + x + 6] = src[srcOffset + 6]; \n \ dst[(y + iy)*width + x + 6] = src[srcOffset + 6]; \n \
dst[(y + iy)*width + x + 7] = src[srcOffset + 7]; \n \ dst[(y + iy)*width + x + 7] = src[srcOffset + 7]; \n \
} \n \ } \n \
} \n \
}\n"; }\n";
// memcpy(dst + (y + iy)*width+x, src, 8); // memcpy(dst + (y + iy)*width+x, src, 8);
const char *KernelOld = " \n \ const char *KernelOld = " \n \
@ -114,21 +117,25 @@ PC_TexFormat TexDecoder_Decode_OpenCL(u8 *dst, const u8 *src, int width, int hei
#endif #endif
} }
switch(texformat) /*switch(texformat)
{ {
case GX_TF_I8: case GX_TF_I8:
{ {
int srcOffset = 0;
for (int y = 0; y < height; y += 4) for (int y = 0; y < height; y += 4)
for (int x = 0; x < width; x += 8) for (int x = 0; x < width; x += 8)
for (int iy = 0; iy < 4; iy++, src += 8) for (int iy = 0; iy < 4; iy++, srcOffset += 8)
memcpy(dst + (y + iy)*width+x, src, 8); {
printf("x: %d y: %d offset: %d\n", x, y, srcOffset);
memcpy(dst + (y + iy)*width+x, src + srcOffset, 8);
}
return PC_TEX_FMT_I8; return PC_TEX_FMT_I8;
} }
break; break;
default: default:
return PC_TEX_FMT_NONE; return PC_TEX_FMT_NONE;
} }
//return PC_TEX_FMT_NONE; //return PC_TEX_FMT_NONE;*/
switch(texformat) switch(texformat)
{ {
case GX_TF_I8: case GX_TF_I8: