diff options
| author | XTra.KrazzY <XTra.KrazzY@gmail.com> | 2009-10-11 22:21:36 +0000 |
|---|---|---|
| committer | XTra.KrazzY <XTra.KrazzY@gmail.com> | 2009-10-11 22:21:36 +0000 |
| commit | 9da93462745b368698afccd7680a685ad603f1ce (patch) | |
| tree | c32c5171c8945522a321ce9042bacdda877a4f76 /Source/Core/VideoCommon/Src/OpenCL/OCLTextureDecoder.cpp | |
| parent | fc1c77b60dbc76b0e5d29836ba944bbc5042cbee (diff) | |
OpenCL: Kernels are now more general...
git-svn-id: https://dolphin-emu.googlecode.com/svn/trunk@4402 8ced0084-cf51-0410-be5f-012b33b47a6e
Diffstat (limited to 'Source/Core/VideoCommon/Src/OpenCL/OCLTextureDecoder.cpp')
| -rw-r--r-- | Source/Core/VideoCommon/Src/OpenCL/OCLTextureDecoder.cpp | 43 |
1 files changed, 35 insertions, 8 deletions
diff --git a/Source/Core/VideoCommon/Src/OpenCL/OCLTextureDecoder.cpp b/Source/Core/VideoCommon/Src/OpenCL/OCLTextureDecoder.cpp index a590f48073..5c97e3733b 100644 --- a/Source/Core/VideoCommon/Src/OpenCL/OCLTextureDecoder.cpp +++ b/Source/Core/VideoCommon/Src/OpenCL/OCLTextureDecoder.cpp @@ -29,7 +29,7 @@ struct sDecoders
{
- cl_program program; // compute program
+ const char name[256]; // kernel name
cl_kernel kernel; // compute kernel
};
@@ -86,11 +86,28 @@ kernel void DecodeIA4(global ushort *dst, \n\ srcOffset++; \n\
} \n\
} \n\
+ \n\
+kernel void DecodeDXT(global ulong *dst, \n\
+ const global ulong *src, int width) \n\
+{ // TODO: PLEASE NOTE THAT THIS CODE DOES NOT WORK \n\
+ int x = get_global_id(0) * 8, y = get_global_id(1) * 8; \n\
+ int srcOffset = ((x * 4) + (y * width)) / 8; \n\
+ for (int iy = 0; iy < 4; iy++) \n\
+ { \n\
+ vstore8(vload8(srcOffset, src), \n\
+ 0, dst + ((y + iy)*width + x)); \n\
+ srcOffset++; \n\
+ } \n\
+} \n\
";
-sDecoders Decoders[] = { {NULL, NULL},
- {NULL, NULL},
- {NULL, NULL},
+cl_program g_program;
+// NULL terminated set of kernels
+sDecoders Decoders[] = { {"DecodeI8", NULL},
+ {"DecodeIA4", NULL},
+ {"DecodeIA8", NULL},
+ {"DecodeDXT", NULL},
+ {"", NULL},
};
bool g_Inited = false;
@@ -109,11 +126,13 @@ PC_TexFormat TexDecoder_Decode_OpenCL(u8 *dst, const u8 *src, int width, int hei return PC_TEX_FMT_NONE;
- Decoders[0].program = OpenCL::CompileProgram(Kernel);
+ g_program = OpenCL::CompileProgram(Kernel);
- Decoders[0].kernel = OpenCL::CompileKernel(Decoders[0].program, "DecodeI8");
- Decoders[1].kernel = OpenCL::CompileKernel(Decoders[0].program, "DecodeIA4");
- Decoders[2].kernel = OpenCL::CompileKernel(Decoders[0].program, "DecodeIA8");
+ int i = 0;
+ while(strlen(Decoders[i].name) > 0) {
+ Decoders[i].kernel = OpenCL::CompileKernel(g_program, Decoders[i].name);
+ i++;
+ }
g_Inited = true;
}
@@ -136,12 +155,20 @@ PC_TexFormat TexDecoder_Decode_OpenCL(u8 *dst, const u8 *src, int width, int hei formatResult = PC_TEX_FMT_I8;
break;
case GX_TF_IA8:
+ //return PC_TEX_FMT_NONE; // <-- TODO: Fix IA8
kernelToRun = Decoders[2].kernel;
sizeOfSrc = sizeOfDst = sizeof(u16);
xSkip = 4;
ySkip = 4;
formatResult = PC_TEX_FMT_IA8;
break;
+ case GX_TF_CMPR:
+ return PC_TEX_FMT_NONE; // <-- TODO: Fix CMPR
+ kernelToRun = Decoders[3].kernel;
+ sizeOfSrc = sizeOfDst = sizeof(u32);
+ xSkip = 8;
+ ySkip = 8;
+ formatResult = PC_TEX_FMT_BGRA32;
default:
return PC_TEX_FMT_NONE;
}
|