Create a stringify example, instead of loading the .cl file from disk, include it as a string.

The kernel in the .cl file is also compiled by the native C++ cpu compiler, when using MiniCL.
When you want to debug the kernel using MiniCL, and want to put breakpoints, it is best to:

1) enabled the define #define DEBUG_MINICL_KERNELS 1 in Bullet/src/BulletMultiThreaded/MiniCL.cpp
2) temporarily remove the stringify lines in the .cl kernel, because it prevents the debugger from finding the right line.

Demos/VectorAdd_OpenCL/AMD/CMakeLists.txt

@@ -35,3 +35,11 @@ IF (UNIX)
   TARGET_LINK_LIBRARIES(AppVectorAdd_AMD pthread)
 ENDIF(UNIX)
 
+IF (NOT INTERNAL_CREATE_DISTRIBUTABLE_MSVC_PROJECTFILES)
+	ADD_CUSTOM_COMMAND(
+		TARGET AppVectorAdd_AMD
+		POST_BUILD
+		COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/Demos/VectorAdd_OpenCL/VectorAddKernels.cl ${CMAKE_CURRENT_BINARY_DIR}
+	)
+ENDIF()
+				

Demos/VectorAdd_OpenCL/MiniCL/CMakeLists.txt

@@ -20,3 +20,12 @@ IF (UNIX)
   TARGET_LINK_LIBRARIES(AppVectorAdd_Mini pthread)
 ENDIF(UNIX)
 
+
+IF (NOT INTERNAL_CREATE_DISTRIBUTABLE_MSVC_PROJECTFILES)
+	ADD_CUSTOM_COMMAND(
+		TARGET AppVectorAdd_Mini
+		POST_BUILD
+		COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/Demos/VectorAdd_OpenCL/VectorAddKernels.cl ${CMAKE_CURRENT_BINARY_DIR}
+	)
+ENDIF()
+				

Demos/VectorAdd_OpenCL/MiniCL_VectorAdd.cpp

@@ -5,15 +5,17 @@
 ///Instead of #include <CL/cl.h> we include <MiniCL/cl.h>
 ///Apart from this include file, all other code should compile and work on OpenCL compliant implementation
 
-#ifdef USE_MINICL
 
-#include "MiniCL/cl.h"
+//#define LOAD_FROM_FILE
+
+#ifdef USE_MINICL
+	#include "MiniCL/cl.h"
 #else //USE_MINICL
-#ifdef __APPLE__
-#include <OpenCL/OpenCL.h>
-#else
-#include <CL/cl.h>
-#endif //__APPLE__
+	#ifdef __APPLE__
+		#include <OpenCL/OpenCL.h>
+	#else
+		#include <CL/cl.h>
+	#endif //__APPLE__
 #endif//USE_MINICL
 
 #include <stdio.h>
@@ -25,8 +27,13 @@
 size_t wgSize;
 
 
-
-
+#ifndef USE_MINICL
+#define MSTRINGIFY(A) #A
+char* stringifiedSourceCL = 
+#include "VectorAddKernels.cl"
+#else
+char* stringifiedSourceCL = "";
+#endif
 
 
 
@@ -224,8 +231,15 @@ int main(int argc, char **argv)
 	
     printf("loadProgSource (%s)...\n", cSourceFile); 
     const char* cPathAndName = cSourceFile;
+#ifdef LOAD_FROM_FILE
 	size_t szKernelLength;
     char* cSourceCL = loadProgSource(cPathAndName, "", &szKernelLength);
+#else
+	char* cSourceCL = stringifiedSourceCL;
+	size_t szKernelLength = strlen(stringifiedSourceCL);
+#endif //LOAD_FROM_FILE
+
+
 	
     // Create the program
     cpProgram = clCreateProgramWithSource(cxGPUContext, 1, (const char **)&cSourceCL, &szKernelLength, &ciErr1);
@@ -238,9 +252,9 @@ int main(int argc, char **argv)
 	
     // Build the program with 'mad' Optimization option
 #ifdef MAC
-	char* flags = "-cl-mad-enable -DMAC";
+	char* flags = "-cl-mad-enable -DMAC -DGUID_ARG";
 #else
-	const char* flags = "";//"-cl-mad-enable";
+	const char* flags = "-DGUID_ARG=";
 #endif
     ciErr1 = clBuildProgram(cpProgram, 0, NULL, flags, NULL, NULL);
     printf("clBuildProgram...\n"); 
@@ -312,20 +326,6 @@ int main(int argc, char **argv)
 		localWorkSize[1] = 1;
 		globalWorkSize[1] = 1;
 
-/*		size_t localWorkSize[2], globalWorkSize[2];
-		workgroupSize = workgroupSize < actualGlobalSize ? workgroupSize : actualGlobalSize;
-		int num_t = actualGlobalSize / workgroupSize;
-		int num_g = num_t * workgroupSize;
-		if(num_g < actualGlobalSize)
-		{
-			num_t++;
-		}
-		localWorkSize[0]  = workgroupSize;
-		globalWorkSize[0] = num_t * workgroupSize;
-		localWorkSize[1] = 1;
-		globalWorkSize[1] = 1;
-*/
-
 		// Copy input data from host to GPU and launch kernel 
 		ciErr1 |= clEnqueueNDRangeKernel(cqCommandQue, ckKernel, 1, NULL, globalThreads, localThreads, 0, NULL, NULL);
 
@@ -380,10 +380,17 @@ int main(int argc, char **argv)
 
 
 #ifdef USE_MINICL
+
 #include "MiniCL/cl_MiniCL_Defs.h"
+
 extern "C"
 {
+	///GUID_ARG is only used by MiniCL to pass in the guid used by its get_global_id implementation
+
+
+	#define MSTRINGIFY(A) A
 	#include "VectorAddKernels.cl"
+	#undef MSTRINGIFY
 }
 MINICL_REGISTER(VectorAdd)
-#endif//USE_MINICL
+#endif//USE_MINICL

Demos/VectorAdd_OpenCL/VectorAddKernels.cl

@@ -1,3 +1,16 @@
+
+#ifndef GUID_ARG
+#define GUID_ARG
+#endif
+
+
+#ifndef MSTRINGIFY
+#define MSTRINGIFY(A) A
+#endif
+
+
+MSTRINGIFY(
+
 /*
 Bullet Continuous Collision Detection and Physics Library, http://bulletphysics.org
 Copyright (C) 2006 - 2009 Sony Computer Entertainment Inc. 
@@ -13,10 +26,7 @@ subject to the following restrictions:
 3. This notice may not be removed or altered from any source distribution.
 */
 
-///GUID_ARG is only used by MiniCL to pass in the guid used by its get_global_id implementation
-#ifndef GUID_ARG
-#define GUID_ARG
-#endif
+
 
 ///////////////////////////////////////////////////
 // OpenCL Kernel Function for element by element vector addition
@@ -47,3 +57,4 @@ __kernel void VectorAdd(__global const float8* a, __global const float8* b, __gl
     c[get_global_id(0)] = f8Out;
 }
 
+);

src/BulletMultiThreaded/SequentialThreadSupport.h

@@ -85,6 +85,10 @@ public:
 	{
 		return 1;
 	}
+	virtual btBarrier*	createBarrier() { return 0;}
+
+	virtual btCriticalSection* createCriticalSection()  { return 0;};
+
 
 };