10 files changed, 3176 insertions, 0 deletions
diff --git a/tests/bullet/src/MiniCL/CMakeLists.txt b/tests/bullet/src/MiniCL/CMakeLists.txt
new file mode 100644
index 00000000..bfaaf6b7
--- /dev/null
+++ b/tests/bullet/src/MiniCL/CMakeLists.txt
@@ -0,0 +1,44 @@
+#MiniCL provides a small subset of OpenCL
+
+INCLUDE_DIRECTORIES(
+	${BULLET_PHYSICS_SOURCE_DIR}/src
+	${VECTOR_MATH_INCLUDE}
+)
+
+ADD_LIBRARY(MiniCL
+		MiniCL.cpp
+		MiniCLTaskScheduler.cpp
+		MiniCLTaskScheduler.h
+		MiniCLTask/MiniCLTask.cpp
+		MiniCLTask/MiniCLTask.h
+		../MiniCL/cl.h
+		../MiniCL/cl_gl.h
+		../MiniCL/cl_platform.h
+		../MiniCL/cl_MiniCL_Defs.h
+)
+SET_TARGET_PROPERTIES(MiniCL PROPERTIES VERSION ${BULLET_VERSION})
+SET_TARGET_PROPERTIES(MiniCL PROPERTIES SOVERSION ${BULLET_VERSION})
+
+
+IF (BUILD_SHARED_LIBS)
+	TARGET_LINK_LIBRARIES(MiniCL BulletMultiThreaded BulletDynamics BulletCollision)
+ENDIF (BUILD_SHARED_LIBS)
+
+IF (INSTALL_LIBS)
+	IF (NOT INTERNAL_CREATE_DISTRIBUTABLE_MSVC_PROJECTFILES)
+		#INSTALL of other files requires CMake 2.6
+		IF (${CMAKE_MAJOR_VERSION}.${CMAKE_MINOR_VERSION} GREATER 2.5)
+#			IF(INSTALL_EXTRA_LIBS)
+				IF (APPLE AND BUILD_SHARED_LIBS AND FRAMEWORK)
+					INSTALL(TARGETS MiniCL DESTINATION .)
+				ELSE (APPLE AND BUILD_SHARED_LIBS AND FRAMEWORK)
+				INSTALL(TARGETS MiniCL DESTINATION lib${LIB_SUFFIX})
+					INSTALL(DIRECTORY
+${CMAKE_CURRENT_SOURCE_DIR} DESTINATION ${INCLUDE_INSTALL_DIR} FILES_MATCHING
+PATTERN "*.h"  PATTERN ".svn" EXCLUDE PATTERN "CMakeFiles" EXCLUDE)
+				ENDIF (APPLE AND BUILD_SHARED_LIBS AND FRAMEWORK)
+#			ENDIF (INSTALL_EXTRA_LIBS)
+		ENDIF (${CMAKE_MAJOR_VERSION}.${CMAKE_MINOR_VERSION} GREATER 2.5)
+	ENDIF (NOT INTERNAL_CREATE_DISTRIBUTABLE_MSVC_PROJECTFILES)
+ENDIF (INSTALL_LIBS)
+
diff --git a/tests/bullet/src/MiniCL/MiniCL.cpp b/tests/bullet/src/MiniCL/MiniCL.cpp
new file mode 100644
index 00000000..1e0823a1
--- /dev/null
+++ b/tests/bullet/src/MiniCL/MiniCL.cpp
@@ -0,0 +1,720 @@
+/*
+   Copyright (C) 2010 Sony Computer Entertainment Inc.
+   All rights reserved.
+
+This software is provided 'as-is', without any express or implied warranty.
+In no event will the authors be held liable for any damages arising from the use of this software.
+Permission is granted to anyone to use this software for any purpose, 
+including commercial applications, and to alter it and redistribute it freely, 
+subject to the following restrictions:
+
+1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required.
+2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software.
+3. This notice may not be removed or altered from any source distribution.
+
+*/
+
+
+#include "MiniCL/cl.h"
+#define __PHYSICS_COMMON_H__ 1
+#ifdef _WIN32
+#include "BulletMultiThreaded/Win32ThreadSupport.h"
+#endif
+
+#include "BulletMultiThreaded/SequentialThreadSupport.h"
+#include "MiniCLTaskScheduler.h"
+#include "MiniCLTask/MiniCLTask.h"
+#include "LinearMath/btMinMax.h"
+#include <stdio.h>
+
+//#define DEBUG_MINICL_KERNELS 1
+
+static char* spPlatformID = "MiniCL, SCEA";
+static char* spDriverVersion= "1.0";
+
+CL_API_ENTRY cl_int CL_API_CALL clGetPlatformIDs(
+	cl_uint           num_entries,
+    cl_platform_id *  platforms,
+    cl_uint *         num_platforms ) CL_API_SUFFIX__VERSION_1_0
+{
+	if(platforms != NULL)
+	{
+		if(num_entries <= 0)
+		{
+			return CL_INVALID_VALUE; 
+		}
+		*((char**)platforms) = spPlatformID;
+	}
+	if(num_platforms != NULL)
+	{
+		*num_platforms = 1;
+	}
+	return CL_SUCCESS;
+}
+
+
+CL_API_ENTRY cl_int CL_API_CALL clGetPlatformInfo(
+	cl_platform_id   platform, 
+	cl_platform_info param_name,
+	size_t           param_value_size, 
+	void *           param_value,
+	size_t *         param_value_size_ret) CL_API_SUFFIX__VERSION_1_0
+{
+	char* pId = (char*)platform;
+	if(strcmp(pId, spPlatformID))
+	{
+			return CL_INVALID_PLATFORM; 
+	}
+	switch(param_name)
+	{
+		case CL_PLATFORM_VENDOR	:
+			if(param_value_size < (strlen(spPlatformID) + 1))
+			{
+				return CL_INVALID_VALUE; 
+			}
+			strcpy((char*)param_value, spPlatformID);
+			if(param_value_size_ret != NULL)
+			{
+				*param_value_size_ret = strlen(spPlatformID) + 1;
+			}
+			break;
+		default : 
+			return CL_INVALID_VALUE; 
+	}
+	return CL_SUCCESS;
+}
+
+
+
+
+CL_API_ENTRY cl_int CL_API_CALL clGetDeviceInfo(
+	cl_device_id            device ,
+	cl_device_info          param_name ,
+	size_t                  param_value_size ,
+	void *                  param_value ,
+	size_t *                param_value_size_ret) CL_API_SUFFIX__VERSION_1_0
+{
+
+	switch (param_name)
+	{
+	case CL_DEVICE_NAME:
+		{
+			char deviceName[] = "MiniCL CPU";
+			unsigned int nameLen = strlen(deviceName)+1;
+			btAssert(param_value_size>strlen(deviceName));
+			if (nameLen < param_value_size)
+			{
+				const char* cpuName = "MiniCL CPU";
+				sprintf((char*)param_value,"%s",cpuName);
+			} else
+			{
+				printf("error: param_value_size should be at least %d, but it is %d\n",nameLen,param_value_size);
+				return CL_INVALID_VALUE; 
+			}
+			break;
+		}
+	case CL_DEVICE_TYPE:
+		{
+			if (param_value_size>=sizeof(cl_device_type))
+			{
+				cl_device_type* deviceType = (cl_device_type*)param_value;
+				*deviceType = CL_DEVICE_TYPE_CPU;
+			} else
+			{
+				printf("error: param_value_size should be at least %d\n",sizeof(cl_device_type));
+				return CL_INVALID_VALUE; 
+			}
+			break;
+		}
+	case CL_DEVICE_MAX_COMPUTE_UNITS:
+		{
+			if (param_value_size>=sizeof(cl_uint))
+			{
+				cl_uint* numUnits = (cl_uint*)param_value;
+				*numUnits= 4;
+			} else
+			{
+				printf("error: param_value_size should be at least %d\n",sizeof(cl_uint));
+				return CL_INVALID_VALUE; 
+			}
+
+			break;
+		}
+	case CL_DEVICE_MAX_WORK_ITEM_SIZES:
+		{
+			size_t workitem_size[3];
+
+			if (param_value_size>=sizeof(workitem_size))
+			{
+				size_t* workItemSize = (size_t*)param_value;
+				workItemSize[0] = 64;
+				workItemSize[1] = 24;
+				workItemSize[2] = 16;
+			} else
+			{
+				printf("error: param_value_size should be at least %d\n",sizeof(cl_uint));
+				return CL_INVALID_VALUE; 
+			}
+			break;
+		}
+	case CL_DEVICE_MAX_CLOCK_FREQUENCY:
+		{
+			 cl_uint* clock_frequency = (cl_uint*)param_value;
+			 *clock_frequency = 3*1024;
+			break;
+		}
+
+	case CL_DEVICE_VENDOR	:
+		{
+			if(param_value_size < (strlen(spPlatformID) + 1))
+			{
+				return CL_INVALID_VALUE; 
+			}
+			strcpy((char*)param_value, spPlatformID);
+			if(param_value_size_ret != NULL)
+			{
+				*param_value_size_ret = strlen(spPlatformID) + 1;
+			}
+			break;
+		}
+	case CL_DRIVER_VERSION:
+		{
+			if(param_value_size < (strlen(spDriverVersion) + 1))
+			{
+				return CL_INVALID_VALUE; 
+			}
+			strcpy((char*)param_value, spDriverVersion);
+			if(param_value_size_ret != NULL)
+			{
+				*param_value_size_ret = strlen(spDriverVersion) + 1;
+			}
+
+			break;
+		}
+	case CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS:
+		{
+			 cl_uint* maxDimensions = (cl_uint*)param_value;
+			 *maxDimensions = 1;
+			 break;
+		}
+		case CL_DEVICE_MAX_WORK_GROUP_SIZE:
+		{
+			 cl_uint* maxWorkGroupSize = (cl_uint*)param_value;
+			 *maxWorkGroupSize = 128;//1;
+			 break;
+		}
+		case CL_DEVICE_ADDRESS_BITS:
+		{
+			 cl_uint* addressBits = (cl_uint*)param_value;
+			 *addressBits= 32; //@todo: should this be 64 for 64bit builds?
+			 break;
+		}
+		case CL_DEVICE_MAX_MEM_ALLOC_SIZE:
+			{
+				cl_ulong* maxMemAlloc = (cl_ulong*)param_value;
+				*maxMemAlloc= 512*1024*1024; //this "should be enough for everyone" ?
+			 break;
+			}
+		case CL_DEVICE_GLOBAL_MEM_SIZE:
+			{
+				cl_ulong* maxMemAlloc = (cl_ulong*)param_value;
+				*maxMemAlloc= 1024*1024*1024; //this "should be enough for everyone" ?
+			 break;
+			}
+
+		case CL_DEVICE_ERROR_CORRECTION_SUPPORT:
+			{
+			cl_bool* error_correction_support = (cl_bool*)param_value;
+			*error_correction_support = CL_FALSE;
+			break;
+			}
+
+		case CL_DEVICE_LOCAL_MEM_TYPE:
+			{
+			cl_device_local_mem_type* local_mem_type = (cl_device_local_mem_type*)param_value;
+			*local_mem_type = CL_GLOBAL;
+			break;
+			}
+		case CL_DEVICE_LOCAL_MEM_SIZE:
+			{
+				cl_ulong* localmem = (cl_ulong*) param_value;
+				*localmem = 32*1024;
+				break;
+			}
+
+		case CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE:
+			{
+				cl_ulong* localmem = (cl_ulong*) param_value;
+				*localmem = 64*1024;
+				break;
+			}
+		case CL_DEVICE_QUEUE_PROPERTIES:
+			{
+				cl_command_queue_properties* queueProp = (cl_command_queue_properties*) param_value;
+				memset(queueProp,0,param_value_size);
+
+				break;
+			}
+		case CL_DEVICE_IMAGE_SUPPORT:
+			{
+				cl_bool* imageSupport = (cl_bool*) param_value;
+				*imageSupport = CL_FALSE;
+				break;
+			}
+
+		case CL_DEVICE_MAX_WRITE_IMAGE_ARGS:
+		case CL_DEVICE_MAX_READ_IMAGE_ARGS:
+			{
+				cl_uint* imageArgs = (cl_uint*) param_value;
+				*imageArgs = 0;
+				break;
+			}
+		case CL_DEVICE_IMAGE3D_MAX_DEPTH:
+		case CL_DEVICE_IMAGE3D_MAX_HEIGHT:
+		case CL_DEVICE_IMAGE2D_MAX_HEIGHT:
+		case CL_DEVICE_IMAGE3D_MAX_WIDTH:
+		case CL_DEVICE_IMAGE2D_MAX_WIDTH:
+			{
+				size_t* maxSize = (size_t*) param_value;
+				*maxSize = 0;
+				break;
+			}
+
+		case CL_DEVICE_EXTENSIONS:
+			{
+				char* extensions = (char*) param_value;
+				*extensions = 0;
+				break;
+			}
+
+		case CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE:
+		case CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT:
+		case CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG:
+		case CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT:
+		case CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT:
+		case CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR:
+			{
+				cl_uint* width  = (cl_uint*) param_value;
+				*width = 1;
+				break;
+			}
+			
+	default:
+		{
+			printf("error: unsupported param_name:%d\n",param_name);
+		}
+	}
+
+
+	return 0;
+}
+
+CL_API_ENTRY cl_int CL_API_CALL clReleaseMemObject(cl_mem /* memobj */) CL_API_SUFFIX__VERSION_1_0
+{
+	return 0;
+}
+
+
+
+CL_API_ENTRY cl_int CL_API_CALL clReleaseCommandQueue(cl_command_queue /* command_queue */) CL_API_SUFFIX__VERSION_1_0
+{
+	return 0;
+}
+
+CL_API_ENTRY cl_int CL_API_CALL clReleaseProgram(cl_program /* program */) CL_API_SUFFIX__VERSION_1_0
+{
+	return 0;
+}
+
+CL_API_ENTRY cl_int CL_API_CALL clReleaseKernel(cl_kernel   /* kernel */) CL_API_SUFFIX__VERSION_1_0
+{
+	return 0;
+}
+
+
+// Enqueued Commands APIs
+CL_API_ENTRY cl_int CL_API_CALL clEnqueueReadBuffer(cl_command_queue     command_queue ,
+                    cl_mem               buffer ,
+                    cl_bool             /* blocking_read */,
+                    size_t               offset ,
+                    size_t               cb , 
+                    void *               ptr ,
+                    cl_uint             /* num_events_in_wait_list */,
+                    const cl_event *    /* event_wait_list */,
+                    cl_event *          /* event */) CL_API_SUFFIX__VERSION_1_0
+{
+	MiniCLTaskScheduler* scheduler = (MiniCLTaskScheduler*) command_queue;
+
+	///wait for all work items to be completed
+	scheduler->flush();
+
+	memcpy(ptr,(char*)buffer + offset,cb);
+	return 0;
+}
+
+
+CL_API_ENTRY cl_int clGetProgramBuildInfo(cl_program            /* program */,
+                      cl_device_id          /* device */,
+                      cl_program_build_info /* param_name */,
+                      size_t                /* param_value_size */,
+                      void *                /* param_value */,
+                      size_t *              /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0
+{
+
+	return 0;
+}
+
+
+// Program Object APIs
+CL_API_ENTRY cl_program
+clCreateProgramWithSource(cl_context         context ,
+                          cl_uint           /* count */,
+                          const char **     /* strings */,
+                          const size_t *    /* lengths */,
+                          cl_int *          errcode_ret ) CL_API_SUFFIX__VERSION_1_0
+{
+	*errcode_ret = CL_SUCCESS;
+	return (cl_program)context;
+}
+
+CL_API_ENTRY cl_int CL_API_CALL clEnqueueWriteBuffer(cl_command_queue     command_queue ,
+                    cl_mem               buffer ,
+                    cl_bool             /* blocking_read */,
+                    size_t              offset,
+                    size_t               cb , 
+                    const void *         ptr ,
+                    cl_uint             /* num_events_in_wait_list */,
+                    const cl_event *    /* event_wait_list */,
+                    cl_event *          /* event */) CL_API_SUFFIX__VERSION_1_0
+{
+	MiniCLTaskScheduler* scheduler = (MiniCLTaskScheduler*) command_queue;
+
+	///wait for all work items to be completed
+	scheduler->flush();
+
+	memcpy((char*)buffer + offset, ptr,cb);
+	return 0;
+}
+
+CL_API_ENTRY cl_int CL_API_CALL clFlush(cl_command_queue  command_queue)
+{
+	MiniCLTaskScheduler* scheduler = (MiniCLTaskScheduler*) command_queue;
+	///wait for all work items to be completed
+	scheduler->flush();
+	return 0;
+}
+
+
+CL_API_ENTRY cl_int CL_API_CALL clEnqueueNDRangeKernel(cl_command_queue /* command_queue */,
+                       cl_kernel         clKernel ,
+                       cl_uint           work_dim ,
+                       const size_t *   /* global_work_offset */,
+                       const size_t *    global_work_size ,
+                       const size_t *   /* local_work_size */,
+                       cl_uint          /* num_events_in_wait_list */,
+                       const cl_event * /* event_wait_list */,
+                       cl_event *       /* event */) CL_API_SUFFIX__VERSION_1_0
+{
+
+	
+	MiniCLKernel* kernel = (MiniCLKernel*) clKernel;
+	for (unsigned int ii=0;ii<work_dim;ii++)
+	{
+		int maxTask = kernel->m_scheduler->getMaxNumOutstandingTasks();
+		int numWorkItems = global_work_size[ii];
+
+//		//at minimum 64 work items per task
+//		int numWorkItemsPerTask = btMax(64,numWorkItems / maxTask);
+		int numWorkItemsPerTask = numWorkItems / maxTask;
+		if (!numWorkItemsPerTask) numWorkItemsPerTask = 1;
+
+		for (int t=0;t<numWorkItems;)
+		{
+			//Performance Hint: tweak this number during benchmarking
+			int endIndex = (t+numWorkItemsPerTask) < numWorkItems ? t+numWorkItemsPerTask : numWorkItems;
+			kernel->m_scheduler->issueTask(t, endIndex, kernel);
+			t = endIndex;
+		}
+	}
+/*
+
+	void* bla = 0;
+
+	scheduler->issueTask(bla,2,3);
+	scheduler->flush();
+
+	*/
+
+	return 0;
+}
+
+#define LOCAL_BUF_SIZE 32768
+static int sLocalMemBuf[LOCAL_BUF_SIZE * 4 + 16];
+static int* spLocalBufCurr = NULL;
+static int sLocalBufUsed = LOCAL_BUF_SIZE; // so it will be reset at the first call
+static void* localBufMalloc(int size)
+{
+	int size16 = (size + 15) >> 4; // in 16-byte units
+	if((sLocalBufUsed + size16) > LOCAL_BUF_SIZE)
+	{ // reset
+		spLocalBufCurr = sLocalMemBuf;
+		while((unsigned long)spLocalBufCurr & 0x0F) spLocalBufCurr++; // align to 16 bytes
+		sLocalBufUsed = 0;
+	}
+	void* ret = spLocalBufCurr;
+	spLocalBufCurr += size16 * 4;
+	sLocalBufUsed += size;
+	return ret;
+}
+
+
+
+CL_API_ENTRY cl_int CL_API_CALL clSetKernelArg(cl_kernel    clKernel ,
+               cl_uint      arg_index ,
+               size_t       arg_size ,
+               const void *  arg_value ) CL_API_SUFFIX__VERSION_1_0
+{
+	MiniCLKernel* kernel = (MiniCLKernel* ) clKernel;
+	btAssert(arg_size <= MINICL_MAX_ARGLENGTH);
+	if (arg_index>MINI_CL_MAX_ARG)
+	{
+		printf("error: clSetKernelArg arg_index (%d) exceeds %d\n",arg_index,MINI_CL_MAX_ARG);
+	} else
+	{
+		if (arg_size>MINICL_MAX_ARGLENGTH)
+		//if (arg_size != MINICL_MAX_ARGLENGTH)
+		{
+			printf("error: clSetKernelArg argdata too large: %d (maximum is %d)\n",arg_size,MINICL_MAX_ARGLENGTH);
+		} 
+		else
+		{
+			if(arg_value == NULL)
+			{	// this is only for __local memory qualifier
+				void* ptr = localBufMalloc(arg_size);
+				kernel->m_argData[arg_index] = ptr;
+			}
+			else
+			{
+				memcpy(&(kernel->m_argData[arg_index]), arg_value, arg_size);
+			}
+			kernel->m_argSizes[arg_index] = arg_size;
+			if(arg_index >= kernel->m_numArgs)
+			{
+				kernel->m_numArgs = arg_index + 1;
+				kernel->updateLauncher();
+			}
+		}
+	}
+	return 0;
+}
+
+// Kernel Object APIs
+CL_API_ENTRY cl_kernel CL_API_CALL clCreateKernel(cl_program       program ,
+               const char *     kernel_name ,
+               cl_int *         errcode_ret ) CL_API_SUFFIX__VERSION_1_0
+{
+	MiniCLTaskScheduler* scheduler = (MiniCLTaskScheduler*) program;
+	MiniCLKernel* kernel = new MiniCLKernel();
+	int nameLen = strlen(kernel_name);
+	if(nameLen >= MINI_CL_MAX_KERNEL_NAME)
+	{
+		*errcode_ret = CL_INVALID_KERNEL_NAME;
+		return NULL;
+	}
+	strcpy(kernel->m_name, kernel_name);
+	kernel->m_numArgs = 0;
+
+	//kernel->m_kernelProgramCommandId = scheduler->findProgramCommandIdByName(kernel_name);
+	//if (kernel->m_kernelProgramCommandId>=0)
+	//{
+	//	*errcode_ret = CL_SUCCESS;
+	//} else
+	//{
+	//	*errcode_ret = CL_INVALID_KERNEL_NAME;
+	//}
+	kernel->m_scheduler = scheduler;
+	if(kernel->registerSelf() == NULL)
+	{
+		*errcode_ret = CL_INVALID_KERNEL_NAME;
+		return NULL;
+	}
+	else
+	{
+		*errcode_ret = CL_SUCCESS;
+	}
+
+	return (cl_kernel)kernel;
+
+}
+
+
+CL_API_ENTRY cl_int CL_API_CALL clBuildProgram(cl_program           /* program */,
+               cl_uint              /* num_devices */,
+               const cl_device_id * /* device_list */,
+               const char *         /* options */, 
+               void (*pfn_notify)(cl_program /* program */, void * /* user_data */),
+               void *               /* user_data */) CL_API_SUFFIX__VERSION_1_0
+{
+	return CL_SUCCESS;
+}
+
+CL_API_ENTRY cl_program CL_API_CALL clCreateProgramWithBinary(cl_context                     context ,
+                          cl_uint                        /* num_devices */,
+                          const cl_device_id *           /* device_list */,
+                          const size_t *                 /* lengths */,
+                          const unsigned char **         /* binaries */,
+                          cl_int *                       /* binary_status */,
+                          cl_int *                       /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0
+{
+	return (cl_program)context;
+}
+
+
+// Memory Object APIs
+CL_API_ENTRY cl_mem CL_API_CALL clCreateBuffer(cl_context   /* context */,
+               cl_mem_flags flags ,
+               size_t       size,
+               void *       host_ptr ,
+               cl_int *     errcode_ret ) CL_API_SUFFIX__VERSION_1_0
+{
+	cl_mem buf = (cl_mem)malloc(size);
+	if ((flags&CL_MEM_COPY_HOST_PTR) && host_ptr)
+	{
+		memcpy(buf,host_ptr,size);
+	}
+	*errcode_ret = 0;
+	return buf;
+}
+
+// Command Queue APIs
+CL_API_ENTRY cl_command_queue CL_API_CALL clCreateCommandQueue(cl_context                      context , 
+                     cl_device_id                   /* device */, 
+                     cl_command_queue_properties    /* properties */,
+                     cl_int *                        errcode_ret ) CL_API_SUFFIX__VERSION_1_0
+{
+	*errcode_ret = 0;
+	return (cl_command_queue) context;
+}
+
+extern CL_API_ENTRY cl_int CL_API_CALL clGetContextInfo(cl_context         /* context */, 
+                 cl_context_info    param_name , 
+                 size_t             param_value_size , 
+                 void *             param_value, 
+                 size_t *           param_value_size_ret ) CL_API_SUFFIX__VERSION_1_0
+{
+
+	switch (param_name)
+	{
+	case CL_CONTEXT_DEVICES:
+		{
+			if (!param_value_size)
+			{
+				*param_value_size_ret = 13;
+			} else
+			{
+				const char* testName = "MiniCL_Test.";
+				sprintf((char*)param_value,"%s",testName);
+			}
+			break;
+		};
+	default:
+		{
+			printf("unsupported\n");
+		}
+	}
+	
+	return 0;
+}
+
+CL_API_ENTRY cl_context CL_API_CALL clCreateContextFromType(cl_context_properties * /* properties */,
+                        cl_device_type           device_type ,
+                        void (*pfn_notify)(const char *, const void *, size_t, void *) /* pfn_notify */,
+                        void *                  /* user_data */,
+                        cl_int *                 errcode_ret ) CL_API_SUFFIX__VERSION_1_0
+{
+	int maxNumOutstandingTasks = 4;
+//	int maxNumOutstandingTasks = 2;
+//	int maxNumOutstandingTasks = 1;
+	gMiniCLNumOutstandingTasks = maxNumOutstandingTasks;
+	const int maxNumOfThreadSupports = 8;
+	static int sUniqueThreadSupportIndex = 0;
+	static char* sUniqueThreadSupportName[maxNumOfThreadSupports] = 
+	{
+		"MiniCL_0", "MiniCL_1", "MiniCL_2", "MiniCL_3", "MiniCL_4", "MiniCL_5", "MiniCL_6", "MiniCL_7" 
+	};
+
+	btThreadSupportInterface* threadSupport = 0;
+
+	if (device_type==CL_DEVICE_TYPE_DEBUG)
+	{
+		SequentialThreadSupport::SequentialThreadConstructionInfo stc("MiniCL",processMiniCLTask,createMiniCLLocalStoreMemory);
+		threadSupport = new SequentialThreadSupport(stc);
+	} else
+	{
+
+#if _WIN32
+	btAssert(sUniqueThreadSupportIndex < maxNumOfThreadSupports);
+	threadSupport = new Win32ThreadSupport(Win32ThreadSupport::Win32ThreadConstructionInfo(
+//								"MiniCL",
+								sUniqueThreadSupportName[sUniqueThreadSupportIndex++],
+								processMiniCLTask, //processCollisionTask,
+								createMiniCLLocalStoreMemory,//createCollisionLocalStoreMemory,
+								maxNumOutstandingTasks));
+#else
+	///todo: add posix thread support for other platforms
+	SequentialThreadSupport::SequentialThreadConstructionInfo stc("MiniCL",processMiniCLTask,createMiniCLLocalStoreMemory);
+	threadSupport = new SequentialThreadSupport(stc);
+#endif
+
+	}
+	
+	
+	MiniCLTaskScheduler* scheduler = new MiniCLTaskScheduler(threadSupport,maxNumOutstandingTasks);
+
+	*errcode_ret = 0;
+	return (cl_context)scheduler;
+}
+
+CL_API_ENTRY cl_int CL_API_CALL clReleaseContext(cl_context  context ) CL_API_SUFFIX__VERSION_1_0
+{
+
+	MiniCLTaskScheduler* scheduler = (MiniCLTaskScheduler*) context;
+	
+	btThreadSupportInterface* threadSupport = scheduler->getThreadSupportInterface();
+	delete scheduler;
+	delete threadSupport;
+	
+	return 0;
+}
+extern CL_API_ENTRY cl_int CL_API_CALL
+clFinish(cl_command_queue command_queue ) CL_API_SUFFIX__VERSION_1_0
+{
+	MiniCLTaskScheduler* scheduler = (MiniCLTaskScheduler*) command_queue;
+	///wait for all work items to be completed
+	scheduler->flush();
+	return CL_SUCCESS;
+}
+
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+clGetKernelWorkGroupInfo(cl_kernel                   kernel ,
+                         cl_device_id               /* device */,
+                         cl_kernel_work_group_info  wgi/* param_name */,
+                         size_t   sz                  /* param_value_size */,
+                         void *     ptr                /* param_value */,
+                         size_t *                   /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0
+{
+	if((wgi == CL_KERNEL_WORK_GROUP_SIZE)
+	 &&(sz == sizeof(size_t))
+	 &&(ptr != NULL))
+	{
+		MiniCLKernel* miniCLKernel = (MiniCLKernel*)kernel;
+		MiniCLTaskScheduler* scheduler = miniCLKernel->m_scheduler;
+		*((size_t*)ptr) = scheduler->getMaxNumOutstandingTasks();
+		return CL_SUCCESS;
+	}
+	else
+	{
+		return CL_INVALID_VALUE;
+	}
+}
diff --git a/tests/bullet/src/MiniCL/MiniCLTask/MiniCLTask.cpp b/tests/bullet/src/MiniCL/MiniCLTask/MiniCLTask.cpp
new file mode 100644
index 00000000..a56e96a0
--- /dev/null
+++ b/tests/bullet/src/MiniCL/MiniCLTask/MiniCLTask.cpp
@@ -0,0 +1,74 @@
+/*
+Bullet Continuous Collision Detection and Physics Library, Copyright (c) 2007 Erwin Coumans
+
+This software is provided 'as-is', without any express or implied warranty.
+In no event will the authors be held liable for any damages arising from the use of this software.
+Permission is granted to anyone to use this software for any purpose, 
+including commercial applications, and to alter it and redistribute it freely, 
+subject to the following restrictions:
+
+1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required.
+2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software.
+3. This notice may not be removed or altered from any source distribution.
+
+*/
+
+
+#include "MiniCLTask.h"
+#include "BulletMultiThreaded/PlatformDefinitions.h"
+#include "BulletMultiThreaded/SpuFakeDma.h"
+#include "LinearMath/btMinMax.h"
+#include "MiniCLTask.h"
+#include "MiniCL/MiniCLTaskScheduler.h"
+
+
+#ifdef __SPU__
+#include <spu_printf.h>
+#else
+#include <stdio.h>
+#define spu_printf printf
+#endif
+
+int gMiniCLNumOutstandingTasks = 0;
+
+struct MiniCLTask_LocalStoreMemory
+{
+	
+};
+
+
+//-- MAIN METHOD
+void processMiniCLTask(void* userPtr, void* lsMemory)
+{
+	//	BT_PROFILE("processSampleTask");
+
+	MiniCLTask_LocalStoreMemory* localMemory = (MiniCLTask_LocalStoreMemory*)lsMemory;
+
+	MiniCLTaskDesc* taskDescPtr = (MiniCLTaskDesc*)userPtr;
+	MiniCLTaskDesc& taskDesc = *taskDescPtr;
+
+	for (unsigned int i=taskDesc.m_firstWorkUnit;i<taskDesc.m_lastWorkUnit;i++)
+	{
+		taskDesc.m_kernel->m_launcher(&taskDesc, i);
+	}
+
+//	printf("Compute Unit[%d] executed kernel %d work items [%d..%d)\n",taskDesc.m_taskId,taskDesc.m_kernelProgramId,taskDesc.m_firstWorkUnit,taskDesc.m_lastWorkUnit);
+	
+}
+
+
+#if defined(__CELLOS_LV2__) || defined (LIBSPE2)
+
+ATTRIBUTE_ALIGNED16(MiniCLTask_LocalStoreMemory	gLocalStoreMemory);
+
+void* createMiniCLLocalStoreMemory()
+{
+	return &gLocalStoreMemory;
+}
+#else
+void* createMiniCLLocalStoreMemory()
+{
+	return new MiniCLTask_LocalStoreMemory;
+};
+
+#endif
diff --git a/tests/bullet/src/MiniCL/MiniCLTask/MiniCLTask.h b/tests/bullet/src/MiniCL/MiniCLTask/MiniCLTask.h
new file mode 100644
index 00000000..7e78be08
--- /dev/null
+++ b/tests/bullet/src/MiniCL/MiniCLTask/MiniCLTask.h
@@ -0,0 +1,62 @@
+/*
+Bullet Continuous Collision Detection and Physics Library, Copyright (c) 2007 Erwin Coumans
+
+This software is provided 'as-is', without any express or implied warranty.
+In no event will the authors be held liable for any damages arising from the use of this software.
+Permission is granted to anyone to use this software for any purpose, 
+including commercial applications, and to alter it and redistribute it freely, 
+subject to the following restrictions:
+
+1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required.
+2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software.
+3. This notice may not be removed or altered from any source distribution.
+
+*/
+
+#ifndef MINICL__TASK_H
+#define MINICL__TASK_H
+
+#include "BulletMultiThreaded/PlatformDefinitions.h"
+#include "LinearMath/btScalar.h"
+
+#include "LinearMath/btAlignedAllocator.h"
+
+
+#define MINICL_MAX_ARGLENGTH (sizeof(void*))
+#define MINI_CL_MAX_ARG 16
+#define MINI_CL_MAX_KERNEL_NAME 256
+
+struct MiniCLKernel;
+
+ATTRIBUTE_ALIGNED16(struct) MiniCLTaskDesc
+{
+	BT_DECLARE_ALIGNED_ALLOCATOR();
+
+	MiniCLTaskDesc()
+	{
+		for (int i=0;i<MINI_CL_MAX_ARG;i++)
+		{
+			m_argSizes[i]=0;
+		}
+	}
+
+	uint32_t		m_taskId;
+
+	uint32_t		m_firstWorkUnit;
+	uint32_t		m_lastWorkUnit;
+
+	MiniCLKernel*	m_kernel;
+
+	void*			m_argData[MINI_CL_MAX_ARG];
+	int				m_argSizes[MINI_CL_MAX_ARG];
+};
+
+extern "C" int gMiniCLNumOutstandingTasks;
+
+
+void	processMiniCLTask(void* userPtr, void* lsMemory);
+void*	createMiniCLLocalStoreMemory();
+
+
+#endif //MINICL__TASK_H
+
diff --git a/tests/bullet/src/MiniCL/MiniCLTaskScheduler.cpp b/tests/bullet/src/MiniCL/MiniCLTaskScheduler.cpp
new file mode 100644
index 00000000..7adee88d
--- /dev/null
+++ b/tests/bullet/src/MiniCL/MiniCLTaskScheduler.cpp
@@ -0,0 +1,519 @@
+/*
+Bullet Continuous Collision Detection and Physics Library
+Copyright (c) 2003-2007 Erwin Coumans  http://bulletphysics.com
+
+This software is provided 'as-is', without any express or implied warranty.
+In no event will the authors be held liable for any damages arising from the use of this software.
+Permission is granted to anyone to use this software for any purpose, 
+including commercial applications, and to alter it and redistribute it freely, 
+subject to the following restrictions:
+
+1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required.
+2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software.
+3. This notice may not be removed or altered from any source distribution.
+*/
+
+//#define __CELLOS_LV2__ 1
+#define __BT_SKIP_UINT64_H 1
+
+#define USE_SAMPLE_PROCESS 1
+#ifdef USE_SAMPLE_PROCESS
+
+
+#include "MiniCLTaskScheduler.h"
+#include <stdio.h>
+
+#ifdef __SPU__
+
+
+
+void	SampleThreadFunc(void* userPtr,void* lsMemory)
+{
+	//do nothing
+	printf("hello world\n");
+}
+
+
+void*	SamplelsMemoryFunc()
+{
+	//don't create local store memory, just return 0
+	return 0;
+}
+
+
+#else
+
+
+#include "BulletMultiThreaded/btThreadSupportInterface.h"
+
+//#	include "SPUAssert.h"
+#include <string.h>
+
+#include "MiniCL/cl_platform.h"
+
+extern "C" {
+	extern char SPU_SAMPLE_ELF_SYMBOL[];
+}
+
+
+MiniCLTaskScheduler::MiniCLTaskScheduler(btThreadSupportInterface*	threadInterface,  int maxNumOutstandingTasks)
+:m_threadInterface(threadInterface),
+m_maxNumOutstandingTasks(maxNumOutstandingTasks)
+{
+
+	m_taskBusy.resize(m_maxNumOutstandingTasks);
+	m_spuSampleTaskDesc.resize(m_maxNumOutstandingTasks);
+
+	m_kernels.resize(0);
+
+	for (int i = 0; i < m_maxNumOutstandingTasks; i++)
+	{
+		m_taskBusy[i] = false;
+	}
+	m_numBusyTasks = 0;
+	m_currentTask = 0;
+
+	m_initialized = false;
+
+	m_threadInterface->startSPU();
+
+
+}
+
+MiniCLTaskScheduler::~MiniCLTaskScheduler()
+{
+	m_threadInterface->stopSPU();
+	
+}
+
+
+
+void	MiniCLTaskScheduler::initialize()
+{
+#ifdef DEBUG_SPU_TASK_SCHEDULING
+	printf("MiniCLTaskScheduler::initialize()\n");
+#endif //DEBUG_SPU_TASK_SCHEDULING
+	
+	for (int i = 0; i < m_maxNumOutstandingTasks; i++)
+	{
+		m_taskBusy[i] = false;
+	}
+	m_numBusyTasks = 0;
+	m_currentTask = 0;
+	m_initialized = true;
+
+}
+
+
+void MiniCLTaskScheduler::issueTask(int firstWorkUnit, int lastWorkUnit, MiniCLKernel* kernel)
+{
+
+#ifdef DEBUG_SPU_TASK_SCHEDULING
+	printf("MiniCLTaskScheduler::issueTask (m_currentTask= %d\)n", m_currentTask);
+#endif //DEBUG_SPU_TASK_SCHEDULING
+
+	m_taskBusy[m_currentTask] = true;
+	m_numBusyTasks++;
+
+	MiniCLTaskDesc& taskDesc = m_spuSampleTaskDesc[m_currentTask];
+	{
+		// send task description in event message
+		taskDesc.m_firstWorkUnit = firstWorkUnit;
+		taskDesc.m_lastWorkUnit = lastWorkUnit;
+		taskDesc.m_kernel = kernel;
+		//some bookkeeping to recognize finished tasks
+		taskDesc.m_taskId = m_currentTask;
+		
+//		for (int i=0;i<MINI_CL_MAX_ARG;i++)
+		for (unsigned int i=0; i < kernel->m_numArgs; i++)
+		{
+			taskDesc.m_argSizes[i] = kernel->m_argSizes[i];
+			if (taskDesc.m_argSizes[i])
+			{
+				taskDesc.m_argData[i] = kernel->m_argData[i];
+//				memcpy(&taskDesc.m_argData[i],&argData[MINICL_MAX_ARGLENGTH*i],taskDesc.m_argSizes[i]);
+			}
+		}
+	}
+