piwalker/dolphin
1
0
Fork 0
mirror of https://github.com/dolphin-emu/dolphin.git synced 2025-07-23 06:09:50 -06:00
Code Issues Packages Projects Releases Wiki Activity

Dolphin now uses the first OpenCL device found on the system.

Browse Source 
Added cast in TextureDecoder.cl to satisfy the OpenCL compiler (thanks, Orphis).

git-svn-id: https://dolphin-emu.googlecode.com/svn/trunk@4718 8ced0084-cf51-0410-be5f-012b33b47a6e
This commit is contained in:
skidau
2009-12-23 01:31:42 +00:00
parent 24d8839793
commit 043681a37f
2 changed files with 163 additions and 153 deletions
Download Patch File Download Diff File Expand all files Collapse all files

103
Data/User/OpenCL/TextureDecoder.cl
View File

@ -24,9 +24,9 @@ kernel void DecodeI4(global uchar *dst,
{
uchar4 val = vload4(srcOffset, src);
uchar8 res;
res.even = (val >> 4) & 0x0F;
res.odd = val & 0x0F;
res |= res << 4;
res.even = (val >> (uchar4)4) & (uchar4)0x0F;
res.odd = val & (uchar4)0x0F;
res |= res << (uchar8)4;
vstore8(res, 0, dst + ((y + iy)*width + x));
srcOffset++;
}
@ -70,9 +70,9 @@ kernel void DecodeIA4(global uchar *dst,
{
uchar8 val = vload8(srcOffset, src);
uchar16 res;
res.odd = (val >> 4) & 0x0F;
res.even = val & 0x0F;
res |= res << 4;
res.odd = (val >> (uchar8)4) & (uchar8)0x0F;
res.even = val & (uchar8)0x0F;
res |= res << (uchar16)4;
vstore16(res, 0, dst + ((y + iy)*width + x) * 2);
srcOffset++;
}
@ -105,7 +105,7 @@ kernel void DecodeRGB565(global ushort *dst,
for (int iy = 0; iy < 4; iy++)
{
ushort4 val = vload4(srcOffset, src);
val = (val >> 8) | (val << 8);
val = (val >> (ushort4)8) | (val << (ushort4)8);
vstore4(val, 0, dst + ((y + iy)*width + x));
srcOffset++;
}
@ -119,32 +119,32 @@ kernel void DecodeRGB5A3(global uchar *dst,
for (int iy = 0; iy < 4; iy++)
{
ushort8 val = convert_ushort8(vload8(srcOffset, src));
ushort4 vs = val.odd | (val.even << 8);
ushort4 vs = val.odd | (ushort4)(val.even << (ushort4)8);
uchar16 resNoAlpha;
resNoAlpha.s26AE = convert_uchar4(vs >> 7); // R
resNoAlpha.s159D = convert_uchar4(vs >> 2); // G
resNoAlpha.s048C = convert_uchar4(vs << 3); // B
resNoAlpha &= 0xF8;
resNoAlpha |= (resNoAlpha >> 5) & 3; // 5 -> 8
resNoAlpha.s26AE = convert_uchar4(vs >> (ushort4)7); // R
resNoAlpha.s159D = convert_uchar4(vs >> (ushort4)2); // G
resNoAlpha.s048C = convert_uchar4(vs << (ushort4)3); // B
resNoAlpha &= (uchar16)0xF8;
resNoAlpha |= (uchar16)(resNoAlpha >> (uchar16)5) & (uchar16)3; // 5 -> 8
resNoAlpha.s37BF = (uchar4)(0xFF);
uchar16 resAlpha;
resAlpha.s26AE = convert_uchar4(vs >> 8); // R
resAlpha.s159D = convert_uchar4(vs >> 4); // G
resAlpha.s26AE = convert_uchar4(vs >> (ushort4)8); // R
resAlpha.s159D = convert_uchar4(vs >> (ushort4)4); // G
resAlpha.s048C = convert_uchar4(vs); // B
resAlpha &= 0x0F;
resAlpha |= (resAlpha << 4);
resAlpha.s37BF = convert_uchar4(vs >> 7) & 0xE0;
resAlpha.s37BF |= ((resAlpha.s37BF >> 3) & 0x1C)
| ((resAlpha.s37BF >> 6) & 0x3);
resAlpha &= (uchar16)0x0F;
resAlpha |= (resAlpha << (uchar16)4);
resAlpha.s37BF = convert_uchar4(vs >> (ushort4)7) & (uchar4)0xE0;
resAlpha.s37BF |= ((resAlpha.s37BF >> (uchar4)3) & (uchar4)0x1C)
| ((resAlpha.s37BF >> (uchar4)6) & (uchar4)0x3);
uchar16 choice = (uchar16)((uchar4)(vs.s0 >> 8),
(uchar4)(vs.s1 >> 8),
(uchar4)(vs.s2 >> 8),
(uchar4)(vs.s3 >> 8));
uchar16 res;
res = select(resAlpha, resNoAlpha, choice);
vstore16(res, 0, dst + ((y + iy)*width + x) * 4);
vstore16(res, 0, dst + ((y + iy) * width + x) * 4);
srcOffset++;
}
}
@ -157,47 +157,33 @@ uint4 unpack2bits(uchar b)
b & 3);
}
/*
Lots of debug code there that I'm using to find the problems with CMPR decoding
I think blocks having no alpha are properly decoded, only the blocks with alpha
are problematic. This is WIP !
*/
kernel void decodeCMPRBlock(global uchar *dst,
const global uchar *src, int width)
{
int x = get_global_id(0) * 4, y = get_global_id(1) * 4;
int srcOffset = (x * 2 + y * width / 2) / 8; //x / 4 + y * width / 16; //(x * 4) + (y * width) / 16;
uchar8 val = vload8(0, src);
ushort2 color565 = (ushort2)((val.s1 & 0xFF) | (val.s0 << 8), (val.s3 & 0xFF) | (val.s2 << 8));
uchar8 color32 = convert_uchar8((ushort8)
(((color565 << 3) & 0xF8) | ((color565 >> 2) & 0x7),
((color565 >> 3) & 0xFC) | ((color565 >> 9) & 0x3),
((color565 >> 8) & 0xF8) | ((color565 >> 13) & 0x7),
(((color565 << (ushort2)3) & (ushort2)0xF8) | ((color565 >> (ushort2)2) & (ushort2)0x7),
((color565 >> (ushort2)3) & (ushort2)0xFC) | ((color565 >> (ushort2)9) & (ushort2)0x3),
((color565 >> (ushort2)8) & (ushort2)0xF8) | ((color565 >> (ushort2)13) & (ushort2)0x7),
0xFF, 0xFF));
uint4 colors;
//uint4 choice = (uint4)((color565.s0 - color565.s1) << 16);
uint4 colorNoAlpha;
//uchar4 frac = (color32.odd - color32.even) / 2;
//frac = frac - (frac / 4);
uchar4 frac = convert_uchar4((((convert_ushort4(color32.even) & 0xFF) - (convert_ushort4(color32.odd) & 0xFF)) * 3) / 8);
uchar4 frac = convert_uchar4((((convert_ushort4(color32.even) & (ushort4)0xFF) - (convert_ushort4(color32.odd) & (ushort4)0xFF)) * (ushort4)3) / (ushort4)8);
colorNoAlpha = convert_uint4(color32.odd + frac);
colorNoAlpha = (colorNoAlpha << 8) | convert_uint4(color32.even - frac);
colorNoAlpha = (colorNoAlpha << 8) | convert_uint4(color32.odd);
colorNoAlpha = (colorNoAlpha << 8) | convert_uint4(color32.even);
colorNoAlpha = (colorNoAlpha << (uint4)8) | convert_uint4(color32.even - frac);
colorNoAlpha = (colorNoAlpha << (uint4)8) | convert_uint4(color32.odd);
colorNoAlpha = (colorNoAlpha << (uint4)8) | convert_uint4(color32.even);
uint4 colorAlpha;
//uchar4 midpoint = rhadd(color32.odd, color32.even);
uchar4 midpoint = convert_uchar4((convert_ushort4(color32.odd) + convert_ushort4(color32.even) + 1) / 2);
uchar4 midpoint = convert_uchar4((convert_ushort4(color32.odd) + convert_ushort4(color32.even) + (ushort4)1) / (ushort4)2);
midpoint.s3 = 0xFF;
colorAlpha = convert_uint4((uchar4)(0, 0, 0, 0));
colorAlpha = (colorAlpha << 8) | convert_uint4(midpoint);
colorAlpha = (colorAlpha << 8) | convert_uint4(color32.odd);
colorAlpha = (colorAlpha << 8) | convert_uint4(color32.even);
colorAlpha = (colorAlpha << (uint4)8) | convert_uint4(midpoint);
colorAlpha = (colorAlpha << (uint4)8) | convert_uint4(color32.odd);
colorAlpha = (colorAlpha << (uint4)8) | convert_uint4(color32.even);
//colorNoAlpha = (uint4)(0xFFFFFFFF);
//colorAlpha = (uint4)(0, 0, 0, 0xFFFFFFFF);
//colors = select(colorAlpha, colorNoAlpha, choice);
colors = color565.s0 > color565.s1 ? colorNoAlpha : colorAlpha;
uint16 colorsFull = (uint16)(colors, colors, colors, colors);
@ -206,31 +192,18 @@ kernel void decodeCMPRBlock(global uchar *dst,
uint4 shift1 = unpack2bits(val.s5);
uint4 shift2 = unpack2bits(val.s6);
uint4 shift3 = unpack2bits(val.s7);
uint16 shifts = (uint16)((uint4)(shift3.s0), (uint4)(shift3.s1), (uint4)(shift3.s2), (uint4)(shift3.s3));
shifts = (shifts << 8) | (uint16)((uint4)(shift2.s0), (uint4)(shift2.s1), (uint4)(shift2.s2), (uint4)(shift2.s3));
shifts = (shifts << 8) | (uint16)((uint4)(shift1.s0), (uint4)(shift1.s1), (uint4)(shift1.s2), (uint4)(shift1.s3));
shifts = (shifts << 8) | (uint16)((uint4)(shift0.s0), (uint4)(shift0.s1), (uint4)(shift0.s2), (uint4)(shift0.s3));
shifts <<= 3;
uint16 shifts = (uint16)((uint4)(shift3.s0), (uint4)(shift3.s1), (uint4)(shift3.s2), (uint4)(shift3.s3));
shifts = (shifts << (uint16)8) | (uint16)((uint4)(shift2.s0), (uint4)(shift2.s1), (uint4)(shift2.s2), (uint4)(shift2.s3));
shifts = (shifts << (uint16)8) | (uint16)((uint4)(shift1.s0), (uint4)(shift1.s1), (uint4)(shift1.s2), (uint4)(shift1.s3));
shifts = (shifts << (uint16)8) | (uint16)((uint4)(shift0.s0), (uint4)(shift0.s1), (uint4)(shift0.s2), (uint4)(shift0.s3)) << (uint16)3;
for (int iy = 0; iy < 4; iy++)
{
uchar16 res;
res = convert_uchar16(colorsFull >> (shifts & 0xFF));
shifts >>= 8;
//uchar4 t = convert_uchar4((ushort4)(color565.s0 >> 8, color565.s0 & 0xFF, color565.s1 >> 8, color565.s1 & 0xFF));
//res = (uchar16)(t, t, t, t);
//res = (uchar16)(frac, color32.even - color32.odd, (color32.even - color32.odd) / 2, (color32.even - color32.odd) / 2 - ((color32.even - color32.odd) / 8));
//res = (uchar16)(color32.even, color32.odd, frac, convert_uchar4(choice));
//res = convert_uchar16((uint16)(colorNoAlpha >> 24, colorNoAlpha >> 16, colorNoAlpha >> 8, colorNoAlpha));
//res = convert_uchar16((uint16)(colorAlpha >> 24, colorAlpha >> 16, colorAlpha >> 8, colorAlpha));
//res = convert_uchar16((uint16)(colors >> 24, colors >> 16, colors >> 8, colors));
//res = convert_uchar16(shifts & 0xFF);
//res = convert_uchar16((uint16)(shift0, shift1, shift2, shift3));
//res = (uchar16)(((x))) + (uchar16)(0, 0, 0, 0, 1, 1, 1, 1, 2, 2, 2, 2, 3, 3, 3, 3);
//res.lo = val; res.s8 = x >> 8; res.s9 = x; res.sA = (iy + y) >> 8; res.sB = y + iy; res.sC = width >> 8; res.sD = width; res.sE = srcOffset >> 8; res.sF = srcOffset;
res = convert_uchar16(colorsFull >> (shifts & (uint16)0xFF)) >> (uchar16)8;
vstore16(res, 0, dst);
dst += width * 4;
}
}
}
kernel void DecodeCMPR(global uchar *dst,