OpenCL: More windows support and some unfinished work

git-svn-id: https://dolphin-emu.googlecode.com/svn/trunk@4380 8ced0084-cf51-0410-be5f-012b33b47a6e

Source/Core/Common/Src/OpenCL.cpp

@@ -27,7 +27,11 @@ namespace OpenCL {
     cl_command_queue g_cmdq = NULL;
 	#endif
 
+bool g_bInitialized = false;
+
 bool Initialize() {
+	if(g_bInitialized)
+		return true;
 
 #if defined(HAVE_OPENCL) && HAVE_OPENCL
 	if(g_context)
@@ -41,7 +45,7 @@ bool Initialize() {
     err = clGetDeviceIDs(NULL, gpu ? CL_DEVICE_TYPE_GPU : CL_DEVICE_TYPE_CPU, 1, &device_id, NULL);
     if (err != CL_SUCCESS)
     {
-        PanicAlert("Error: Failed to create a device group!\n");
+        PanicAlert("Error: Failed to create a device group!");
         return false;
     }
   
@@ -50,7 +54,7 @@ bool Initialize() {
     g_context = clCreateContext(0, 1, &device_id, NULL, NULL, &err);
     if (!g_context)
     {
-        PanicAlert("Error: Failed to create a compute context!\n");
+        PanicAlert("Error: Failed to create a compute context!");
         return false;
     }
 
@@ -59,17 +63,18 @@ bool Initialize() {
     g_cmdq = clCreateCommandQueue(g_context, device_id, 0, &err);
     if (!g_cmdq)
     {
-        PanicAlert("Error: Failed to create a command commands!\n");
+        PanicAlert("Error: Failed to create a command commands!");
         return false;
     }
-	printf("Initialized OpenCL fine!\n");
+	//PanicAlert("Initialized OpenCL fine!");
+	g_bInitialized = true;
 	return true;
 #else
 	return false;
 #endif
 }
 #if defined(HAVE_OPENCL) && HAVE_OPENCL
-cl_context GetInstance() {
+cl_context GetContext() {
 	return g_context;
 }
 
@@ -83,21 +88,18 @@ cl_program CompileProgram(const char *Kernel) {
 		program = clCreateProgramWithSource(OpenCL::g_context, 1, (const char **) & Kernel, NULL, &err);
 		if (!program)
 		{
-			printf("Error: Failed to create compute program!\n");
+			printf("Error: Failed to create compute program!");
 			return NULL;
 		}
 
 		// Build the program executable
 		//
 		err = clBuildProgram(program , 0, NULL, NULL, NULL, NULL);
-		if (err != CL_SUCCESS)
+		if(err != CL_SUCCESS) {
-		{
+			char *errors[16384] = {0};
-			size_t len;
+			err = clGetProgramBuildInfo(program, OpenCL::device_id, CL_PROGRAM_BUILD_LOG, sizeof(errors),
-			char buffer[2048];
+									    errors, NULL);
-
+			PanicAlert("Error log:\n%s\n", errors);
-			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;
 		}
 
@@ -112,7 +114,7 @@ cl_kernel CompileKernel(cl_program program, const char *Function)
 		cl_kernel kernel = clCreateKernel(program, Function, &err);
 		if (!kernel || err != CL_SUCCESS)
 		{
-			printf("Error: Failed to create compute kernel!\n");
+			PanicAlert("Error: Failed to create compute kernel!");
 			return NULL;
 		} 
 		return kernel;

Source/Core/Core/Src/HW/SystemTimers.cpp

@@ -238,11 +238,11 @@ void Init()
 		CPU_CORE_CLOCK = 729000000u;
 		SI_PERIOD = GetTicksPerSecond() / 60; // once a frame is good for controllers
 
-		if (!Core::GetStartupParameter().bDSPThread) {
+		//if (!Core::GetStartupParameter().bDSPThread) {
-			DSP_PERIOD = 12000;  // TO BE TWEAKED
+		//	DSP_PERIOD = 12000;  // TO BE TWEAKED
-		} else {
+		//} else {
 			DSP_PERIOD = (int)(GetTicksPerSecond() * 0.003f);
-		}
+		//}
 
 		// This is the biggest question mark.
 		AI_PERIOD = GetTicksPerSecond() / 80;
@@ -254,11 +254,11 @@ void Init()
 		CPU_CORE_CLOCK = 486000000;
 		SI_PERIOD = GetTicksPerSecond() / 60; // once a frame is good for controllers
 	
-		if (!Core::GetStartupParameter().bDSPThread) {
+		//if (!Core::GetStartupParameter().bDSPThread) {
-			DSP_PERIOD = 12000;  // TO BE TWEAKED
+		//	DSP_PERIOD = 12000;  // TO BE TWEAKED
-		} else {
+		//} else {
 			DSP_PERIOD = (int)(GetTicksPerSecond() * 0.005f);
-		}
+		//}
 
 		// This is the biggest question mark.
 		AI_PERIOD = GetTicksPerSecond() / 80;

Source/Core/VideoCommon/Src/OpenCL/OCLTextureDecoder.cpp

@@ -15,32 +15,50 @@
 // Official SVN repository and contact information can be found at
 // http://code.google.com/p/dolphin-emu/
 
-#include "TextureDecoder.h"
+#include "OCLTextureDecoder.h"
 
 #include "OpenCL.h"
-#include <CL/cl.h>
 
 #include <fcntl.h>
 #include <stdio.h>
 #include <stdlib.h>
 #include <string.h>
 #include <math.h>
-#include <unistd.h>
 #include <sys/types.h>
 #include <sys/stat.h>
     
-    
 struct sDecoders
 {
     cl_program program;                 // compute program
     cl_kernel kernel;                   // compute kernel
-    const char **cKernel;
+	cl_mem src, dst;                    // texture buffer memory objects
 };
 
+// XK's neatly aligned kernel
+const char *XKernel = "                                        \n\
+kernel void DecodeI8(global uchar *dst,                       \n\
+                     const global uchar *src, int width)      \n\
+{                                                             \n\
+    int x = get_global_id(0) * 8, y = get_global_id(1) * 4;   \n\
+    int srcOffset = 0;                                        \n\
+    for (int iy = 0; iy < 4; iy++)                            \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\
+		srcOffset += 8;                                       \n\
+    }                                                         \n\
+}\n";
+
 const char *Kernel = "         \n                               \
-__kernel void Decode(__global unsigned char *dst,               \n		\
+__kernel void DecodeI8(__global unsigned char *dst,               \n              \
                      const __global unsigned char *src,       \n                        \
-		     const __global int width, const __global int height)       \n          		\
+                     const __global int width)       \n                      \
 {                                    \n                         \
         int x = get_global_id(0) % width, y = get_global_id(0) / width;   \n            \
         if((y % 4) == 0 && (x % 8) == 0) \n \
@@ -61,81 +79,72 @@ __kernel void Decode(__global unsigned char *dst,               \n		\
         } \n \
 } \n \
 }\n";
-//						memcpy(dst + (y + iy)*width+x, src, 8);			
+
-const char *KernelOld = "         \n                     		\
+
-__kernel void Decode(__global uchar *dst,               \n		\
+sDecoders Decoders[] = { {NULL, NULL, NULL, NULL}, };
-		     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 defined(HAVE_OPENCL) && HAVE_OPENCL
+    cl_int err;
     if(!g_Inited)
     {
-		g_Inited = true;
+		if(!OpenCL::Initialize())
-
-#if defined(HAVE_OPENCL) && HAVE_OPENCL
-		// TODO: Switch this over to the OpenCl.h backend
-	    // Create the compute program from the source buffer
-		//
-		
-		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 PC_TEX_FMT_NONE;
-		}
 
-		// Build the program executable
-		//
-		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");
+		Decoders[0].program = OpenCL::CompileProgram(Kernel);
-			clGetProgramBuildInfo(Decoders[0].program , OpenCL::device_id, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len);
+		Decoders[0].kernel = OpenCL::CompileKernel(Decoders[0].program, "DecodeI8");
-			printf("%s\n", buffer);
-			exit(1);
-		}
 
-		// Create the compute kernel in the program we wish to run
+		g_Inited = true;
-		//
-		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);
 	}
-    
+    switch(texformat)
-#endif
-	}
-	    /*switch(texformat)
 	{
 		case GX_TF_I8:
 			{
-				int srcOffset = 0;
+				// TODO: Optimize
+				//PanicAlert("Really calling the OCL version");
+				Decoders[0].src = clCreateBuffer(OpenCL::GetContext(), CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, width * height * sizeof(u8), (void *)src, NULL);
+				Decoders[0].dst = clCreateBuffer(OpenCL::GetContext(), CL_MEM_WRITE_ONLY, width * height * sizeof(u8), NULL, NULL);
+
+				clSetKernelArg(Decoders[0].kernel, 0, sizeof(cl_mem), &Decoders[0].dst);
+				clSetKernelArg(Decoders[0].kernel, 1, sizeof(cl_mem), &Decoders[0].src);
+				clSetKernelArg(Decoders[0].kernel, 2, sizeof(cl_int), &width);
+
+				//size_t global[] = { width, height, 0 }, local[3] = {0}; // Later on we'll get the max work-group size and actually use them
+				size_t global = width * height, local;
+
+				clGetKernelWorkGroupInfo(Decoders[0].kernel, OpenCL::device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(local), &local, NULL);
+
+				err = clEnqueueNDRangeKernel(OpenCL::GetCommandQueue(), Decoders[0].kernel, 1 /* 2 */, NULL, &global, &local, 0, NULL, NULL);
+				if(err)
+					PanicAlert("Error queueing kernel");
+
+				clFinish(OpenCL::GetCommandQueue());
+				
+				clEnqueueReadBuffer(OpenCL::GetCommandQueue(), Decoders[0].dst, CL_TRUE, 0, width * height * sizeof(u8), dst, 0, NULL, NULL);
+
+				clReleaseMemObject(Decoders[0].src);
+				clReleaseMemObject(Decoders[0].dst);
+				
+				/*
 				for (int y = 0; y < height; y += 4)
 					for (int x = 0; x < width; x += 8)
-					for (int iy = 0; iy < 4; iy++, srcOffset += 8)
+						for (int iy = 0; iy < 4; iy++, src += 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;
 			}
-			break;
+
 		default:
 			return PC_TEX_FMT_NONE;
 	}
-		//return PC_TEX_FMT_NONE;*/
+#else
+	return PC_TEX_FMT_NONE;
+#endif
+	/* OLD CODE
     switch(texformat)
     {
         case GX_TF_I8:
@@ -211,7 +220,7 @@ PC_TexFormat TexDecoder_Decode_OpenCL(u8 *dst, const u8 *src, int width, int hei
             return PC_TEX_FMT_NONE;
     }
 	// TODO: clEnqueueNDRangeKernel
-
+*/
 
   /*  switch (texformat)
     {

Source/Core/VideoCommon/Src/SConscript

@@ -41,7 +41,7 @@ env_vcommon = env.Clone()
 
 if env_vcommon['HAVE_OPENCL']:
     files += [
-        'OpenCL/TextureDecoder.cpp',
+        'OpenCL/OCLTextureDecoder.cpp',
         ]
 
 env_vcommon.Append(CXXFLAGS = [ '-fPIC' ])

Source/Core/VideoCommon/Src/TextureDecoder.cpp

@@ -21,8 +21,10 @@
 
 #include "CPUDetect.h"
 #include "TextureDecoder.h"
+
+#include "OpenCL.h"
 #if defined(HAVE_OPENCL) && HAVE_OPENCL
-#include "OpenCL/TextureDecoder.h"
+#include "OpenCL/OCLTextureDecoder.h"
 #endif
 
 #include "LookUpTables.h"

Source/Core/VideoCommon/VideoCommon.vcproj

@@ -692,11 +692,11 @@
 				Name="OpenCL"
 				>
 				<File
-					RelativePath=".\Src\OpenCL\TextureDecoder.cpp"
+					RelativePath=".\Src\OpenCL\OCLTextureDecoder.cpp"
 					>
 				</File>
 				<File
-					RelativePath=".\Src\OpenCL\TextureDecoder.h"
+					RelativePath=".\Src\OpenCL\OCLTextureDecoder.h"
 					>
 				</File>
 			</Filter>