From 8bf91f735cb4626b0f5805a337cc227bdae6801d Mon Sep 17 00:00:00 2001 From: "erwin.coumans" Date: Fri, 25 Jun 2010 22:21:18 +0000 Subject: [PATCH] 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 | 8 +++ Demos/VectorAdd_OpenCL/MiniCL/CMakeLists.txt | 9 +++ Demos/VectorAdd_OpenCL/MiniCL_VectorAdd.cpp | 59 +++++++++++-------- Demos/VectorAdd_OpenCL/VectorAddKernels.cl | 19 ++++-- .../SequentialThreadSupport.h | 4 ++ 5 files changed, 69 insertions(+), 30 deletions(-) diff --git a/Demos/VectorAdd_OpenCL/AMD/CMakeLists.txt b/Demos/VectorAdd_OpenCL/AMD/CMakeLists.txt index 6371394ce..8b0f989ea 100644 --- a/Demos/VectorAdd_OpenCL/AMD/CMakeLists.txt +++ b/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() + \ No newline at end of file diff --git a/Demos/VectorAdd_OpenCL/MiniCL/CMakeLists.txt b/Demos/VectorAdd_OpenCL/MiniCL/CMakeLists.txt index 59740c7f0..7bea05fdc 100644 --- a/Demos/VectorAdd_OpenCL/MiniCL/CMakeLists.txt +++ b/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() + \ No newline at end of file diff --git a/Demos/VectorAdd_OpenCL/MiniCL_VectorAdd.cpp b/Demos/VectorAdd_OpenCL/MiniCL_VectorAdd.cpp index 6220140e8..92db86225 100644 --- a/Demos/VectorAdd_OpenCL/MiniCL_VectorAdd.cpp +++ b/Demos/VectorAdd_OpenCL/MiniCL_VectorAdd.cpp @@ -5,15 +5,17 @@ ///Instead of #include we include ///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 -#else -#include -#endif //__APPLE__ + #ifdef __APPLE__ + #include + #else + #include + #endif //__APPLE__ #endif//USE_MINICL #include @@ -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 \ No newline at end of file +#endif//USE_MINICL diff --git a/Demos/VectorAdd_OpenCL/VectorAddKernels.cl b/Demos/VectorAdd_OpenCL/VectorAddKernels.cl index 70d96a93d..e224eb6ff 100644 --- a/Demos/VectorAdd_OpenCL/VectorAddKernels.cl +++ b/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; } +); \ No newline at end of file diff --git a/src/BulletMultiThreaded/SequentialThreadSupport.h b/src/BulletMultiThreaded/SequentialThreadSupport.h index 4256ebd2a..514e54ee9 100644 --- a/src/BulletMultiThreaded/SequentialThreadSupport.h +++ b/src/BulletMultiThreaded/SequentialThreadSupport.h @@ -85,6 +85,10 @@ public: { return 1; } + virtual btBarrier* createBarrier() { return 0;} + + virtual btCriticalSection* createCriticalSection() { return 0;}; + };