summaryrefslogtreecommitdiff
path: root/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives
diff options
context:
space:
mode:
authorOussama <o.boukhelf@gmail.com>2019-01-03 14:26:51 +0100
committerRĂ©mi Verschelde <rverschelde@gmail.com>2019-01-07 12:30:35 +0100
commit22b7c9dfa80d0f7abca40f061865c2ab3c136a74 (patch)
tree311cd3f22b012329160f9d43810aea429994af48 /thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives
parenta6722cf36251ddcb538e6ebed9fa4950342b68ba (diff)
Update Bullet to the latest commit 126b676
Diffstat (limited to 'thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives')
-rw-r--r--thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3BoundSearchCL.cpp162
-rw-r--r--thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3BoundSearchCL.h55
-rw-r--r--thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3BufferInfoCL.h7
-rw-r--r--thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3FillCL.cpp83
-rw-r--r--thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3FillCL.h63
-rw-r--r--thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3LauncherCL.cpp376
-rw-r--r--thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3LauncherCL.h165
-rw-r--r--thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3OpenCLArray.h174
-rw-r--r--thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3PrefixScanCL.cpp84
-rw-r--r--thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3PrefixScanCL.h14
-rw-r--r--thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3PrefixScanFloat4CL.cpp86
-rw-r--r--thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3PrefixScanFloat4CL.h12
-rw-r--r--thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3RadixSort32CL.cpp730
-rw-r--r--thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3RadixSort32CL.h103
-rw-r--r--thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/BoundSearchKernelsCL.h171
-rw-r--r--thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/CopyKernelsCL.h261
-rw-r--r--thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/FillKernelsCL.h179
-rw-r--r--thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/PrefixScanKernelsCL.h255
-rw-r--r--thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/PrefixScanKernelsFloat4CL.h255
-rw-r--r--thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/RadixSort32KernelsCL.h1817
20 files changed, 2449 insertions, 2603 deletions
diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3BoundSearchCL.cpp b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3BoundSearchCL.cpp
index a4980f71e1..c0e11bfb26 100644
--- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3BoundSearchCL.cpp
+++ b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3BoundSearchCL.cpp
@@ -19,149 +19,139 @@ subject to the following restrictions:
#define KERNEL1 "SearchSortDataUpperKernel"
#define KERNEL2 "SubtractKernel"
-
#include "b3BoundSearchCL.h"
#include "Bullet3OpenCL/Initialize/b3OpenCLUtils.h"
#include "b3LauncherCL.h"
#include "kernels/BoundSearchKernelsCL.h"
b3BoundSearchCL::b3BoundSearchCL(cl_context ctx, cl_device_id device, cl_command_queue queue, int maxSize)
- :m_context(ctx),
- m_device(device),
- m_queue(queue)
+ : m_context(ctx),
+ m_device(device),
+ m_queue(queue)
{
-
const char* additionalMacros = "";
//const char* srcFileNameForCaching="";
cl_int pErrNum;
const char* kernelSource = boundSearchKernelsCL;
- cl_program boundSearchProg = b3OpenCLUtils::compileCLProgramFromString( ctx, device, kernelSource, &pErrNum,additionalMacros, BOUNDSEARCH_PATH);
+ cl_program boundSearchProg = b3OpenCLUtils::compileCLProgramFromString(ctx, device, kernelSource, &pErrNum, additionalMacros, BOUNDSEARCH_PATH);
b3Assert(boundSearchProg);
- m_lowerSortDataKernel = b3OpenCLUtils::compileCLKernelFromString( ctx, device, kernelSource, "SearchSortDataLowerKernel", &pErrNum, boundSearchProg,additionalMacros );
- b3Assert(m_lowerSortDataKernel );
+ m_lowerSortDataKernel = b3OpenCLUtils::compileCLKernelFromString(ctx, device, kernelSource, "SearchSortDataLowerKernel", &pErrNum, boundSearchProg, additionalMacros);
+ b3Assert(m_lowerSortDataKernel);
- m_upperSortDataKernel= b3OpenCLUtils::compileCLKernelFromString( ctx, device, kernelSource, "SearchSortDataUpperKernel", &pErrNum, boundSearchProg,additionalMacros );
+ m_upperSortDataKernel = b3OpenCLUtils::compileCLKernelFromString(ctx, device, kernelSource, "SearchSortDataUpperKernel", &pErrNum, boundSearchProg, additionalMacros);
b3Assert(m_upperSortDataKernel);
m_subtractKernel = 0;
- if( maxSize )
+ if (maxSize)
{
- m_subtractKernel= b3OpenCLUtils::compileCLKernelFromString( ctx, device, kernelSource, "SubtractKernel", &pErrNum, boundSearchProg,additionalMacros );
+ m_subtractKernel = b3OpenCLUtils::compileCLKernelFromString(ctx, device, kernelSource, "SubtractKernel", &pErrNum, boundSearchProg, additionalMacros);
b3Assert(m_subtractKernel);
}
//m_constBuffer = new b3OpenCLArray<b3Int4>( device, 1, BufferBase::BUFFER_CONST );
-
- m_lower = (maxSize == 0)? 0: new b3OpenCLArray<unsigned int>(ctx,queue,maxSize );
- m_upper = (maxSize == 0)? 0: new b3OpenCLArray<unsigned int>(ctx,queue, maxSize );
- m_filler = new b3FillCL(ctx,device,queue);
+ m_lower = (maxSize == 0) ? 0 : new b3OpenCLArray<unsigned int>(ctx, queue, maxSize);
+ m_upper = (maxSize == 0) ? 0 : new b3OpenCLArray<unsigned int>(ctx, queue, maxSize);
+
+ m_filler = new b3FillCL(ctx, device, queue);
}
b3BoundSearchCL::~b3BoundSearchCL()
{
-
delete m_lower;
delete m_upper;
delete m_filler;
-
+
clReleaseKernel(m_lowerSortDataKernel);
clReleaseKernel(m_upperSortDataKernel);
clReleaseKernel(m_subtractKernel);
-
-
}
-
-void b3BoundSearchCL::execute(b3OpenCLArray<b3SortData>& src, int nSrc, b3OpenCLArray<unsigned int>& dst, int nDst, Option option )
+void b3BoundSearchCL::execute(b3OpenCLArray<b3SortData>& src, int nSrc, b3OpenCLArray<unsigned int>& dst, int nDst, Option option)
{
b3Int4 constBuffer;
constBuffer.x = nSrc;
constBuffer.y = nDst;
- if( option == BOUND_LOWER )
+ if (option == BOUND_LOWER)
{
- b3BufferInfoCL bInfo[] = { b3BufferInfoCL( src.getBufferCL(), true ), b3BufferInfoCL( dst.getBufferCL()) };
-
- b3LauncherCL launcher( m_queue, m_lowerSortDataKernel,"m_lowerSortDataKernel" );
- launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
- launcher.setConst( nSrc );
- launcher.setConst( nDst );
-
- launcher.launch1D( nSrc, 64 );
+ b3BufferInfoCL bInfo[] = {b3BufferInfoCL(src.getBufferCL(), true), b3BufferInfoCL(dst.getBufferCL())};
+
+ b3LauncherCL launcher(m_queue, m_lowerSortDataKernel, "m_lowerSortDataKernel");
+ launcher.setBuffers(bInfo, sizeof(bInfo) / sizeof(b3BufferInfoCL));
+ launcher.setConst(nSrc);
+ launcher.setConst(nDst);
+
+ launcher.launch1D(nSrc, 64);
}
- else if( option == BOUND_UPPER )
+ else if (option == BOUND_UPPER)
{
- b3BufferInfoCL bInfo[] = { b3BufferInfoCL( src.getBufferCL(), true ), b3BufferInfoCL( dst.getBufferCL() ) };
+ b3BufferInfoCL bInfo[] = {b3BufferInfoCL(src.getBufferCL(), true), b3BufferInfoCL(dst.getBufferCL())};
- b3LauncherCL launcher(m_queue, m_upperSortDataKernel,"m_upperSortDataKernel" );
- launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
- launcher.setConst( nSrc );
- launcher.setConst( nDst );
+ b3LauncherCL launcher(m_queue, m_upperSortDataKernel, "m_upperSortDataKernel");
+ launcher.setBuffers(bInfo, sizeof(bInfo) / sizeof(b3BufferInfoCL));
+ launcher.setConst(nSrc);
+ launcher.setConst(nDst);
- launcher.launch1D( nSrc, 64 );
+ launcher.launch1D(nSrc, 64);
}
- else if( option == COUNT )
+ else if (option == COUNT)
{
- b3Assert( m_lower );
- b3Assert( m_upper );
- b3Assert( m_lower->capacity() <= (int)nDst );
- b3Assert( m_upper->capacity() <= (int)nDst );
+ b3Assert(m_lower);
+ b3Assert(m_upper);
+ b3Assert(m_lower->capacity() <= (int)nDst);
+ b3Assert(m_upper->capacity() <= (int)nDst);
int zero = 0;
- m_filler->execute( *m_lower, zero, nDst );
- m_filler->execute( *m_upper, zero, nDst );
+ m_filler->execute(*m_lower, zero, nDst);
+ m_filler->execute(*m_upper, zero, nDst);
- execute( src, nSrc, *m_lower, nDst, BOUND_LOWER );
- execute( src, nSrc, *m_upper, nDst, BOUND_UPPER );
+ execute(src, nSrc, *m_lower, nDst, BOUND_LOWER);
+ execute(src, nSrc, *m_upper, nDst, BOUND_UPPER);
{
- b3BufferInfoCL bInfo[] = { b3BufferInfoCL( m_upper->getBufferCL(), true ), b3BufferInfoCL( m_lower->getBufferCL(), true ), b3BufferInfoCL( dst.getBufferCL() ) };
+ b3BufferInfoCL bInfo[] = {b3BufferInfoCL(m_upper->getBufferCL(), true), b3BufferInfoCL(m_lower->getBufferCL(), true), b3BufferInfoCL(dst.getBufferCL())};
- b3LauncherCL launcher( m_queue, m_subtractKernel ,"m_subtractKernel");
- launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
- launcher.setConst( nSrc );
- launcher.setConst( nDst );
+ b3LauncherCL launcher(m_queue, m_subtractKernel, "m_subtractKernel");
+ launcher.setBuffers(bInfo, sizeof(bInfo) / sizeof(b3BufferInfoCL));
+ launcher.setConst(nSrc);
+ launcher.setConst(nDst);
- launcher.launch1D( nDst, 64 );
+ launcher.launch1D(nDst, 64);
}
}
else
{
- b3Assert( 0 );
+ b3Assert(0);
}
-
}
-
-void b3BoundSearchCL::executeHost( b3AlignedObjectArray<b3SortData>& src, int nSrc,
- b3AlignedObjectArray<unsigned int>& dst, int nDst, Option option )
+void b3BoundSearchCL::executeHost(b3AlignedObjectArray<b3SortData>& src, int nSrc,
+ b3AlignedObjectArray<unsigned int>& dst, int nDst, Option option)
{
+ for (int i = 0; i < nSrc - 1; i++)
+ b3Assert(src[i].m_key <= src[i + 1].m_key);
-
- for(int i=0; i<nSrc-1; i++)
- b3Assert( src[i].m_key <= src[i+1].m_key );
-
- b3SortData minData,zeroData,maxData;
+ b3SortData minData, zeroData, maxData;
minData.m_key = -1;
minData.m_value = -1;
- zeroData.m_key=0;
- zeroData.m_value=0;
+ zeroData.m_key = 0;
+ zeroData.m_value = 0;
maxData.m_key = nDst;
maxData.m_value = nDst;
- if( option == BOUND_LOWER )
+ if (option == BOUND_LOWER)
{
- for(int i=0; i<nSrc; i++)
+ for (int i = 0; i < nSrc; i++)
{
- b3SortData& iData = (i==0)? minData: src[i-1];
- b3SortData& jData = (i==nSrc)? maxData: src[i];
+ b3SortData& iData = (i == 0) ? minData : src[i - 1];
+ b3SortData& jData = (i == nSrc) ? maxData : src[i];
- if( iData.m_key != jData.m_key )
+ if (iData.m_key != jData.m_key)
{
int k = jData.m_key;
{
@@ -170,14 +160,14 @@ void b3BoundSearchCL::executeHost( b3AlignedObjectArray<b3SortData>& src, int nS
}
}
}
- else if( option == BOUND_UPPER )
+ else if (option == BOUND_UPPER)
{
- for(int i=1; i<nSrc+1; i++)
+ for (int i = 1; i < nSrc + 1; i++)
{
- b3SortData& iData = src[i-1];
- b3SortData& jData = (i==nSrc)? maxData: src[i];
+ b3SortData& iData = src[i - 1];
+ b3SortData& jData = (i == nSrc) ? maxData : src[i];
- if( iData.m_key != jData.m_key )
+ if (iData.m_key != jData.m_key)
{
int k = iData.m_key;
{
@@ -186,28 +176,28 @@ void b3BoundSearchCL::executeHost( b3AlignedObjectArray<b3SortData>& src, int nS
}
}
}
- else if( option == COUNT )
+ else if (option == COUNT)
{
b3AlignedObjectArray<unsigned int> lower;
- lower.resize(nDst );
+ lower.resize(nDst);
b3AlignedObjectArray<unsigned int> upper;
- upper.resize(nDst );
+ upper.resize(nDst);
- for(int i=0; i<nDst; i++)
- {
- lower[i] = upper[i] = 0;
+ for (int i = 0; i < nDst; i++)
+ {
+ lower[i] = upper[i] = 0;
}
- executeHost( src, nSrc, lower, nDst, BOUND_LOWER );
- executeHost( src, nSrc, upper, nDst, BOUND_UPPER );
+ executeHost(src, nSrc, lower, nDst, BOUND_LOWER);
+ executeHost(src, nSrc, upper, nDst, BOUND_UPPER);
- for( int i=0; i<nDst; i++)
- {
- dst[i] = upper[i] - lower[i];
+ for (int i = 0; i < nDst; i++)
+ {
+ dst[i] = upper[i] - lower[i];
}
}
else
{
- b3Assert( 0 );
+ b3Assert(0);
}
}
diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3BoundSearchCL.h b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3BoundSearchCL.h
index 7e2940965c..0d633e3d23 100644
--- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3BoundSearchCL.h
+++ b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3BoundSearchCL.h
@@ -26,42 +26,39 @@ subject to the following restrictions:
#include "b3OpenCLArray.h"
#include "b3FillCL.h"
-#include "b3RadixSort32CL.h" //for b3SortData (perhaps move it?)
+#include "b3RadixSort32CL.h" //for b3SortData (perhaps move it?)
class b3BoundSearchCL
{
- public:
+public:
+ enum Option
+ {
+ BOUND_LOWER,
+ BOUND_UPPER,
+ COUNT,
+ };
- enum Option
- {
- BOUND_LOWER,
- BOUND_UPPER,
- COUNT,
- };
+ cl_context m_context;
+ cl_device_id m_device;
+ cl_command_queue m_queue;
- cl_context m_context;
- cl_device_id m_device;
- cl_command_queue m_queue;
+ cl_kernel m_lowerSortDataKernel;
+ cl_kernel m_upperSortDataKernel;
+ cl_kernel m_subtractKernel;
-
- cl_kernel m_lowerSortDataKernel;
- cl_kernel m_upperSortDataKernel;
- cl_kernel m_subtractKernel;
-
- b3OpenCLArray<b3Int4>* m_constbtOpenCLArray;
- b3OpenCLArray<unsigned int>* m_lower;
- b3OpenCLArray<unsigned int>* m_upper;
-
- b3FillCL* m_filler;
-
- b3BoundSearchCL(cl_context context, cl_device_id device, cl_command_queue queue, int size);
+ b3OpenCLArray<b3Int4>* m_constbtOpenCLArray;
+ b3OpenCLArray<unsigned int>* m_lower;
+ b3OpenCLArray<unsigned int>* m_upper;
- virtual ~b3BoundSearchCL();
+ b3FillCL* m_filler;
- // src has to be src[i].m_key <= src[i+1].m_key
- void execute( b3OpenCLArray<b3SortData>& src, int nSrc, b3OpenCLArray<unsigned int>& dst, int nDst, Option option = BOUND_LOWER );
+ b3BoundSearchCL(cl_context context, cl_device_id device, cl_command_queue queue, int size);
- void executeHost( b3AlignedObjectArray<b3SortData>& src, int nSrc, b3AlignedObjectArray<unsigned int>& dst, int nDst, Option option = BOUND_LOWER);
-};
+ virtual ~b3BoundSearchCL();
+
+ // src has to be src[i].m_key <= src[i+1].m_key
+ void execute(b3OpenCLArray<b3SortData>& src, int nSrc, b3OpenCLArray<unsigned int>& dst, int nDst, Option option = BOUND_LOWER);
+ void executeHost(b3AlignedObjectArray<b3SortData>& src, int nSrc, b3AlignedObjectArray<unsigned int>& dst, int nDst, Option option = BOUND_LOWER);
+};
-#endif //B3_BOUNDSEARCH_H
+#endif //B3_BOUNDSEARCH_H
diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3BufferInfoCL.h b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3BufferInfoCL.h
index 52f219ae3f..35fc467b20 100644
--- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3BufferInfoCL.h
+++ b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3BufferInfoCL.h
@@ -4,16 +4,15 @@
#include "b3OpenCLArray.h"
-
struct b3BufferInfoCL
{
//b3BufferInfoCL(){}
-// template<typename T>
- b3BufferInfoCL(cl_mem buff, bool isReadOnly = false): m_clBuffer(buff), m_isReadOnly(isReadOnly){}
+ // template<typename T>
+ b3BufferInfoCL(cl_mem buff, bool isReadOnly = false) : m_clBuffer(buff), m_isReadOnly(isReadOnly) {}
cl_mem m_clBuffer;
bool m_isReadOnly;
};
-#endif //B3_BUFFER_INFO_CL_H
+#endif //B3_BUFFER_INFO_CL_H
diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3FillCL.cpp b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3FillCL.cpp
index f05c2648f1..bd25bb2101 100644
--- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3FillCL.cpp
+++ b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3FillCL.cpp
@@ -8,29 +8,26 @@
#include "kernels/FillKernelsCL.h"
b3FillCL::b3FillCL(cl_context ctx, cl_device_id device, cl_command_queue queue)
-:m_commandQueue(queue)
+ : m_commandQueue(queue)
{
const char* kernelSource = fillKernelsCL;
cl_int pErrNum;
const char* additionalMacros = "";
- cl_program fillProg = b3OpenCLUtils::compileCLProgramFromString( ctx, device, kernelSource, &pErrNum,additionalMacros, FILL_CL_PROGRAM_PATH);
+ cl_program fillProg = b3OpenCLUtils::compileCLProgramFromString(ctx, device, kernelSource, &pErrNum, additionalMacros, FILL_CL_PROGRAM_PATH);
b3Assert(fillProg);
- m_fillIntKernel = b3OpenCLUtils::compileCLKernelFromString( ctx, device, kernelSource, "FillIntKernel", &pErrNum, fillProg,additionalMacros );
+ m_fillIntKernel = b3OpenCLUtils::compileCLKernelFromString(ctx, device, kernelSource, "FillIntKernel", &pErrNum, fillProg, additionalMacros);
b3Assert(m_fillIntKernel);
- m_fillUnsignedIntKernel = b3OpenCLUtils::compileCLKernelFromString( ctx, device, kernelSource, "FillUnsignedIntKernel", &pErrNum, fillProg,additionalMacros );
+ m_fillUnsignedIntKernel = b3OpenCLUtils::compileCLKernelFromString(ctx, device, kernelSource, "FillUnsignedIntKernel", &pErrNum, fillProg, additionalMacros);
b3Assert(m_fillIntKernel);
- m_fillFloatKernel = b3OpenCLUtils::compileCLKernelFromString( ctx, device, kernelSource, "FillFloatKernel", &pErrNum, fillProg,additionalMacros );
+ m_fillFloatKernel = b3OpenCLUtils::compileCLKernelFromString(ctx, device, kernelSource, "FillFloatKernel", &pErrNum, fillProg, additionalMacros);
b3Assert(m_fillFloatKernel);
-
-
- m_fillKernelInt2 = b3OpenCLUtils::compileCLKernelFromString( ctx, device, kernelSource, "FillInt2Kernel", &pErrNum, fillProg,additionalMacros );
+ m_fillKernelInt2 = b3OpenCLUtils::compileCLKernelFromString(ctx, device, kernelSource, "FillInt2Kernel", &pErrNum, fillProg, additionalMacros);
b3Assert(m_fillKernelInt2);
-
}
b3FillCL::~b3FillCL()
@@ -39,88 +36,84 @@ b3FillCL::~b3FillCL()
clReleaseKernel(m_fillIntKernel);
clReleaseKernel(m_fillUnsignedIntKernel);
clReleaseKernel(m_fillFloatKernel);
-
}
void b3FillCL::execute(b3OpenCLArray<float>& src, const float value, int n, int offset)
{
- b3Assert( n>0 );
+ b3Assert(n > 0);
{
- b3LauncherCL launcher( m_commandQueue, m_fillFloatKernel,"m_fillFloatKernel" );
- launcher.setBuffer( src.getBufferCL());
- launcher.setConst( n );
- launcher.setConst( value );
- launcher.setConst( offset);
+ b3LauncherCL launcher(m_commandQueue, m_fillFloatKernel, "m_fillFloatKernel");
+ launcher.setBuffer(src.getBufferCL());
+ launcher.setConst(n);
+ launcher.setConst(value);
+ launcher.setConst(offset);
- launcher.launch1D( n );
+ launcher.launch1D(n);
}
}
void b3FillCL::execute(b3OpenCLArray<int>& src, const int value, int n, int offset)
{
- b3Assert( n>0 );
-
+ b3Assert(n > 0);
{
- b3LauncherCL launcher( m_commandQueue, m_fillIntKernel ,"m_fillIntKernel");
+ b3LauncherCL launcher(m_commandQueue, m_fillIntKernel, "m_fillIntKernel");
launcher.setBuffer(src.getBufferCL());
- launcher.setConst( n);
- launcher.setConst( value);
- launcher.setConst( offset);
- launcher.launch1D( n );
+ launcher.setConst(n);
+ launcher.setConst(value);
+ launcher.setConst(offset);
+ launcher.launch1D(n);
}
}
-
void b3FillCL::execute(b3OpenCLArray<unsigned int>& src, const unsigned int value, int n, int offset)
{
- b3Assert( n>0 );
+ b3Assert(n > 0);
{
- b3BufferInfoCL bInfo[] = { b3BufferInfoCL( src.getBufferCL() ) };
+ b3BufferInfoCL bInfo[] = {b3BufferInfoCL(src.getBufferCL())};
- b3LauncherCL launcher( m_commandQueue, m_fillUnsignedIntKernel,"m_fillUnsignedIntKernel" );
- launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
- launcher.setConst( n );
- launcher.setConst(value);
+ b3LauncherCL launcher(m_commandQueue, m_fillUnsignedIntKernel, "m_fillUnsignedIntKernel");
+ launcher.setBuffers(bInfo, sizeof(bInfo) / sizeof(b3BufferInfoCL));
+ launcher.setConst(n);
+ launcher.setConst(value);
launcher.setConst(offset);
- launcher.launch1D( n );
+ launcher.launch1D(n);
}
}
-void b3FillCL::executeHost(b3AlignedObjectArray<b3Int2> &src, const b3Int2 &value, int n, int offset)
+void b3FillCL::executeHost(b3AlignedObjectArray<b3Int2>& src, const b3Int2& value, int n, int offset)
{
- for (int i=0;i<n;i++)
+ for (int i = 0; i < n; i++)
{
- src[i+offset]=value;
+ src[i + offset] = value;
}
}
-void b3FillCL::executeHost(b3AlignedObjectArray<int> &src, const int value, int n, int offset)
+void b3FillCL::executeHost(b3AlignedObjectArray<int>& src, const int value, int n, int offset)
{
- for (int i=0;i<n;i++)
+ for (int i = 0; i < n; i++)
{
- src[i+offset]=value;
+ src[i + offset] = value;
}
}
-void b3FillCL::execute(b3OpenCLArray<b3Int2> &src, const b3Int2 &value, int n, int offset)
+void b3FillCL::execute(b3OpenCLArray<b3Int2>& src, const b3Int2& value, int n, int offset)
{
- b3Assert( n>0 );
-
+ b3Assert(n > 0);
{
- b3BufferInfoCL bInfo[] = { b3BufferInfoCL( src.getBufferCL() ) };
+ b3BufferInfoCL bInfo[] = {b3BufferInfoCL(src.getBufferCL())};
- b3LauncherCL launcher(m_commandQueue, m_fillKernelInt2,"m_fillKernelInt2");
- launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
+ b3LauncherCL launcher(m_commandQueue, m_fillKernelInt2, "m_fillKernelInt2");
+ launcher.setBuffers(bInfo, sizeof(bInfo) / sizeof(b3BufferInfoCL));
launcher.setConst(n);
launcher.setConst(value);
launcher.setConst(offset);
//( constBuffer );
- launcher.launch1D( n );
+ launcher.launch1D(n);
}
}
diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3FillCL.h b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3FillCL.h
index 1609676b9d..c92c3e5119 100644
--- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3FillCL.h
+++ b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3FillCL.h
@@ -7,57 +7,46 @@
#include "Bullet3Common/shared/b3Int2.h"
#include "Bullet3Common/shared/b3Int4.h"
-
class b3FillCL
{
-
- cl_command_queue m_commandQueue;
-
- cl_kernel m_fillKernelInt2;
- cl_kernel m_fillIntKernel;
- cl_kernel m_fillUnsignedIntKernel;
- cl_kernel m_fillFloatKernel;
-
- public:
-
- struct b3ConstData
- {
- union
- {
- b3Int4 m_data;
- b3UnsignedInt4 m_UnsignedData;
- };
- int m_offset;
- int m_n;
- int m_padding[2];
+ cl_command_queue m_commandQueue;
+
+ cl_kernel m_fillKernelInt2;
+ cl_kernel m_fillIntKernel;
+ cl_kernel m_fillUnsignedIntKernel;
+ cl_kernel m_fillFloatKernel;
+
+public:
+ struct b3ConstData
+ {
+ union {
+ b3Int4 m_data;
+ b3UnsignedInt4 m_UnsignedData;
};
+ int m_offset;
+ int m_n;
+ int m_padding[2];
+ };
protected:
-
public:
+ b3FillCL(cl_context ctx, cl_device_id device, cl_command_queue queue);
- b3FillCL(cl_context ctx, cl_device_id device, cl_command_queue queue);
+ virtual ~b3FillCL();
- virtual ~b3FillCL();
+ void execute(b3OpenCLArray<unsigned int>& src, const unsigned int value, int n, int offset = 0);
- void execute(b3OpenCLArray<unsigned int>& src, const unsigned int value, int n, int offset = 0);
-
- void execute(b3OpenCLArray<int>& src, const int value, int n, int offset = 0);
+ void execute(b3OpenCLArray<int>& src, const int value, int n, int offset = 0);
- void execute(b3OpenCLArray<float>& src, const float value, int n, int offset = 0);
+ void execute(b3OpenCLArray<float>& src, const float value, int n, int offset = 0);
- void execute(b3OpenCLArray<b3Int2>& src, const b3Int2& value, int n, int offset = 0);
+ void execute(b3OpenCLArray<b3Int2>& src, const b3Int2& value, int n, int offset = 0);
- void executeHost(b3AlignedObjectArray<b3Int2> &src, const b3Int2 &value, int n, int offset);
+ void executeHost(b3AlignedObjectArray<b3Int2>& src, const b3Int2& value, int n, int offset);
- void executeHost(b3AlignedObjectArray<int> &src, const int value, int n, int offset);
+ void executeHost(b3AlignedObjectArray<int>& src, const int value, int n, int offset);
// void execute(b3OpenCLArray<b3Int4>& src, const b3Int4& value, int n, int offset = 0);
-
};
-
-
-
-
-#endif //B3_FILL_CL_H
+#endif //B3_FILL_CL_H
diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3LauncherCL.cpp b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3LauncherCL.cpp
index 94590d11ca..c97d02eb45 100644
--- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3LauncherCL.cpp
+++ b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3LauncherCL.cpp
@@ -1,13 +1,13 @@
#include "b3LauncherCL.h"
bool gDebugLauncherCL = false;
-
+
b3LauncherCL::b3LauncherCL(cl_command_queue queue, cl_kernel kernel, const char* name)
-:m_commandQueue(queue),
-m_kernel(kernel),
-m_idx(0),
-m_enableSerialization(false),
-m_name(name)
+ : m_commandQueue(queue),
+ m_kernel(kernel),
+ m_idx(0),
+ m_enableSerialization(false),
+ m_name(name)
{
if (gDebugLauncherCL)
{
@@ -15,59 +15,58 @@ m_name(name)
printf("[%d] Prepare to launch OpenCL kernel %s\n", counter++, name);
}
- m_serializationSizeInBytes = sizeof(int);
+ m_serializationSizeInBytes = sizeof(int);
}
-
+
b3LauncherCL::~b3LauncherCL()
- {
- for (int i=0;i<m_arrays.size();i++)
- {
- delete (m_arrays[i]);
- }
-
- m_arrays.clear();
- if (gDebugLauncherCL)
- {
+{
+ for (int i = 0; i < m_arrays.size(); i++)
+ {
+ delete (m_arrays[i]);
+ }
+
+ m_arrays.clear();
+ if (gDebugLauncherCL)
+ {
static int counter = 0;
- printf("[%d] Finished launching OpenCL kernel %s\n", counter++,m_name);
- }
- }
+ printf("[%d] Finished launching OpenCL kernel %s\n", counter++, m_name);
+ }
+}
-void b3LauncherCL::setBuffer( cl_mem clBuffer)
+void b3LauncherCL::setBuffer(cl_mem clBuffer)
{
- if (m_enableSerialization)
- {
- b3KernelArgData kernelArg;
- kernelArg.m_argIndex = m_idx;
- kernelArg.m_isBuffer = 1;
- kernelArg.m_clBuffer = clBuffer;
-
- cl_mem_info param_name = CL_MEM_SIZE;
- size_t param_value;
- size_t sizeInBytes = sizeof(size_t);
- size_t actualSizeInBytes;
- cl_int err;
- err = clGetMemObjectInfo ( kernelArg.m_clBuffer,
- param_name,
- sizeInBytes,
- &param_value,
- &actualSizeInBytes);
-
- b3Assert( err == CL_SUCCESS );
- kernelArg.m_argSizeInBytes = param_value;
-
- m_kernelArguments.push_back(kernelArg);
- m_serializationSizeInBytes+= sizeof(b3KernelArgData);
- m_serializationSizeInBytes+=param_value;
- }
- cl_int status = clSetKernelArg( m_kernel, m_idx++, sizeof(cl_mem), &clBuffer);
- b3Assert( status == CL_SUCCESS );
-}
+ if (m_enableSerialization)
+ {
+ b3KernelArgData kernelArg;
+ kernelArg.m_argIndex = m_idx;
+ kernelArg.m_isBuffer = 1;
+ kernelArg.m_clBuffer = clBuffer;
+ cl_mem_info param_name = CL_MEM_SIZE;
+ size_t param_value;
+ size_t sizeInBytes = sizeof(size_t);
+ size_t actualSizeInBytes;
+ cl_int err;
+ err = clGetMemObjectInfo(kernelArg.m_clBuffer,
+ param_name,
+ sizeInBytes,
+ &param_value,
+ &actualSizeInBytes);
+
+ b3Assert(err == CL_SUCCESS);
+ kernelArg.m_argSizeInBytes = param_value;
+
+ m_kernelArguments.push_back(kernelArg);
+ m_serializationSizeInBytes += sizeof(b3KernelArgData);
+ m_serializationSizeInBytes += param_value;
+ }
+ cl_int status = clSetKernelArg(m_kernel, m_idx++, sizeof(cl_mem), &clBuffer);
+ b3Assert(status == CL_SUCCESS);
+}
-void b3LauncherCL::setBuffers( b3BufferInfoCL* buffInfo, int n )
+void b3LauncherCL::setBuffers(b3BufferInfoCL* buffInfo, int n)
{
- for(int i=0; i<n; i++)
+ for (int i = 0; i < n; i++)
{
if (m_enableSerialization)
{
@@ -75,106 +74,103 @@ void b3LauncherCL::setBuffers( b3BufferInfoCL* buffInfo, int n )
kernelArg.m_argIndex = m_idx;
kernelArg.m_isBuffer = 1;
kernelArg.m_clBuffer = buffInfo[i].m_clBuffer;
-
+
cl_mem_info param_name = CL_MEM_SIZE;
size_t param_value;
size_t sizeInBytes = sizeof(size_t);
size_t actualSizeInBytes;
cl_int err;
- err = clGetMemObjectInfo ( kernelArg.m_clBuffer,
- param_name,
- sizeInBytes,
- &param_value,
- &actualSizeInBytes);
-
- b3Assert( err == CL_SUCCESS );
+ err = clGetMemObjectInfo(kernelArg.m_clBuffer,
+ param_name,
+ sizeInBytes,
+ &param_value,
+ &actualSizeInBytes);
+
+ b3Assert(err == CL_SUCCESS);
kernelArg.m_argSizeInBytes = param_value;
-
+
m_kernelArguments.push_back(kernelArg);
- m_serializationSizeInBytes+= sizeof(b3KernelArgData);
- m_serializationSizeInBytes+=param_value;
- }
- cl_int status = clSetKernelArg( m_kernel, m_idx++, sizeof(cl_mem), &buffInfo[i].m_clBuffer);
- b3Assert( status == CL_SUCCESS );
- }
+ m_serializationSizeInBytes += sizeof(b3KernelArgData);
+ m_serializationSizeInBytes += param_value;
+ }
+ cl_int status = clSetKernelArg(m_kernel, m_idx++, sizeof(cl_mem), &buffInfo[i].m_clBuffer);
+ b3Assert(status == CL_SUCCESS);
+ }
}
struct b3KernelArgDataUnaligned
{
- int m_isBuffer;
- int m_argIndex;
- int m_argSizeInBytes;
+ int m_isBuffer;
+ int m_argIndex;
+ int m_argSizeInBytes;
int m_unusedPadding;
- union
- {
- cl_mem m_clBuffer;
- unsigned char m_argData[B3_CL_MAX_ARG_SIZE];
- };
-
+ union {
+ cl_mem m_clBuffer;
+ unsigned char m_argData[B3_CL_MAX_ARG_SIZE];
+ };
};
#include <string.h>
-
-
int b3LauncherCL::deserializeArgs(unsigned char* buf, int bufSize, cl_context ctx)
{
- int index=0;
-
- int numArguments = *(int*) &buf[index];
- index+=sizeof(int);
-
- for (int i=0;i<numArguments;i++)
- {
- b3KernelArgDataUnaligned* arg = (b3KernelArgDataUnaligned*)&buf[index];
-
- index+=sizeof(b3KernelArgData);
- if (arg->m_isBuffer)
- {
- b3OpenCLArray<unsigned char>* clData = new b3OpenCLArray<unsigned char>(ctx,m_commandQueue, arg->m_argSizeInBytes);
- clData->resize(arg->m_argSizeInBytes);
-
- clData->copyFromHostPointer(&buf[index], arg->m_argSizeInBytes);
-
- arg->m_clBuffer = clData->getBufferCL();
-
- m_arrays.push_back(clData);
-
- cl_int status = clSetKernelArg( m_kernel, m_idx++, sizeof(cl_mem), &arg->m_clBuffer);
- b3Assert( status == CL_SUCCESS );
- index+=arg->m_argSizeInBytes;
- } else
- {
- cl_int status = clSetKernelArg( m_kernel, m_idx++, arg->m_argSizeInBytes, &arg->m_argData);
- b3Assert( status == CL_SUCCESS );
- }
+ int index = 0;
+
+ int numArguments = *(int*)&buf[index];
+ index += sizeof(int);
+
+ for (int i = 0; i < numArguments; i++)
+ {
+ b3KernelArgDataUnaligned* arg = (b3KernelArgDataUnaligned*)&buf[index];
+
+ index += sizeof(b3KernelArgData);
+ if (arg->m_isBuffer)
+ {
+ b3OpenCLArray<unsigned char>* clData = new b3OpenCLArray<unsigned char>(ctx, m_commandQueue, arg->m_argSizeInBytes);
+ clData->resize(arg->m_argSizeInBytes);
+
+ clData->copyFromHostPointer(&buf[index], arg->m_argSizeInBytes);
+
+ arg->m_clBuffer = clData->getBufferCL();
+
+ m_arrays.push_back(clData);
+
+ cl_int status = clSetKernelArg(m_kernel, m_idx++, sizeof(cl_mem), &arg->m_clBuffer);
+ b3Assert(status == CL_SUCCESS);
+ index += arg->m_argSizeInBytes;
+ }
+ else
+ {
+ cl_int status = clSetKernelArg(m_kernel, m_idx++, arg->m_argSizeInBytes, &arg->m_argData);
+ b3Assert(status == CL_SUCCESS);
+ }
b3KernelArgData b;
- memcpy(&b,arg,sizeof(b3KernelArgDataUnaligned));
- m_kernelArguments.push_back(b);
- }
-m_serializationSizeInBytes = index;
- return index;
+ memcpy(&b, arg, sizeof(b3KernelArgDataUnaligned));
+ m_kernelArguments.push_back(b);
+ }
+ m_serializationSizeInBytes = index;
+ return index;
}
int b3LauncherCL::validateResults(unsigned char* goldBuffer, int goldBufferCapacity, cl_context ctx)
- {
- int index=0;
-
- int numArguments = *(int*) &goldBuffer[index];
- index+=sizeof(int);
+{
+ int index = 0;
+
+ int numArguments = *(int*)&goldBuffer[index];
+ index += sizeof(int);
if (numArguments != m_kernelArguments.size())
{
- printf("failed validation: expected %d arguments, found %d\n",numArguments, m_kernelArguments.size());
+ printf("failed validation: expected %d arguments, found %d\n", numArguments, m_kernelArguments.size());
return -1;
}
-
- for (int ii=0;ii<numArguments;ii++)
- {
- b3KernelArgData* argGold = (b3KernelArgData*)&goldBuffer[index];
+
+ for (int ii = 0; ii < numArguments; ii++)
+ {
+ b3KernelArgData* argGold = (b3KernelArgData*)&goldBuffer[index];
if (m_kernelArguments[ii].m_argSizeInBytes != argGold->m_argSizeInBytes)
{
- printf("failed validation: argument %d sizeInBytes expected: %d, found %d\n",ii, argGold->m_argSizeInBytes, m_kernelArguments[ii].m_argSizeInBytes);
+ printf("failed validation: argument %d sizeInBytes expected: %d, found %d\n", ii, argGold->m_argSizeInBytes, m_kernelArguments[ii].m_argSizeInBytes);
return -2;
}
@@ -184,125 +180,117 @@ int b3LauncherCL::validateResults(unsigned char* goldBuffer, int goldBufferCapac
if (expected != found)
{
- printf("failed validation: argument %d isBuffer expected: %d, found %d\n",ii,expected, found);
+ printf("failed validation: argument %d isBuffer expected: %d, found %d\n", ii, expected, found);
return -3;
}
}
- index+=sizeof(b3KernelArgData);
+ index += sizeof(b3KernelArgData);
if (argGold->m_isBuffer)
- {
-
- unsigned char* memBuf= (unsigned char*) malloc(m_kernelArguments[ii].m_argSizeInBytes);
+ {
+ unsigned char* memBuf = (unsigned char*)malloc(m_kernelArguments[ii].m_argSizeInBytes);
unsigned char* goldBuf = &goldBuffer[index];
- for (int j=0;j<m_kernelArguments[j].m_argSizeInBytes;j++)
+ for (int j = 0; j < m_kernelArguments[j].m_argSizeInBytes; j++)
{
memBuf[j] = 0xaa;
}
cl_int status = 0;
- status = clEnqueueReadBuffer( m_commandQueue, m_kernelArguments[ii].m_clBuffer, CL_TRUE, 0, m_kernelArguments[ii].m_argSizeInBytes,
- memBuf, 0,0,0 );
- b3Assert( status==CL_SUCCESS );
- clFinish(m_commandQueue);
+ status = clEnqueueReadBuffer(m_commandQueue, m_kernelArguments[ii].m_clBuffer, CL_TRUE, 0, m_kernelArguments[ii].m_argSizeInBytes,
+ memBuf, 0, 0, 0);
+ b3Assert(status == CL_SUCCESS);
+ clFinish(m_commandQueue);
- for (int b=0;b<m_kernelArguments[ii].m_argSizeInBytes;b++)
+ for (int b = 0; b < m_kernelArguments[ii].m_argSizeInBytes; b++)
{
int expected = goldBuf[b];
int found = memBuf[b];
if (expected != found)
{
printf("failed validation: argument %d OpenCL data at byte position %d expected: %d, found %d\n",
- ii, b, expected, found);
+ ii, b, expected, found);
return -4;
}
}
-
- index+=argGold->m_argSizeInBytes;
- } else
- {
-
+ index += argGold->m_argSizeInBytes;
+ }
+ else
+ {
//compare content
- for (int b=0;b<m_kernelArguments[ii].m_argSizeInBytes;b++)
+ for (int b = 0; b < m_kernelArguments[ii].m_argSizeInBytes; b++)
{
int expected = argGold->m_argData[b];
- int found =m_kernelArguments[ii].m_argData[b];
+ int found = m_kernelArguments[ii].m_argData[b];
if (expected != found)
{
printf("failed validation: argument %d const data at byte position %d expected: %d, found %d\n",
- ii, b, expected, found);
+ ii, b, expected, found);
return -5;
}
}
-
- }
- }
- return index;
-
+ }
+ }
+ return index;
}
int b3LauncherCL::serializeArguments(unsigned char* destBuffer, int destBufferCapacity)
{
-//initialize to known values
-for (int i=0;i<destBufferCapacity;i++)
- destBuffer[i] = 0xec;
-
- assert(destBufferCapacity>=m_serializationSizeInBytes);
-
- //todo: use the b3Serializer for this to allow for 32/64bit, endianness etc
- int numArguments = m_kernelArguments.size();
- int curBufferSize = 0;
- int* dest = (int*)&destBuffer[curBufferSize];
- *dest = numArguments;
- curBufferSize += sizeof(int);
-
-
-
- for (int i=0;i<this->m_kernelArguments.size();i++)
- {
- b3KernelArgData* arg = (b3KernelArgData*) &destBuffer[curBufferSize];
- *arg = m_kernelArguments[i];
- curBufferSize+=sizeof(b3KernelArgData);
- if (arg->m_isBuffer==1)
- {
- //copy the OpenCL buffer content
- cl_int status = 0;
- status = clEnqueueReadBuffer( m_commandQueue, arg->m_clBuffer, 0, 0, arg->m_argSizeInBytes,
- &destBuffer[curBufferSize], 0,0,0 );
- b3Assert( status==CL_SUCCESS );
- clFinish(m_commandQueue);
- curBufferSize+=arg->m_argSizeInBytes;
- }
-
- }
- return curBufferSize;
+ //initialize to known values
+ for (int i = 0; i < destBufferCapacity; i++)
+ destBuffer[i] = 0xec;
+
+ assert(destBufferCapacity >= m_serializationSizeInBytes);
+
+ //todo: use the b3Serializer for this to allow for 32/64bit, endianness etc
+ int numArguments = m_kernelArguments.size();
+ int curBufferSize = 0;
+ int* dest = (int*)&destBuffer[curBufferSize];
+ *dest = numArguments;
+ curBufferSize += sizeof(int);
+
+ for (int i = 0; i < this->m_kernelArguments.size(); i++)
+ {
+ b3KernelArgData* arg = (b3KernelArgData*)&destBuffer[curBufferSize];
+ *arg = m_kernelArguments[i];
+ curBufferSize += sizeof(b3KernelArgData);
+ if (arg->m_isBuffer == 1)
+ {
+ //copy the OpenCL buffer content
+ cl_int status = 0;
+ status = clEnqueueReadBuffer(m_commandQueue, arg->m_clBuffer, 0, 0, arg->m_argSizeInBytes,
+ &destBuffer[curBufferSize], 0, 0, 0);
+ b3Assert(status == CL_SUCCESS);
+ clFinish(m_commandQueue);
+ curBufferSize += arg->m_argSizeInBytes;
+ }
+ }
+ return curBufferSize;
}
void b3LauncherCL::serializeToFile(const char* fileName, int numWorkItems)
{
int num = numWorkItems;
int buffSize = getSerializationBufferSize();
- unsigned char* buf = new unsigned char[buffSize+sizeof(int)];
- for (int i=0;i<buffSize+1;i++)
+ unsigned char* buf = new unsigned char[buffSize + sizeof(int)];
+ for (int i = 0; i < buffSize + 1; i++)
{
unsigned char* ptr = (unsigned char*)&buf[i];
*ptr = 0xff;
}
-// int actualWrite = serializeArguments(buf,buffSize);
-
-// unsigned char* cptr = (unsigned char*)&buf[buffSize];
-// printf("buf[buffSize] = %d\n",*cptr);
-
- assert(buf[buffSize]==0xff);//check for buffer overrun
+ // int actualWrite = serializeArguments(buf,buffSize);
+
+ // unsigned char* cptr = (unsigned char*)&buf[buffSize];
+ // printf("buf[buffSize] = %d\n",*cptr);
+
+ assert(buf[buffSize] == 0xff); //check for buffer overrun
int* ptr = (int*)&buf[buffSize];
-
+
*ptr = num;
-
- FILE* f = fopen(fileName,"wb");
- fwrite(buf,buffSize+sizeof(int),1,f);
+
+ FILE* f = fopen(fileName, "wb");
+ fwrite(buf, buffSize + sizeof(int), 1, f);
fclose(f);
delete[] buf;
-}
-
+}
diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3LauncherCL.h b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3LauncherCL.h
index 1b267b31ef..18e9c1db2b 100644
--- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3LauncherCL.h
+++ b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3LauncherCL.h
@@ -9,60 +9,57 @@
#define B3_DEBUG_SERIALIZE_CL
-
#ifdef _WIN32
-#pragma warning(disable :4996)
+#pragma warning(disable : 4996)
#endif
#define B3_CL_MAX_ARG_SIZE 16
-B3_ATTRIBUTE_ALIGNED16(struct) b3KernelArgData
+B3_ATTRIBUTE_ALIGNED16(struct)
+b3KernelArgData
{
- int m_isBuffer;
- int m_argIndex;
- int m_argSizeInBytes;
+ int m_isBuffer;
+ int m_argIndex;
+ int m_argSizeInBytes;
int m_unusedPadding;
- union
- {
- cl_mem m_clBuffer;
- unsigned char m_argData[B3_CL_MAX_ARG_SIZE];
- };
-
+ union {
+ cl_mem m_clBuffer;
+ unsigned char m_argData[B3_CL_MAX_ARG_SIZE];
+ };
};
class b3LauncherCL
{
-
cl_command_queue m_commandQueue;
cl_kernel m_kernel;
int m_idx;
-
- b3AlignedObjectArray<b3KernelArgData> m_kernelArguments;
- int m_serializationSizeInBytes;
- bool m_enableSerialization;
+
+ b3AlignedObjectArray<b3KernelArgData> m_kernelArguments;
+ int m_serializationSizeInBytes;
+ bool m_enableSerialization;
const char* m_name;
- public:
-
- b3AlignedObjectArray<b3OpenCLArray<unsigned char>* > m_arrays;
-
- b3LauncherCL(cl_command_queue queue, cl_kernel kernel, const char* name);
-
- virtual ~b3LauncherCL();
-
- void setBuffer( cl_mem clBuffer);
-
- void setBuffers( b3BufferInfoCL* buffInfo, int n );
-
- int getSerializationBufferSize() const
- {
- return m_serializationSizeInBytes;
- }
-
- int deserializeArgs(unsigned char* buf, int bufSize, cl_context ctx);
+
+public:
+ b3AlignedObjectArray<b3OpenCLArray<unsigned char>*> m_arrays;
+
+ b3LauncherCL(cl_command_queue queue, cl_kernel kernel, const char* name);
+
+ virtual ~b3LauncherCL();
+
+ void setBuffer(cl_mem clBuffer);
+
+ void setBuffers(b3BufferInfoCL* buffInfo, int n);
+
+ int getSerializationBufferSize() const
+ {
+ return m_serializationSizeInBytes;
+ }
+
+ int deserializeArgs(unsigned char* buf, int bufSize, cl_context ctx);
inline int validateResults(unsigned char* goldBuffer, int goldBufferCapacity, cl_context ctx);
- int serializeArguments(unsigned char* destBuffer, int destBufferCapacity);
-
+ int serializeArguments(unsigned char* destBuffer, int destBufferCapacity);
+
int getNumArguments() const
{
return m_kernelArguments.size();
@@ -75,61 +72,57 @@ class b3LauncherCL
void serializeToFile(const char* fileName, int numWorkItems);
- template<typename T>
- inline void setConst( const T& consts )
- {
- int sz=sizeof(T);
- b3Assert(sz<=B3_CL_MAX_ARG_SIZE);
-
- if (m_enableSerialization)
- {
- b3KernelArgData kernelArg;
- kernelArg.m_argIndex = m_idx;
- kernelArg.m_isBuffer = 0;
- T* destArg = (T*)kernelArg.m_argData;
- *destArg = consts;
- kernelArg.m_argSizeInBytes = sizeof(T);
- m_kernelArguments.push_back(kernelArg);
- m_serializationSizeInBytes+=sizeof(b3KernelArgData);
- }
-
- cl_int status = clSetKernelArg( m_kernel, m_idx++, sz, &consts );
- b3Assert( status == CL_SUCCESS );
- }
+ template <typename T>
+ inline void setConst(const T& consts)
+ {
+ int sz = sizeof(T);
+ b3Assert(sz <= B3_CL_MAX_ARG_SIZE);
- inline void launch1D( int numThreads, int localSize = 64)
+ if (m_enableSerialization)
{
- launch2D( numThreads, 1, localSize, 1 );
+ b3KernelArgData kernelArg;
+ kernelArg.m_argIndex = m_idx;
+ kernelArg.m_isBuffer = 0;
+ T* destArg = (T*)kernelArg.m_argData;
+ *destArg = consts;
+ kernelArg.m_argSizeInBytes = sizeof(T);
+ m_kernelArguments.push_back(kernelArg);
+ m_serializationSizeInBytes += sizeof(b3KernelArgData);
}
- inline void launch2D( int numThreadsX, int numThreadsY, int localSizeX, int localSizeY )
- {
- size_t gRange[3] = {1,1,1};
- size_t lRange[3] = {1,1,1};
- lRange[0] = localSizeX;
- lRange[1] = localSizeY;
- gRange[0] = b3Max((size_t)1, (numThreadsX/lRange[0])+(!(numThreadsX%lRange[0])?0:1));
- gRange[0] *= lRange[0];
- gRange[1] = b3Max((size_t)1, (numThreadsY/lRange[1])+(!(numThreadsY%lRange[1])?0:1));
- gRange[1] *= lRange[1];
-
- cl_int status = clEnqueueNDRangeKernel( m_commandQueue,
- m_kernel, 2, NULL, gRange, lRange, 0,0,0 );
- if (status != CL_SUCCESS)
- {
- printf("Error: OpenCL status = %d\n",status);
- }
- b3Assert( status == CL_SUCCESS );
+ cl_int status = clSetKernelArg(m_kernel, m_idx++, sz, &consts);
+ b3Assert(status == CL_SUCCESS);
+ }
- }
-
- void enableSerialization(bool serialize)
+ inline void launch1D(int numThreads, int localSize = 64)
+ {
+ launch2D(numThreads, 1, localSize, 1);
+ }
+
+ inline void launch2D(int numThreadsX, int numThreadsY, int localSizeX, int localSizeY)
+ {
+ size_t gRange[3] = {1, 1, 1};
+ size_t lRange[3] = {1, 1, 1};
+ lRange[0] = localSizeX;
+ lRange[1] = localSizeY;
+ gRange[0] = b3Max((size_t)1, (numThreadsX / lRange[0]) + (!(numThreadsX % lRange[0]) ? 0 : 1));
+ gRange[0] *= lRange[0];
+ gRange[1] = b3Max((size_t)1, (numThreadsY / lRange[1]) + (!(numThreadsY % lRange[1]) ? 0 : 1));
+ gRange[1] *= lRange[1];
+
+ cl_int status = clEnqueueNDRangeKernel(m_commandQueue,
+ m_kernel, 2, NULL, gRange, lRange, 0, 0, 0);
+ if (status != CL_SUCCESS)
{
- m_enableSerialization = serialize;
+ printf("Error: OpenCL status = %d\n", status);
}
-
-};
-
+ b3Assert(status == CL_SUCCESS);
+ }
+ void enableSerialization(bool serialize)
+ {
+ m_enableSerialization = serialize;
+ }
+};
-#endif //B3_LAUNCHER_CL_H
+#endif //B3_LAUNCHER_CL_H
diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3OpenCLArray.h b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3OpenCLArray.h
index d70c30f53f..e837cceb66 100644
--- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3OpenCLArray.h
+++ b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3OpenCLArray.h
@@ -7,16 +7,16 @@
template <typename T>
class b3OpenCLArray
{
- size_t m_size;
- size_t m_capacity;
- cl_mem m_clBuffer;
+ size_t m_size;
+ size_t m_capacity;
+ cl_mem m_clBuffer;
- cl_context m_clContext;
+ cl_context m_clContext;
cl_command_queue m_commandQueue;
- bool m_ownsMemory;
+ bool m_ownsMemory;
- bool m_allowGrowingCapacity;
+ bool m_allowGrowingCapacity;
void deallocate()
{
@@ -25,22 +25,19 @@ class b3OpenCLArray
clReleaseMemObject(m_clBuffer);
}
m_clBuffer = 0;
- m_capacity=0;
+ m_capacity = 0;
}
b3OpenCLArray<T>& operator=(const b3OpenCLArray<T>& src);
- B3_FORCE_INLINE size_t allocSize(size_t size)
- {
- return (size ? size*2 : 1);
- }
+ B3_FORCE_INLINE size_t allocSize(size_t size)
+ {
+ return (size ? size * 2 : 1);
+ }
public:
-
- b3OpenCLArray(cl_context ctx, cl_command_queue queue, size_t initialCapacity=0, bool allowGrowingCapacity=true)
- :m_size(0), m_capacity(0),m_clBuffer(0),
- m_clContext(ctx),m_commandQueue(queue),
- m_ownsMemory(true),m_allowGrowingCapacity(true)
+ b3OpenCLArray(cl_context ctx, cl_command_queue queue, size_t initialCapacity = 0, bool allowGrowingCapacity = true)
+ : m_size(0), m_capacity(0), m_clBuffer(0), m_clContext(ctx), m_commandQueue(queue), m_ownsMemory(true), m_allowGrowingCapacity(true)
{
if (initialCapacity)
{
@@ -60,34 +57,32 @@ public:
m_capacity = sizeInElements;
}
-// we could enable this assignment, but need to make sure to avoid accidental deep copies
-// b3OpenCLArray<T>& operator=(const b3AlignedObjectArray<T>& src)
-// {
-// copyFromArray(src);
-// return *this;
-// }
+ // we could enable this assignment, but need to make sure to avoid accidental deep copies
+ // b3OpenCLArray<T>& operator=(const b3AlignedObjectArray<T>& src)
+ // {
+ // copyFromArray(src);
+ // return *this;
+ // }
-
- cl_mem getBufferCL() const
+ cl_mem getBufferCL() const
{
return m_clBuffer;
}
-
virtual ~b3OpenCLArray()
{
deallocate();
- m_size=0;
- m_capacity=0;
+ m_size = 0;
+ m_capacity = 0;
}
- B3_FORCE_INLINE bool push_back(const T& _Val,bool waitForCompletion=true)
+ B3_FORCE_INLINE bool push_back(const T& _Val, bool waitForCompletion = true)
{
bool result = true;
size_t sz = size();
- if( sz == capacity() )
+ if (sz == capacity())
{
- result = reserve( allocSize(size()) );
+ result = reserve(allocSize(size()));
}
copyFromHostPointer(&_Val, 1, sz, waitForCompletion);
m_size++;
@@ -96,23 +91,23 @@ public:
B3_FORCE_INLINE T forcedAt(size_t n) const
{
- b3Assert(n>=0);
- b3Assert(n<capacity());
+ b3Assert(n >= 0);
+ b3Assert(n < capacity());
T elem;
- copyToHostPointer(&elem,1,n,true);
+ copyToHostPointer(&elem, 1, n, true);
return elem;
}
B3_FORCE_INLINE T at(size_t n) const
{
- b3Assert(n>=0);
- b3Assert(n<size());
+ b3Assert(n >= 0);
+ b3Assert(n < size());
T elem;
- copyToHostPointer(&elem,1,n,true);
+ copyToHostPointer(&elem, 1, n, true);
return elem;
}
- B3_FORCE_INLINE bool resize(size_t newsize, bool copyOldContents=true)
+ B3_FORCE_INLINE bool resize(size_t newsize, bool copyOldContents = true)
{
bool result = true;
size_t curSize = size();
@@ -120,11 +115,12 @@ public:
if (newsize < curSize)
{
//leave the OpenCL memory for now
- } else
+ }
+ else
{
if (newsize > size())
{
- result = reserve(newsize,copyOldContents);
+ result = reserve(newsize, copyOldContents);
}
//leave new data uninitialized (init in debug mode?)
@@ -134,7 +130,8 @@ public:
if (result)
{
m_size = newsize;
- } else
+ }
+ else
{
m_size = 0;
}
@@ -146,25 +143,25 @@ public:
return m_size;
}
- B3_FORCE_INLINE size_t capacity() const
+ B3_FORCE_INLINE size_t capacity() const
{
return m_capacity;
}
- B3_FORCE_INLINE bool reserve(size_t _Count, bool copyOldContents=true)
+ B3_FORCE_INLINE bool reserve(size_t _Count, bool copyOldContents = true)
{
- bool result=true;
+ bool result = true;
// determine new minimum length of allocated storage
if (capacity() < _Count)
- { // not enough room, reallocate
+ { // not enough room, reallocate
if (m_allowGrowingCapacity)
{
cl_int ciErrNum;
//create a new OpenCL buffer
- size_t memSizeInBytes = sizeof(T)*_Count;
+ size_t memSizeInBytes = sizeof(T) * _Count;
cl_mem buf = clCreateBuffer(m_clContext, CL_MEM_READ_WRITE, memSizeInBytes, NULL, &ciErrNum);
- if (ciErrNum!=CL_SUCCESS)
+ if (ciErrNum != CL_SUCCESS)
{
b3Error("OpenCL out-of-memory\n");
_Count = 0;
@@ -173,13 +170,13 @@ public:
//#define B3_ALWAYS_INITIALIZE_OPENCL_BUFFERS
#ifdef B3_ALWAYS_INITIALIZE_OPENCL_BUFFERS
unsigned char* src = (unsigned char*)malloc(memSizeInBytes);
- for (size_t i=0;i<memSizeInBytes;i++)
+ for (size_t i = 0; i < memSizeInBytes; i++)
src[i] = 0xbb;
- ciErrNum = clEnqueueWriteBuffer( m_commandQueue, buf, CL_TRUE, 0, memSizeInBytes, src, 0,0,0 );
- b3Assert(ciErrNum==CL_SUCCESS);
+ ciErrNum = clEnqueueWriteBuffer(m_commandQueue, buf, CL_TRUE, 0, memSizeInBytes, src, 0, 0, 0);
+ b3Assert(ciErrNum == CL_SUCCESS);
clFinish(m_commandQueue);
free(src);
-#endif //B3_ALWAYS_INITIALIZE_OPENCL_BUFFERS
+#endif //B3_ALWAYS_INITIALIZE_OPENCL_BUFFERS
if (result)
{
@@ -193,21 +190,21 @@ public:
m_clBuffer = buf;
m_capacity = _Count;
- } else
+ }
+ else
{
//fail: assert and
b3Assert(0);
deallocate();
- result=false;
+ result = false;
}
}
return result;
}
-
- void copyToCL(cl_mem destination, size_t numElements, size_t firstElem=0, size_t dstOffsetInElems=0) const
+ void copyToCL(cl_mem destination, size_t numElements, size_t firstElem = 0, size_t dstOffsetInElems = 0) const
{
- if (numElements<=0)
+ if (numElements <= 0)
return;
b3Assert(m_clBuffer);
@@ -216,75 +213,74 @@ public:
//likely some error, destination is same as source
b3Assert(m_clBuffer != destination);
- b3Assert((firstElem+numElements)<=m_size);
+ b3Assert((firstElem + numElements) <= m_size);
cl_int status = 0;
+ b3Assert(numElements > 0);
+ b3Assert(numElements <= m_size);
- b3Assert(numElements>0);
- b3Assert(numElements<=m_size);
-
- size_t srcOffsetBytes = sizeof(T)*firstElem;
- size_t dstOffsetInBytes = sizeof(T)*dstOffsetInElems;
+ size_t srcOffsetBytes = sizeof(T) * firstElem;
+ size_t dstOffsetInBytes = sizeof(T) * dstOffsetInElems;
- status = clEnqueueCopyBuffer( m_commandQueue, m_clBuffer, destination,
- srcOffsetBytes, dstOffsetInBytes, sizeof(T)*numElements, 0, 0, 0 );
+ status = clEnqueueCopyBuffer(m_commandQueue, m_clBuffer, destination,
+ srcOffsetBytes, dstOffsetInBytes, sizeof(T) * numElements, 0, 0, 0);
- b3Assert( status == CL_SUCCESS );
+ b3Assert(status == CL_SUCCESS);
}
- void copyFromHost(const b3AlignedObjectArray<T>& srcArray, bool waitForCompletion=true)
+ void copyFromHost(const b3AlignedObjectArray<T>& srcArray, bool waitForCompletion = true)
{
size_t newSize = srcArray.size();
bool copyOldContents = false;
- resize (newSize,copyOldContents);
+ resize(newSize, copyOldContents);
if (newSize)
- copyFromHostPointer(&srcArray[0],newSize,0,waitForCompletion);
-
+ copyFromHostPointer(&srcArray[0], newSize, 0, waitForCompletion);
}
- void copyFromHostPointer(const T* src, size_t numElems, size_t destFirstElem= 0, bool waitForCompletion=true)
+ void copyFromHostPointer(const T* src, size_t numElems, size_t destFirstElem = 0, bool waitForCompletion = true)
{
- b3Assert(numElems+destFirstElem <= capacity());
+ b3Assert(numElems + destFirstElem <= capacity());
- if (numElems+destFirstElem)
+ if (numElems + destFirstElem)
{
cl_int status = 0;
- size_t sizeInBytes=sizeof(T)*numElems;
- status = clEnqueueWriteBuffer( m_commandQueue, m_clBuffer, 0, sizeof(T)*destFirstElem, sizeInBytes,
- src, 0,0,0 );
- b3Assert(status == CL_SUCCESS );
+ size_t sizeInBytes = sizeof(T) * numElems;
+ status = clEnqueueWriteBuffer(m_commandQueue, m_clBuffer, 0, sizeof(T) * destFirstElem, sizeInBytes,
+ src, 0, 0, 0);
+ b3Assert(status == CL_SUCCESS);
if (waitForCompletion)
clFinish(m_commandQueue);
- } else
+ }
+ else
{
b3Error("copyFromHostPointer invalid range\n");
}
}
-
- void copyToHost(b3AlignedObjectArray<T>& destArray, bool waitForCompletion=true) const
+ void copyToHost(b3AlignedObjectArray<T>& destArray, bool waitForCompletion = true) const
{
destArray.resize(this->size());
if (size())
- copyToHostPointer(&destArray[0], size(),0,waitForCompletion);
+ copyToHostPointer(&destArray[0], size(), 0, waitForCompletion);
}
- void copyToHostPointer(T* destPtr, size_t numElem, size_t srcFirstElem=0, bool waitForCompletion=true) const
+ void copyToHostPointer(T* destPtr, size_t numElem, size_t srcFirstElem = 0, bool waitForCompletion = true) const
{
- b3Assert(numElem+srcFirstElem <= capacity());
+ b3Assert(numElem + srcFirstElem <= capacity());
- if(numElem+srcFirstElem <= capacity())
+ if (numElem + srcFirstElem <= capacity())
{
cl_int status = 0;
- status = clEnqueueReadBuffer( m_commandQueue, m_clBuffer, 0, sizeof(T)*srcFirstElem, sizeof(T)*numElem,
- destPtr, 0,0,0 );
- b3Assert( status==CL_SUCCESS );
+ status = clEnqueueReadBuffer(m_commandQueue, m_clBuffer, 0, sizeof(T) * srcFirstElem, sizeof(T) * numElem,
+ destPtr, 0, 0, 0);
+ b3Assert(status == CL_SUCCESS);
if (waitForCompletion)
clFinish(m_commandQueue);
- } else
+ }
+ else
{
b3Error("copyToHostPointer invalid range\n");
}
@@ -296,11 +292,9 @@ public:
resize(newSize);
if (size())
{
- src.copyToCL(m_clBuffer,size());
+ src.copyToCL(m_clBuffer, size());
}
}
-
};
-
-#endif //B3_OPENCL_ARRAY_H
+#endif //B3_OPENCL_ARRAY_H
diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3PrefixScanCL.cpp b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3PrefixScanCL.cpp
index 42cd197740..822b511633 100644
--- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3PrefixScanCL.cpp
+++ b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3PrefixScanCL.cpp
@@ -7,25 +7,24 @@
#include "kernels/PrefixScanKernelsCL.h"
b3PrefixScanCL::b3PrefixScanCL(cl_context ctx, cl_device_id device, cl_command_queue queue, int size)
-:m_commandQueue(queue)
+ : m_commandQueue(queue)
{
const char* scanKernelSource = prefixScanKernelsCL;
cl_int pErrNum;
- char* additionalMacros=0;
+ char* additionalMacros = 0;
- m_workBuffer = new b3OpenCLArray<unsigned int>(ctx,queue,size);
- cl_program scanProg = b3OpenCLUtils::compileCLProgramFromString( ctx, device, scanKernelSource, &pErrNum,additionalMacros, B3_PREFIXSCAN_PROG_PATH);
+ m_workBuffer = new b3OpenCLArray<unsigned int>(ctx, queue, size);
+ cl_program scanProg = b3OpenCLUtils::compileCLProgramFromString(ctx, device, scanKernelSource, &pErrNum, additionalMacros, B3_PREFIXSCAN_PROG_PATH);
b3Assert(scanProg);
- m_localScanKernel = b3OpenCLUtils::compileCLKernelFromString( ctx, device, scanKernelSource, "LocalScanKernel", &pErrNum, scanProg,additionalMacros );
- b3Assert(m_localScanKernel );
- m_blockSumKernel = b3OpenCLUtils::compileCLKernelFromString( ctx, device, scanKernelSource, "TopLevelScanKernel", &pErrNum, scanProg,additionalMacros );
- b3Assert(m_blockSumKernel );
- m_propagationKernel = b3OpenCLUtils::compileCLKernelFromString( ctx, device, scanKernelSource, "AddOffsetKernel", &pErrNum, scanProg,additionalMacros );
- b3Assert(m_propagationKernel );
+ m_localScanKernel = b3OpenCLUtils::compileCLKernelFromString(ctx, device, scanKernelSource, "LocalScanKernel", &pErrNum, scanProg, additionalMacros);
+ b3Assert(m_localScanKernel);
+ m_blockSumKernel = b3OpenCLUtils::compileCLKernelFromString(ctx, device, scanKernelSource, "TopLevelScanKernel", &pErrNum, scanProg, additionalMacros);
+ b3Assert(m_blockSumKernel);
+ m_propagationKernel = b3OpenCLUtils::compileCLKernelFromString(ctx, device, scanKernelSource, "AddOffsetKernel", &pErrNum, scanProg, additionalMacros);
+ b3Assert(m_propagationKernel);
}
-
b3PrefixScanCL::~b3PrefixScanCL()
{
delete m_workBuffer;
@@ -34,20 +33,19 @@ b3PrefixScanCL::~b3PrefixScanCL()
clReleaseKernel(m_propagationKernel);
}
-template<class T>
+template <class T>
T b3NextPowerOf2(T n)
{
n -= 1;
- for(int i=0; i<sizeof(T)*8; i++)
- n = n | (n>>i);
- return n+1;
+ for (int i = 0; i < sizeof(T) * 8; i++)
+ n = n | (n >> i);
+ return n + 1;
}
void b3PrefixScanCL::execute(b3OpenCLArray<unsigned int>& src, b3OpenCLArray<unsigned int>& dst, int n, unsigned int* sum)
{
-
-// b3Assert( data->m_option == EXCLUSIVE );
- const unsigned int numBlocks = (const unsigned int)( (n+BLOCK_SIZE*2-1)/(BLOCK_SIZE*2) );
+ // b3Assert( data->m_option == EXCLUSIVE );
+ const unsigned int numBlocks = (const unsigned int)((n + BLOCK_SIZE * 2 - 1) / (BLOCK_SIZE * 2));
dst.resize(src.size());
m_workBuffer->resize(src.size());
@@ -55,55 +53,51 @@ void b3PrefixScanCL::execute(b3OpenCLArray<unsigned int>& src, b3OpenCLArray<uns
b3Int4 constBuffer;
constBuffer.x = n;
constBuffer.y = numBlocks;
- constBuffer.z = (int)b3NextPowerOf2( numBlocks );
+ constBuffer.z = (int)b3NextPowerOf2(numBlocks);
b3OpenCLArray<unsigned int>* srcNative = &src;
b3OpenCLArray<unsigned int>* dstNative = &dst;
-
+
{
- b3BufferInfoCL bInfo[] = { b3BufferInfoCL( dstNative->getBufferCL() ), b3BufferInfoCL( srcNative->getBufferCL() ), b3BufferInfoCL( m_workBuffer->getBufferCL() ) };
+ b3BufferInfoCL bInfo[] = {b3BufferInfoCL(dstNative->getBufferCL()), b3BufferInfoCL(srcNative->getBufferCL()), b3BufferInfoCL(m_workBuffer->getBufferCL())};
- b3LauncherCL launcher( m_commandQueue, m_localScanKernel,"m_localScanKernel" );
- launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
- launcher.setConst( constBuffer );
- launcher.launch1D( numBlocks*BLOCK_SIZE, BLOCK_SIZE );
+ b3LauncherCL launcher(m_commandQueue, m_localScanKernel, "m_localScanKernel");
+ launcher.setBuffers(bInfo, sizeof(bInfo) / sizeof(b3BufferInfoCL));
+ launcher.setConst(constBuffer);
+ launcher.launch1D(numBlocks * BLOCK_SIZE, BLOCK_SIZE);
}
{
- b3BufferInfoCL bInfo[] = { b3BufferInfoCL( m_workBuffer->getBufferCL() ) };
+ b3BufferInfoCL bInfo[] = {b3BufferInfoCL(m_workBuffer->getBufferCL())};
- b3LauncherCL launcher( m_commandQueue, m_blockSumKernel,"m_blockSumKernel" );
- launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
- launcher.setConst( constBuffer );
- launcher.launch1D( BLOCK_SIZE, BLOCK_SIZE );
+ b3LauncherCL launcher(m_commandQueue, m_blockSumKernel, "m_blockSumKernel");
+ launcher.setBuffers(bInfo, sizeof(bInfo) / sizeof(b3BufferInfoCL));
+ launcher.setConst(constBuffer);
+ launcher.launch1D(BLOCK_SIZE, BLOCK_SIZE);
}
-
- if( numBlocks > 1 )
+ if (numBlocks > 1)
{
- b3BufferInfoCL bInfo[] = { b3BufferInfoCL( dstNative->getBufferCL() ), b3BufferInfoCL( m_workBuffer->getBufferCL() ) };
- b3LauncherCL launcher( m_commandQueue, m_propagationKernel,"m_propagationKernel" );
- launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
- launcher.setConst( constBuffer );
- launcher.launch1D( (numBlocks-1)*BLOCK_SIZE, BLOCK_SIZE );
+ b3BufferInfoCL bInfo[] = {b3BufferInfoCL(dstNative->getBufferCL()), b3BufferInfoCL(m_workBuffer->getBufferCL())};
+ b3LauncherCL launcher(m_commandQueue, m_propagationKernel, "m_propagationKernel");
+ launcher.setBuffers(bInfo, sizeof(bInfo) / sizeof(b3BufferInfoCL));
+ launcher.setConst(constBuffer);
+ launcher.launch1D((numBlocks - 1) * BLOCK_SIZE, BLOCK_SIZE);
}
-
- if( sum )
+ if (sum)
{
clFinish(m_commandQueue);
- dstNative->copyToHostPointer(sum,1,n-1,true);
+ dstNative->copyToHostPointer(sum, 1, n - 1, true);
}
-
}
-
void b3PrefixScanCL::executeHost(b3AlignedObjectArray<unsigned int>& src, b3AlignedObjectArray<unsigned int>& dst, int n, unsigned int* sum)
{
unsigned int s = 0;
//if( data->m_option == EXCLUSIVE )
{
- for(int i=0; i<n; i++)
+ for (int i = 0; i < n; i++)
{
dst[i] = s;
s += src[i];
@@ -119,8 +113,8 @@ void b3PrefixScanCL::executeHost(b3AlignedObjectArray<unsigned int>& src, b3Alig
}
*/
- if( sum )
+ if (sum)
{
- *sum = dst[n-1];
+ *sum = dst[n - 1];
}
} \ No newline at end of file
diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3PrefixScanCL.h b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3PrefixScanCL.h
index a9a2e61b9e..346efa0c73 100644
--- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3PrefixScanCL.h
+++ b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3PrefixScanCL.h
@@ -13,9 +13,9 @@ class b3PrefixScanCL
BLOCK_SIZE = 128
};
-// Option m_option;
+ // Option m_option;
- cl_command_queue m_commandQueue;
+ cl_command_queue m_commandQueue;
cl_kernel m_localScanKernel;
cl_kernel m_blockSumKernel;
@@ -23,15 +23,13 @@ class b3PrefixScanCL
b3OpenCLArray<unsigned int>* m_workBuffer;
-
- public:
-
- b3PrefixScanCL(cl_context ctx, cl_device_id device, cl_command_queue queue,int size=0);
+public:
+ b3PrefixScanCL(cl_context ctx, cl_device_id device, cl_command_queue queue, int size = 0);
virtual ~b3PrefixScanCL();
void execute(b3OpenCLArray<unsigned int>& src, b3OpenCLArray<unsigned int>& dst, int n, unsigned int* sum = 0);
- void executeHost(b3AlignedObjectArray<unsigned int>& src, b3AlignedObjectArray<unsigned int>& dst, int n, unsigned int* sum=0);
+ void executeHost(b3AlignedObjectArray<unsigned int>& src, b3AlignedObjectArray<unsigned int>& dst, int n, unsigned int* sum = 0);
};
-#endif //B3_PREFIX_SCAN_CL_H
+#endif //B3_PREFIX_SCAN_CL_H
diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3PrefixScanFloat4CL.cpp b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3PrefixScanFloat4CL.cpp
index 80560d793d..1cac97c988 100644
--- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3PrefixScanFloat4CL.cpp
+++ b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3PrefixScanFloat4CL.cpp
@@ -7,25 +7,24 @@
#include "kernels/PrefixScanKernelsFloat4CL.h"
b3PrefixScanFloat4CL::b3PrefixScanFloat4CL(cl_context ctx, cl_device_id device, cl_command_queue queue, int size)
-:m_commandQueue(queue)
+ : m_commandQueue(queue)
{
const char* scanKernelSource = prefixScanKernelsFloat4CL;
cl_int pErrNum;
- char* additionalMacros=0;
+ char* additionalMacros = 0;
- m_workBuffer = new b3OpenCLArray<b3Vector3>(ctx,queue,size);
- cl_program scanProg = b3OpenCLUtils::compileCLProgramFromString( ctx, device, scanKernelSource, &pErrNum,additionalMacros, B3_PREFIXSCAN_FLOAT4_PROG_PATH);
+ m_workBuffer = new b3OpenCLArray<b3Vector3>(ctx, queue, size);
+ cl_program scanProg = b3OpenCLUtils::compileCLProgramFromString(ctx, device, scanKernelSource, &pErrNum, additionalMacros, B3_PREFIXSCAN_FLOAT4_PROG_PATH);
b3Assert(scanProg);
- m_localScanKernel = b3OpenCLUtils::compileCLKernelFromString( ctx, device, scanKernelSource, "LocalScanKernel", &pErrNum, scanProg,additionalMacros );
- b3Assert(m_localScanKernel );
- m_blockSumKernel = b3OpenCLUtils::compileCLKernelFromString( ctx, device, scanKernelSource, "TopLevelScanKernel", &pErrNum, scanProg,additionalMacros );
- b3Assert(m_blockSumKernel );
- m_propagationKernel = b3OpenCLUtils::compileCLKernelFromString( ctx, device, scanKernelSource, "AddOffsetKernel", &pErrNum, scanProg,additionalMacros );
- b3Assert(m_propagationKernel );
+ m_localScanKernel = b3OpenCLUtils::compileCLKernelFromString(ctx, device, scanKernelSource, "LocalScanKernel", &pErrNum, scanProg, additionalMacros);
+ b3Assert(m_localScanKernel);
+ m_blockSumKernel = b3OpenCLUtils::compileCLKernelFromString(ctx, device, scanKernelSource, "TopLevelScanKernel", &pErrNum, scanProg, additionalMacros);
+ b3Assert(m_blockSumKernel);
+ m_propagationKernel = b3OpenCLUtils::compileCLKernelFromString(ctx, device, scanKernelSource, "AddOffsetKernel", &pErrNum, scanProg, additionalMacros);
+ b3Assert(m_propagationKernel);
}
-
b3PrefixScanFloat4CL::~b3PrefixScanFloat4CL()
{
delete m_workBuffer;
@@ -34,20 +33,19 @@ b3PrefixScanFloat4CL::~b3PrefixScanFloat4CL()
clReleaseKernel(m_propagationKernel);
}
-template<class T>
+template <class T>
T b3NextPowerOf2(T n)
{
n -= 1;
- for(int i=0; i<sizeof(T)*8; i++)
- n = n | (n>>i);
- return n+1;
+ for (int i = 0; i < sizeof(T) * 8; i++)
+ n = n | (n >> i);
+ return n + 1;
}
void b3PrefixScanFloat4CL::execute(b3OpenCLArray<b3Vector3>& src, b3OpenCLArray<b3Vector3>& dst, int n, b3Vector3* sum)
{
-
-// b3Assert( data->m_option == EXCLUSIVE );
- const unsigned int numBlocks = (const unsigned int)( (n+BLOCK_SIZE*2-1)/(BLOCK_SIZE*2) );
+ // b3Assert( data->m_option == EXCLUSIVE );
+ const unsigned int numBlocks = (const unsigned int)((n + BLOCK_SIZE * 2 - 1) / (BLOCK_SIZE * 2));
dst.resize(src.size());
m_workBuffer->resize(src.size());
@@ -55,55 +53,51 @@ void b3PrefixScanFloat4CL::execute(b3OpenCLArray<b3Vector3>& src, b3OpenCLArray<
b3Int4 constBuffer;
constBuffer.x = n;
constBuffer.y = numBlocks;
- constBuffer.z = (int)b3NextPowerOf2( numBlocks );
+ constBuffer.z = (int)b3NextPowerOf2(numBlocks);
b3OpenCLArray<b3Vector3>* srcNative = &src;
b3OpenCLArray<b3Vector3>* dstNative = &dst;
-
+
{
- b3BufferInfoCL bInfo[] = { b3BufferInfoCL( dstNative->getBufferCL() ), b3BufferInfoCL( srcNative->getBufferCL() ), b3BufferInfoCL( m_workBuffer->getBufferCL() ) };
+ b3BufferInfoCL bInfo[] = {b3BufferInfoCL(dstNative->getBufferCL()), b3BufferInfoCL(srcNative->getBufferCL()), b3BufferInfoCL(m_workBuffer->getBufferCL())};
- b3LauncherCL launcher( m_commandQueue, m_localScanKernel ,"m_localScanKernel");
- launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
- launcher.setConst( constBuffer );
- launcher.launch1D( numBlocks*BLOCK_SIZE, BLOCK_SIZE );
+ b3LauncherCL launcher(m_commandQueue, m_localScanKernel, "m_localScanKernel");
+ launcher.setBuffers(bInfo, sizeof(bInfo) / sizeof(b3BufferInfoCL));
+ launcher.setConst(constBuffer);
+ launcher.launch1D(numBlocks * BLOCK_SIZE, BLOCK_SIZE);
}
{
- b3BufferInfoCL bInfo[] = { b3BufferInfoCL( m_workBuffer->getBufferCL() ) };
+ b3BufferInfoCL bInfo[] = {b3BufferInfoCL(m_workBuffer->getBufferCL())};
- b3LauncherCL launcher( m_commandQueue, m_blockSumKernel ,"m_blockSumKernel");
- launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
- launcher.setConst( constBuffer );
- launcher.launch1D( BLOCK_SIZE, BLOCK_SIZE );
+ b3LauncherCL launcher(m_commandQueue, m_blockSumKernel, "m_blockSumKernel");
+ launcher.setBuffers(bInfo, sizeof(bInfo) / sizeof(b3BufferInfoCL));
+ launcher.setConst(constBuffer);
+ launcher.launch1D(BLOCK_SIZE, BLOCK_SIZE);
}
-
- if( numBlocks > 1 )
+ if (numBlocks > 1)
{
- b3BufferInfoCL bInfo[] = { b3BufferInfoCL( dstNative->getBufferCL() ), b3BufferInfoCL( m_workBuffer->getBufferCL() ) };
- b3LauncherCL launcher( m_commandQueue, m_propagationKernel ,"m_propagationKernel");
- launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
- launcher.setConst( constBuffer );
- launcher.launch1D( (numBlocks-1)*BLOCK_SIZE, BLOCK_SIZE );
+ b3BufferInfoCL bInfo[] = {b3BufferInfoCL(dstNative->getBufferCL()), b3BufferInfoCL(m_workBuffer->getBufferCL())};
+ b3LauncherCL launcher(m_commandQueue, m_propagationKernel, "m_propagationKernel");
+ launcher.setBuffers(bInfo, sizeof(bInfo) / sizeof(b3BufferInfoCL));
+ launcher.setConst(constBuffer);
+ launcher.launch1D((numBlocks - 1) * BLOCK_SIZE, BLOCK_SIZE);
}
-
- if( sum )
+ if (sum)
{
clFinish(m_commandQueue);
- dstNative->copyToHostPointer(sum,1,n-1,true);
+ dstNative->copyToHostPointer(sum, 1, n - 1, true);
}
-
}
-
void b3PrefixScanFloat4CL::executeHost(b3AlignedObjectArray<b3Vector3>& src, b3AlignedObjectArray<b3Vector3>& dst, int n, b3Vector3* sum)
{
- b3Vector3 s=b3MakeVector3(0,0,0);
+ b3Vector3 s = b3MakeVector3(0, 0, 0);
//if( data->m_option == EXCLUSIVE )
{
- for(int i=0; i<n; i++)
+ for (int i = 0; i < n; i++)
{
dst[i] = s;
s += src[i];
@@ -119,8 +113,8 @@ void b3PrefixScanFloat4CL::executeHost(b3AlignedObjectArray<b3Vector3>& src, b3A
}
*/
- if( sum )
+ if (sum)
{
- *sum = dst[n-1];
+ *sum = dst[n - 1];
}
} \ No newline at end of file
diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3PrefixScanFloat4CL.h b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3PrefixScanFloat4CL.h
index 2c8003c1bb..122b0bfd68 100644
--- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3PrefixScanFloat4CL.h
+++ b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3PrefixScanFloat4CL.h
@@ -14,9 +14,9 @@ class b3PrefixScanFloat4CL
BLOCK_SIZE = 128
};
-// Option m_option;
+ // Option m_option;
- cl_command_queue m_commandQueue;
+ cl_command_queue m_commandQueue;
cl_kernel m_localScanKernel;
cl_kernel m_blockSumKernel;
@@ -24,10 +24,8 @@ class b3PrefixScanFloat4CL
b3OpenCLArray<b3Vector3>* m_workBuffer;
-
- public:
-
- b3PrefixScanFloat4CL(cl_context ctx, cl_device_id device, cl_command_queue queue,int size=0);
+public:
+ b3PrefixScanFloat4CL(cl_context ctx, cl_device_id device, cl_command_queue queue, int size = 0);
virtual ~b3PrefixScanFloat4CL();
@@ -35,4 +33,4 @@ class b3PrefixScanFloat4CL
void executeHost(b3AlignedObjectArray<b3Vector3>& src, b3AlignedObjectArray<b3Vector3>& dst, int n, b3Vector3* sum);
};
-#endif //B3_PREFIX_SCAN_CL_H
+#endif //B3_PREFIX_SCAN_CL_H
diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3RadixSort32CL.cpp b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3RadixSort32CL.cpp
index f11ae4bcdb..e86af6583f 100644
--- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3RadixSort32CL.cpp
+++ b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3RadixSort32CL.cpp
@@ -10,21 +10,20 @@
#include "kernels/RadixSort32KernelsCL.h"
b3RadixSort32CL::b3RadixSort32CL(cl_context ctx, cl_device_id device, cl_command_queue queue, int initialCapacity)
-:m_commandQueue(queue)
+ : m_commandQueue(queue)
{
b3OpenCLDeviceInfo info;
- b3OpenCLUtils::getDeviceInfo(device,&info);
- m_deviceCPU = (info.m_deviceType & CL_DEVICE_TYPE_CPU)!=0;
+ b3OpenCLUtils::getDeviceInfo(device, &info);
+ m_deviceCPU = (info.m_deviceType & CL_DEVICE_TYPE_CPU) != 0;
- m_workBuffer1 = new b3OpenCLArray<unsigned int>(ctx,queue);
- m_workBuffer2 = new b3OpenCLArray<unsigned int>(ctx,queue);
- m_workBuffer3 = new b3OpenCLArray<b3SortData>(ctx,queue);
- m_workBuffer3a = new b3OpenCLArray<unsigned int>(ctx,queue);
- m_workBuffer4 = new b3OpenCLArray<b3SortData>(ctx,queue);
- m_workBuffer4a = new b3OpenCLArray<unsigned int>(ctx,queue);
+ m_workBuffer1 = new b3OpenCLArray<unsigned int>(ctx, queue);
+ m_workBuffer2 = new b3OpenCLArray<unsigned int>(ctx, queue);
+ m_workBuffer3 = new b3OpenCLArray<b3SortData>(ctx, queue);
+ m_workBuffer3a = new b3OpenCLArray<unsigned int>(ctx, queue);
+ m_workBuffer4 = new b3OpenCLArray<b3SortData>(ctx, queue);
+ m_workBuffer4a = new b3OpenCLArray<unsigned int>(ctx, queue);
-
- if (initialCapacity>0)
+ if (initialCapacity > 0)
{
m_workBuffer1->resize(initialCapacity);
m_workBuffer3->resize(initialCapacity);
@@ -33,45 +32,40 @@ b3RadixSort32CL::b3RadixSort32CL(cl_context ctx, cl_device_id device, cl_command
m_workBuffer4a->resize(initialCapacity);
}
- m_scan = new b3PrefixScanCL(ctx,device,queue);
- m_fill = new b3FillCL(ctx,device,queue);
-
+ m_scan = new b3PrefixScanCL(ctx, device, queue);
+ m_fill = new b3FillCL(ctx, device, queue);
+
const char* additionalMacros = "";
cl_int pErrNum;
const char* kernelSource = radixSort32KernelsCL;
-
- cl_program sortProg = b3OpenCLUtils::compileCLProgramFromString( ctx, device, kernelSource, &pErrNum,additionalMacros, RADIXSORT32_PATH);
- b3Assert(sortProg);
- m_streamCountSortDataKernel = b3OpenCLUtils::compileCLKernelFromString( ctx, device, kernelSource, "StreamCountSortDataKernel", &pErrNum, sortProg,additionalMacros );
- b3Assert(m_streamCountSortDataKernel );
+ cl_program sortProg = b3OpenCLUtils::compileCLProgramFromString(ctx, device, kernelSource, &pErrNum, additionalMacros, RADIXSORT32_PATH);
+ b3Assert(sortProg);
+ m_streamCountSortDataKernel = b3OpenCLUtils::compileCLKernelFromString(ctx, device, kernelSource, "StreamCountSortDataKernel", &pErrNum, sortProg, additionalMacros);
+ b3Assert(m_streamCountSortDataKernel);
-
- m_streamCountKernel = b3OpenCLUtils::compileCLKernelFromString( ctx, device, kernelSource, "StreamCountKernel", &pErrNum, sortProg,additionalMacros );
+ m_streamCountKernel = b3OpenCLUtils::compileCLKernelFromString(ctx, device, kernelSource, "StreamCountKernel", &pErrNum, sortProg, additionalMacros);
b3Assert(m_streamCountKernel);
-
-
if (m_deviceCPU)
{
-
- m_sortAndScatterSortDataKernel = b3OpenCLUtils::compileCLKernelFromString( ctx, device, kernelSource, "SortAndScatterSortDataKernelSerial", &pErrNum, sortProg,additionalMacros );
+ m_sortAndScatterSortDataKernel = b3OpenCLUtils::compileCLKernelFromString(ctx, device, kernelSource, "SortAndScatterSortDataKernelSerial", &pErrNum, sortProg, additionalMacros);
b3Assert(m_sortAndScatterSortDataKernel);
- m_sortAndScatterKernel = b3OpenCLUtils::compileCLKernelFromString( ctx, device, kernelSource, "SortAndScatterKernelSerial", &pErrNum, sortProg,additionalMacros );
+ m_sortAndScatterKernel = b3OpenCLUtils::compileCLKernelFromString(ctx, device, kernelSource, "SortAndScatterKernelSerial", &pErrNum, sortProg, additionalMacros);
b3Assert(m_sortAndScatterKernel);
- } else
+ }
+ else
{
- m_sortAndScatterSortDataKernel = b3OpenCLUtils::compileCLKernelFromString( ctx, device, kernelSource, "SortAndScatterSortDataKernel", &pErrNum, sortProg,additionalMacros );
+ m_sortAndScatterSortDataKernel = b3OpenCLUtils::compileCLKernelFromString(ctx, device, kernelSource, "SortAndScatterSortDataKernel", &pErrNum, sortProg, additionalMacros);
b3Assert(m_sortAndScatterSortDataKernel);
- m_sortAndScatterKernel = b3OpenCLUtils::compileCLKernelFromString( ctx, device, kernelSource, "SortAndScatterKernel", &pErrNum, sortProg,additionalMacros );
+ m_sortAndScatterKernel = b3OpenCLUtils::compileCLKernelFromString(ctx, device, kernelSource, "SortAndScatterKernel", &pErrNum, sortProg, additionalMacros);
b3Assert(m_sortAndScatterKernel);
}
-
- m_prefixScanKernel = b3OpenCLUtils::compileCLKernelFromString( ctx, device, kernelSource, "PrefixScanKernel", &pErrNum, sortProg,additionalMacros );
+
+ m_prefixScanKernel = b3OpenCLUtils::compileCLKernelFromString(ctx, device, kernelSource, "PrefixScanKernel", &pErrNum, sortProg, additionalMacros);
b3Assert(m_prefixScanKernel);
-
}
b3RadixSort32CL::~b3RadixSort32CL()
@@ -96,8 +90,7 @@ void b3RadixSort32CL::executeHost(b3AlignedObjectArray<b3SortData>& inout, int s
{
int n = inout.size();
const int BITS_PER_PASS = 8;
- const int NUM_TABLES = (1<<BITS_PER_PASS);
-
+ const int NUM_TABLES = (1 << BITS_PER_PASS);
int tables[NUM_TABLES];
int counter[NUM_TABLES];
@@ -107,34 +100,33 @@ void b3RadixSort32CL::executeHost(b3AlignedObjectArray<b3SortData>& inout, int s
workbuffer.resize(inout.size());
b3SortData* dst = &workbuffer[0];
- int count=0;
- for(int startBit=0; startBit<sortBits; startBit+=BITS_PER_PASS)
+ int count = 0;
+ for (int startBit = 0; startBit < sortBits; startBit += BITS_PER_PASS)
{
- for(int i=0; i<NUM_TABLES; i++)
+ for (int i = 0; i < NUM_TABLES; i++)
{
tables[i] = 0;
}
- for(int i=0; i<n; i++)
+ for (int i = 0; i < n; i++)
{
- int tableIdx = (src[i].m_key >> startBit) & (NUM_TABLES-1);
+ int tableIdx = (src[i].m_key >> startBit) & (NUM_TABLES - 1);
tables[tableIdx]++;
}
//#define TEST
#ifdef TEST
- printf("histogram size=%d\n",NUM_TABLES);
- for (int i=0;i<NUM_TABLES;i++)
+ printf("histogram size=%d\n", NUM_TABLES);
+ for (int i = 0; i < NUM_TABLES; i++)
{
- if (tables[i]!=0)
+ if (tables[i] != 0)
{
- printf("tables[%d]=%d]\n",i,tables[i]);
+ printf("tables[%d]=%d]\n", i, tables[i]);
}
-
}
-#endif //TEST
- // prefix scan
+#endif //TEST \
+ // prefix scan
int sum = 0;
- for(int i=0; i<NUM_TABLES; i++)
+ for (int i = 0; i < NUM_TABLES; i++)
{
int iData = tables[i];
tables[i] = sum;
@@ -143,71 +135,65 @@ void b3RadixSort32CL::executeHost(b3AlignedObjectArray<b3SortData>& inout, int s
}
// distribute
- for(int i=0; i<n; i++)
+ for (int i = 0; i < n; i++)
{
- int tableIdx = (src[i].m_key >> startBit) & (NUM_TABLES-1);
-
+ int tableIdx = (src[i].m_key >> startBit) & (NUM_TABLES - 1);
+
dst[tables[tableIdx] + counter[tableIdx]] = src[i];
- counter[tableIdx] ++;
+ counter[tableIdx]++;
}
- b3Swap( src, dst );
+ b3Swap(src, dst);
count++;
}
- if (count&1)
+ if (count & 1)
{
- b3Assert(0);//need to copy
-
+ b3Assert(0); //need to copy
}
}
void b3RadixSort32CL::executeHost(b3OpenCLArray<b3SortData>& keyValuesInOut, int sortBits /* = 32 */)
{
-
b3AlignedObjectArray<b3SortData> inout;
keyValuesInOut.copyToHost(inout);
- executeHost(inout,sortBits);
+ executeHost(inout, sortBits);
keyValuesInOut.copyFromHost(inout);
}
-void b3RadixSort32CL::execute(b3OpenCLArray<unsigned int>& keysIn, b3OpenCLArray<unsigned int>& keysOut, b3OpenCLArray<unsigned int>& valuesIn,
- b3OpenCLArray<unsigned int>& valuesOut, int n, int sortBits)
+void b3RadixSort32CL::execute(b3OpenCLArray<unsigned int>& keysIn, b3OpenCLArray<unsigned int>& keysOut, b3OpenCLArray<unsigned int>& valuesIn,
+ b3OpenCLArray<unsigned int>& valuesOut, int n, int sortBits)
{
-
}
//#define DEBUG_RADIXSORT
//#define DEBUG_RADIXSORT2
-
void b3RadixSort32CL::execute(b3OpenCLArray<b3SortData>& keyValuesInOut, int sortBits /* = 32 */)
{
-
int originalSize = keyValuesInOut.size();
int workingSize = originalSize;
-
-
+
int dataAlignment = DATA_ALIGNMENT;
#ifdef DEBUG_RADIXSORT2
- b3AlignedObjectArray<b3SortData> test2;
- keyValuesInOut.copyToHost(test2);
- printf("numElem = %d\n",test2.size());
- for (int i=0;i<test2.size();i++)
- {
- printf("test2[%d].m_key=%d\n",i,test2[i].m_key);
- printf("test2[%d].m_value=%d\n",i,test2[i].m_value);
- }
-#endif //DEBUG_RADIXSORT2
-
+ b3AlignedObjectArray<b3SortData> test2;
+ keyValuesInOut.copyToHost(test2);
+ printf("numElem = %d\n", test2.size());
+ for (int i = 0; i < test2.size(); i++)
+ {
+ printf("test2[%d].m_key=%d\n", i, test2[i].m_key);
+ printf("test2[%d].m_value=%d\n", i, test2[i].m_value);
+ }
+#endif //DEBUG_RADIXSORT2
+
b3OpenCLArray<b3SortData>* src = 0;
- if (workingSize%dataAlignment)
+ if (workingSize % dataAlignment)
{
- workingSize += dataAlignment-(workingSize%dataAlignment);
+ workingSize += dataAlignment - (workingSize % dataAlignment);
m_workBuffer4->copyFromOpenCLArray(keyValuesInOut);
m_workBuffer4->resize(workingSize);
b3SortData fillValue;
@@ -216,327 +202,301 @@ void b3RadixSort32CL::execute(b3OpenCLArray<b3SortData>& keyValuesInOut, int sor
#define USE_BTFILL
#ifdef USE_BTFILL
- m_fill->execute((b3OpenCLArray<b3Int2>&)*m_workBuffer4,(b3Int2&)fillValue,workingSize-originalSize,originalSize);
+ m_fill->execute((b3OpenCLArray<b3Int2>&)*m_workBuffer4, (b3Int2&)fillValue, workingSize - originalSize, originalSize);
#else
//fill the remaining bits (very slow way, todo: fill on GPU/OpenCL side)
-
- for (int i=originalSize; i<workingSize;i++)
+
+ for (int i = originalSize; i < workingSize; i++)
{
- m_workBuffer4->copyFromHostPointer(&fillValue,1,i);
+ m_workBuffer4->copyFromHostPointer(&fillValue, 1, i);
}
-#endif//USE_BTFILL
+#endif //USE_BTFILL
src = m_workBuffer4;
- } else
+ }
+ else
{
src = &keyValuesInOut;
m_workBuffer4->resize(0);
}
-
- b3Assert( workingSize%DATA_ALIGNMENT == 0 );
- int minCap = NUM_BUCKET*NUM_WGS;
+ b3Assert(workingSize % DATA_ALIGNMENT == 0);
+ int minCap = NUM_BUCKET * NUM_WGS;
int n = workingSize;
m_workBuffer1->resize(minCap);
m_workBuffer3->resize(workingSize);
-
-// ADLASSERT( ELEMENTS_PER_WORK_ITEM == 4 );
- b3Assert( BITS_PER_PASS == 4 );
- b3Assert( WG_SIZE == 64 );
- b3Assert( (sortBits&0x3) == 0 );
+ // ADLASSERT( ELEMENTS_PER_WORK_ITEM == 4 );
+ b3Assert(BITS_PER_PASS == 4);
+ b3Assert(WG_SIZE == 64);
+ b3Assert((sortBits & 0x3) == 0);
-
-
b3OpenCLArray<b3SortData>* dst = m_workBuffer3;
b3OpenCLArray<unsigned int>* srcHisto = m_workBuffer1;
b3OpenCLArray<unsigned int>* destHisto = m_workBuffer2;
-
int nWGs = NUM_WGS;
b3ConstData cdata;
{
- int blockSize = ELEMENTS_PER_WORK_ITEM*WG_SIZE;//set at 256
- int nBlocks = (n+blockSize-1)/(blockSize);
+ int blockSize = ELEMENTS_PER_WORK_ITEM * WG_SIZE; //set at 256
+ int nBlocks = (n + blockSize - 1) / (blockSize);
cdata.m_n = n;
cdata.m_nWGs = NUM_WGS;
cdata.m_startBit = 0;
- cdata.m_nBlocksPerWG = (nBlocks + cdata.m_nWGs - 1)/cdata.m_nWGs;
- if( nBlocks < NUM_WGS )
+ cdata.m_nBlocksPerWG = (nBlocks + cdata.m_nWGs - 1) / cdata.m_nWGs;
+ if (nBlocks < NUM_WGS)
{
cdata.m_nBlocksPerWG = 1;
nWGs = nBlocks;
}
}
- int count=0;
- for(int ib=0; ib<sortBits; ib+=4)
+ int count = 0;
+ for (int ib = 0; ib < sortBits; ib += 4)
{
#ifdef DEBUG_RADIXSORT2
- keyValuesInOut.copyToHost(test2);
- printf("numElem = %d\n",test2.size());
- for (int i=0;i<test2.size();i++)
- {
- if (test2[i].m_key != test2[i].m_value)
- {
- printf("test2[%d].m_key=%d\n",i,test2[i].m_key);
- printf("test2[%d].m_value=%d\n",i,test2[i].m_value);
- }
- }
-#endif //DEBUG_RADIXSORT2
-
+ keyValuesInOut.copyToHost(test2);
+ printf("numElem = %d\n", test2.size());
+ for (int i = 0; i < test2.size(); i++)
+ {
+ if (test2[i].m_key != test2[i].m_value)
+ {
+ printf("test2[%d].m_key=%d\n", i, test2[i].m_key);
+ printf("test2[%d].m_value=%d\n", i, test2[i].m_value);
+ }
+ }
+#endif //DEBUG_RADIXSORT2
+
cdata.m_startBit = ib;
-
+
if (src->size())
{
- b3BufferInfoCL bInfo[] = { b3BufferInfoCL( src->getBufferCL(), true ), b3BufferInfoCL( srcHisto->getBufferCL() ) };
- b3LauncherCL launcher(m_commandQueue, m_streamCountSortDataKernel,"m_streamCountSortDataKernel");
-
- launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
- launcher.setConst( cdata );
-
- int num = NUM_WGS*WG_SIZE;
- launcher.launch1D( num, WG_SIZE );
+ b3BufferInfoCL bInfo[] = {b3BufferInfoCL(src->getBufferCL(), true), b3BufferInfoCL(srcHisto->getBufferCL())};
+ b3LauncherCL launcher(m_commandQueue, m_streamCountSortDataKernel, "m_streamCountSortDataKernel");
+
+ launcher.setBuffers(bInfo, sizeof(bInfo) / sizeof(b3BufferInfoCL));
+ launcher.setConst(cdata);
+
+ int num = NUM_WGS * WG_SIZE;
+ launcher.launch1D(num, WG_SIZE);
}
-
-
#ifdef DEBUG_RADIXSORT
b3AlignedObjectArray<unsigned int> testHist;
srcHisto->copyToHost(testHist);
- printf("ib = %d, testHist size = %d, non zero elements:\n",ib, testHist.size());
- for (int i=0;i<testHist.size();i++)
+ printf("ib = %d, testHist size = %d, non zero elements:\n", ib, testHist.size());
+ for (int i = 0; i < testHist.size(); i++)
{
- if (testHist[i]!=0)
- printf("testHist[%d]=%d\n",i,testHist[i]);
+ if (testHist[i] != 0)
+ printf("testHist[%d]=%d\n", i, testHist[i]);
}
-#endif //DEBUG_RADIXSORT
-
-
+#endif //DEBUG_RADIXSORT
//fast prefix scan is not working properly on Mac OSX yet
#ifdef __APPLE__
- bool fastScan=false;
+ bool fastScan = false;
#else
- bool fastScan=!m_deviceCPU;//only use fast scan on GPU
+ bool fastScan = !m_deviceCPU; //only use fast scan on GPU
#endif
if (fastScan)
- {// prefix scan group histogram
- b3BufferInfoCL bInfo[] = { b3BufferInfoCL( srcHisto->getBufferCL() ) };
- b3LauncherCL launcher( m_commandQueue, m_prefixScanKernel,"m_prefixScanKernel" );
- launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
- launcher.setConst( cdata );
- launcher.launch1D( 128, 128 );
+ { // prefix scan group histogram
+ b3BufferInfoCL bInfo[] = {b3BufferInfoCL(srcHisto->getBufferCL())};
+ b3LauncherCL launcher(m_commandQueue, m_prefixScanKernel, "m_prefixScanKernel");
+ launcher.setBuffers(bInfo, sizeof(bInfo) / sizeof(b3BufferInfoCL));
+ launcher.setConst(cdata);
+ launcher.launch1D(128, 128);
destHisto = srcHisto;
- }else
+ }
+ else
{
//unsigned int sum; //for debugging
- m_scan->execute(*srcHisto,*destHisto,1920,0);//,&sum);
+ m_scan->execute(*srcHisto, *destHisto, 1920, 0); //,&sum);
}
-
#ifdef DEBUG_RADIXSORT
destHisto->copyToHost(testHist);
- printf("ib = %d, testHist size = %d, non zero elements:\n",ib, testHist.size());
- for (int i=0;i<testHist.size();i++)
+ printf("ib = %d, testHist size = %d, non zero elements:\n", ib, testHist.size());
+ for (int i = 0; i < testHist.size(); i++)
{
- if (testHist[i]!=0)
- printf("testHist[%d]=%d\n",i,testHist[i]);
+ if (testHist[i] != 0)
+ printf("testHist[%d]=%d\n", i, testHist[i]);
}
-
- for (int i=0;i<testHist.size();i+=NUM_WGS)
+
+ for (int i = 0; i < testHist.size(); i += NUM_WGS)
{
- printf("testHist[%d]=%d\n",i/NUM_WGS,testHist[i]);
+ printf("testHist[%d]=%d\n", i / NUM_WGS, testHist[i]);
}
-#endif //DEBUG_RADIXSORT
+#endif //DEBUG_RADIXSORT
#define USE_GPU
#ifdef USE_GPU
-
+
if (src->size())
- {// local sort and distribute
- b3BufferInfoCL bInfo[] = { b3BufferInfoCL( src->getBufferCL(), true ), b3BufferInfoCL( destHisto->getBufferCL(), true ), b3BufferInfoCL( dst->getBufferCL() )};
- b3LauncherCL launcher( m_commandQueue, m_sortAndScatterSortDataKernel,"m_sortAndScatterSortDataKernel" );
- launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
- launcher.setConst( cdata );
- launcher.launch1D( nWGs*WG_SIZE, WG_SIZE );
-
+ { // local sort and distribute
+ b3BufferInfoCL bInfo[] = {b3BufferInfoCL(src->getBufferCL(), true), b3BufferInfoCL(destHisto->getBufferCL(), true), b3BufferInfoCL(dst->getBufferCL())};
+ b3LauncherCL launcher(m_commandQueue, m_sortAndScatterSortDataKernel, "m_sortAndScatterSortDataKernel");
+ launcher.setBuffers(bInfo, sizeof(bInfo) / sizeof(b3BufferInfoCL));
+ launcher.setConst(cdata);
+ launcher.launch1D(nWGs * WG_SIZE, WG_SIZE);
}
#else
- {
+ {
#define NUM_TABLES 16
//#define SEQUENTIAL
#ifdef SEQUENTIAL
- int counter2[NUM_TABLES]={0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0};
- int tables[NUM_TABLES];
- int startBit = ib;
-
- destHisto->copyToHost(testHist);
- b3AlignedObjectArray<b3SortData> srcHost;
- b3AlignedObjectArray<b3SortData> dstHost;
- dstHost.resize(src->size());
-
- src->copyToHost(srcHost);
-
- for (int i=0;i<NUM_TABLES;i++)
- {
- tables[i] = testHist[i*NUM_WGS];
- }
-
- // distribute
- for(int i=0; i<n; i++)
- {
- int tableIdx = (srcHost[i].m_key >> startBit) & (NUM_TABLES-1);
-
- dstHost[tables[tableIdx] + counter2[tableIdx]] = srcHost[i];
- counter2[tableIdx] ++;
- }
-
-
+ int counter2[NUM_TABLES] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+ int tables[NUM_TABLES];
+ int startBit = ib;
+
+ destHisto->copyToHost(testHist);
+ b3AlignedObjectArray<b3SortData> srcHost;
+ b3AlignedObjectArray<b3SortData> dstHost;
+ dstHost.resize(src->size());
+
+ src->copyToHost(srcHost);
+
+ for (int i = 0; i < NUM_TABLES; i++)
+ {
+ tables[i] = testHist[i * NUM_WGS];
+ }
+
+ // distribute
+ for (int i = 0; i < n; i++)
+ {
+ int tableIdx = (srcHost[i].m_key >> startBit) & (NUM_TABLES - 1);
+
+ dstHost[tables[tableIdx] + counter2[tableIdx]] = srcHost[i];
+ counter2[tableIdx]++;
+ }
+
#else
-
- int counter2[NUM_TABLES]={0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0};
-
- int tables[NUM_TABLES];
- b3AlignedObjectArray<b3SortData> dstHostOK;
- dstHostOK.resize(src->size());
-
- destHisto->copyToHost(testHist);
- b3AlignedObjectArray<b3SortData> srcHost;
- src->copyToHost(srcHost);
-
- int blockSize = 256;
- int nBlocksPerWG = cdata.m_nBlocksPerWG;
- int startBit = ib;
-
- {
- for (int i=0;i<NUM_TABLES;i++)
- {
- tables[i] = testHist[i*NUM_WGS];
- }
-
- // distribute
- for(int i=0; i<n; i++)
- {
- int tableIdx = (srcHost[i].m_key >> startBit) & (NUM_TABLES-1);
-
- dstHostOK[tables[tableIdx] + counter2[tableIdx]] = srcHost[i];
- counter2[tableIdx] ++;
- }
-
-
- }
-
-
- b3AlignedObjectArray<b3SortData> dstHost;
- dstHost.resize(src->size());
-
-
- int counter[NUM_TABLES]={0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0};
-
-
-
- for (int wgIdx=0;wgIdx<NUM_WGS;wgIdx++)
- {
- int counter[NUM_TABLES]={0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0};
-
- int nBlocks = (n)/blockSize - nBlocksPerWG*wgIdx;
-
- for(int iblock=0; iblock<b3Min(cdata.m_nBlocksPerWG, nBlocks); iblock++)
- {
- for (int lIdx = 0;lIdx < 64;lIdx++)
- {
- int addr = iblock*blockSize + blockSize*cdata.m_nBlocksPerWG*wgIdx + ELEMENTS_PER_WORK_ITEM*lIdx;
-
- // MY_HISTOGRAM( localKeys.x ) ++ is much expensive than atomic add as it requires read and write while atomics can just add on AMD
- // Using registers didn't perform well. It seems like use localKeys to address requires a lot of alu ops
- // AMD: AtomInc performs better while NV prefers ++
- for(int j=0; j<ELEMENTS_PER_WORK_ITEM; j++)
- {
- if( addr+j < n )
- {
- // printf ("addr+j=%d\n", addr+j);
-
- int i = addr+j;
-
- int tableIdx = (srcHost[i].m_key >> startBit) & (NUM_TABLES-1);
-
- int destIndex = testHist[tableIdx*NUM_WGS+wgIdx] + counter[tableIdx];
-
- b3SortData ok = dstHostOK[destIndex];
-
- if (ok.m_key != srcHost[i].m_key)
- {
- printf("ok.m_key = %d, srcHost[i].m_key = %d\n", ok.m_key,srcHost[i].m_key );
- printf("(ok.m_value = %d, srcHost[i].m_value = %d)\n", ok.m_value,srcHost[i].m_value );
- }
- if (ok.m_value != srcHost[i].m_value)
- {
-
- printf("ok.m_value = %d, srcHost[i].m_value = %d\n", ok.m_value,srcHost[i].m_value );
- printf("(ok.m_key = %d, srcHost[i].m_key = %d)\n", ok.m_key,srcHost[i].m_key );
-
- }
-
- dstHost[destIndex] = srcHost[i];
- counter[tableIdx] ++;
-
- }
- }
- }
- }
- }
-
-
-#endif //SEQUENTIAL
-
- dst->copyFromHost(dstHost);
- }
-#endif//USE_GPU
-
-
-
+
+ int counter2[NUM_TABLES] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+
+ int tables[NUM_TABLES];
+ b3AlignedObjectArray<b3SortData> dstHostOK;
+ dstHostOK.resize(src->size());
+
+ destHisto->copyToHost(testHist);
+ b3AlignedObjectArray<b3SortData> srcHost;
+ src->copyToHost(srcHost);
+
+ int blockSize = 256;
+ int nBlocksPerWG = cdata.m_nBlocksPerWG;
+ int startBit = ib;
+
+ {
+ for (int i = 0; i < NUM_TABLES; i++)
+ {
+ tables[i] = testHist[i * NUM_WGS];
+ }
+
+ // distribute
+ for (int i = 0; i < n; i++)
+ {
+ int tableIdx = (srcHost[i].m_key >> startBit) & (NUM_TABLES - 1);
+
+ dstHostOK[tables[tableIdx] + counter2[tableIdx]] = srcHost[i];
+ counter2[tableIdx]++;
+ }
+ }
+
+ b3AlignedObjectArray<b3SortData> dstHost;
+ dstHost.resize(src->size());
+
+ int counter[NUM_TABLES] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+
+ for (int wgIdx = 0; wgIdx < NUM_WGS; wgIdx++)
+ {
+ int counter[NUM_TABLES] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+
+ int nBlocks = (n) / blockSize - nBlocksPerWG * wgIdx;
+
+ for (int iblock = 0; iblock < b3Min(cdata.m_nBlocksPerWG, nBlocks); iblock++)
+ {
+ for (int lIdx = 0; lIdx < 64; lIdx++)
+ {
+ int addr = iblock * blockSize + blockSize * cdata.m_nBlocksPerWG * wgIdx + ELEMENTS_PER_WORK_ITEM * lIdx;
+
+ // MY_HISTOGRAM( localKeys.x ) ++ is much expensive than atomic add as it requires read and write while atomics can just add on AMD
+ // Using registers didn't perform well. It seems like use localKeys to address requires a lot of alu ops
+ // AMD: AtomInc performs better while NV prefers ++
+ for (int j = 0; j < ELEMENTS_PER_WORK_ITEM; j++)
+ {
+ if (addr + j < n)
+ {
+ // printf ("addr+j=%d\n", addr+j);
+
+ int i = addr + j;
+
+ int tableIdx = (srcHost[i].m_key >> startBit) & (NUM_TABLES - 1);
+
+ int destIndex = testHist[tableIdx * NUM_WGS + wgIdx] + counter[tableIdx];
+
+ b3SortData ok = dstHostOK[destIndex];
+
+ if (ok.m_key != srcHost[i].m_key)
+ {
+ printf("ok.m_key = %d, srcHost[i].m_key = %d\n", ok.m_key, srcHost[i].m_key);
+ printf("(ok.m_value = %d, srcHost[i].m_value = %d)\n", ok.m_value, srcHost[i].m_value);
+ }
+ if (ok.m_value != srcHost[i].m_value)
+ {
+ printf("ok.m_value = %d, srcHost[i].m_value = %d\n", ok.m_value, srcHost[i].m_value);
+ printf("(ok.m_key = %d, srcHost[i].m_key = %d)\n", ok.m_key, srcHost[i].m_key);
+ }
+
+ dstHost[destIndex] = srcHost[i];
+ counter[tableIdx]++;
+ }
+ }
+ }
+ }
+ }
+
+#endif //SEQUENTIAL
+
+ dst->copyFromHost(dstHost);
+ }
+#endif //USE_GPU
+
#ifdef DEBUG_RADIXSORT
destHisto->copyToHost(testHist);
- printf("ib = %d, testHist size = %d, non zero elements:\n",ib, testHist.size());
- for (int i=0;i<testHist.size();i++)
+ printf("ib = %d, testHist size = %d, non zero elements:\n", ib, testHist.size());
+ for (int i = 0; i < testHist.size(); i++)
{
- if (testHist[i]!=0)
- printf("testHist[%d]=%d\n",i,testHist[i]);
+ if (testHist[i] != 0)
+ printf("testHist[%d]=%d\n", i, testHist[i]);
}
-#endif //DEBUG_RADIXSORT
- b3Swap(src, dst );
- b3Swap(srcHisto,destHisto);
+#endif //DEBUG_RADIXSORT
+ b3Swap(src, dst);
+ b3Swap(srcHisto, destHisto);
#ifdef DEBUG_RADIXSORT2
- keyValuesInOut.copyToHost(test2);
- printf("numElem = %d\n",test2.size());
- for (int i=0;i<test2.size();i++)
- {
- if (test2[i].m_key != test2[i].m_value)
- {
- printf("test2[%d].m_key=%d\n",i,test2[i].m_key);
- printf("test2[%d].m_value=%d\n",i,test2[i].m_value);
- }
- }
-#endif //DEBUG_RADIXSORT2
-
- count++;
-
-
+ keyValuesInOut.copyToHost(test2);
+ printf("numElem = %d\n", test2.size());
+ for (int i = 0; i < test2.size(); i++)
+ {
+ if (test2[i].m_key != test2[i].m_value)
+ {
+ printf("test2[%d].m_key=%d\n", i, test2[i].m_key);
+ printf("test2[%d].m_value=%d\n", i, test2[i].m_value);
+ }
+ }
+#endif //DEBUG_RADIXSORT2
+
+ count++;
}
-
-
-
- if (count&1)
+
+ if (count & 1)
{
- b3Assert(0);//need to copy from workbuffer to keyValuesInOut
+ b3Assert(0); //need to copy from workbuffer to keyValuesInOut
}
if (m_workBuffer4->size())
@@ -545,153 +505,137 @@ void b3RadixSort32CL::execute(b3OpenCLArray<b3SortData>& keyValuesInOut, int sor
keyValuesInOut.copyFromOpenCLArray(*m_workBuffer4);
}
-
#ifdef DEBUG_RADIXSORT
- keyValuesInOut.copyToHost(test2);
-
- printf("numElem = %d\n",test2.size());
- for (int i=0;i<test2.size();i++)
- {
- printf("test2[%d].m_key=%d\n",i,test2[i].m_key);
- printf("test2[%d].m_value=%d\n",i,test2[i].m_value);
- }
-#endif
-
-}
-
-
-
-
+ keyValuesInOut.copyToHost(test2);
+ printf("numElem = %d\n", test2.size());
+ for (int i = 0; i < test2.size(); i++)
+ {
+ printf("test2[%d].m_key=%d\n", i, test2[i].m_key);
+ printf("test2[%d].m_value=%d\n", i, test2[i].m_value);
+ }
+#endif
+}
void b3RadixSort32CL::execute(b3OpenCLArray<unsigned int>& keysInOut, int sortBits /* = 32 */)
{
int originalSize = keysInOut.size();
int workingSize = originalSize;
-
-
+
int dataAlignment = DATA_ALIGNMENT;
b3OpenCLArray<unsigned int>* src = 0;
- if (workingSize%dataAlignment)
+ if (workingSize % dataAlignment)
{
- workingSize += dataAlignment-(workingSize%dataAlignment);
+ workingSize += dataAlignment - (workingSize % dataAlignment);
m_workBuffer4a->copyFromOpenCLArray(keysInOut);
m_workBuffer4a->resize(workingSize);
unsigned int fillValue = 0xffffffff;
-
- m_fill->execute(*m_workBuffer4a,fillValue,workingSize-originalSize,originalSize);
+
+ m_fill->execute(*m_workBuffer4a, fillValue, workingSize - originalSize, originalSize);
src = m_workBuffer4a;
- } else
+ }
+ else
{
src = &keysInOut;
m_workBuffer4a->resize(0);
}
-
-
-
- b3Assert( workingSize%DATA_ALIGNMENT == 0 );
- int minCap = NUM_BUCKET*NUM_WGS;
+ b3Assert(workingSize % DATA_ALIGNMENT == 0);
+ int minCap = NUM_BUCKET * NUM_WGS;
int n = workingSize;
-
m_workBuffer1->resize(minCap);
m_workBuffer3->resize(workingSize);
m_workBuffer3a->resize(workingSize);
-// ADLASSERT( ELEMENTS_PER_WORK_ITEM == 4 );
- b3Assert( BITS_PER_PASS == 4 );
- b3Assert( WG_SIZE == 64 );
- b3Assert( (sortBits&0x3) == 0 );
+ // ADLASSERT( ELEMENTS_PER_WORK_ITEM == 4 );
+ b3Assert(BITS_PER_PASS == 4);
+ b3Assert(WG_SIZE == 64);
+ b3Assert((sortBits & 0x3) == 0);
-
-
b3OpenCLArray<unsigned int>* dst = m_workBuffer3a;
b3OpenCLArray<unsigned int>* srcHisto = m_workBuffer1;
b3OpenCLArray<unsigned int>* destHisto = m_workBuffer2;
-
int nWGs = NUM_WGS;
b3ConstData cdata;
{
- int blockSize = ELEMENTS_PER_WORK_ITEM*WG_SIZE;//set at 256
- int nBlocks = (n+blockSize-1)/(blockSize);
+ int blockSize = ELEMENTS_PER_WORK_ITEM * WG_SIZE; //set at 256
+ int nBlocks = (n + blockSize - 1) / (blockSize);
cdata.m_n = n;
cdata.m_nWGs = NUM_WGS;
cdata.m_startBit = 0;
- cdata.m_nBlocksPerWG = (nBlocks + cdata.m_nWGs - 1)/cdata.m_nWGs;
- if( nBlocks < NUM_WGS )
+ cdata.m_nBlocksPerWG = (nBlocks + cdata.m_nWGs - 1) / cdata.m_nWGs;
+ if (nBlocks < NUM_WGS)
{
cdata.m_nBlocksPerWG = 1;
nWGs = nBlocks;
}
}
- int count=0;
- for(int ib=0; ib<sortBits; ib+=4)
+ int count = 0;
+ for (int ib = 0; ib < sortBits; ib += 4)
{
cdata.m_startBit = ib;
-
+
if (src->size())
{
- b3BufferInfoCL bInfo[] = { b3BufferInfoCL( src->getBufferCL(), true ), b3BufferInfoCL( srcHisto->getBufferCL() ) };
- b3LauncherCL launcher(m_commandQueue, m_streamCountKernel,"m_streamCountKernel");
-
- launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
- launcher.setConst( cdata );
-
- int num = NUM_WGS*WG_SIZE;
- launcher.launch1D( num, WG_SIZE );
- }
+ b3BufferInfoCL bInfo[] = {b3BufferInfoCL(src->getBufferCL(), true), b3BufferInfoCL(srcHisto->getBufferCL())};
+ b3LauncherCL launcher(m_commandQueue, m_streamCountKernel, "m_streamCountKernel");
+
+ launcher.setBuffers(bInfo, sizeof(bInfo) / sizeof(b3BufferInfoCL));
+ launcher.setConst(cdata);
-
+ int num = NUM_WGS * WG_SIZE;
+ launcher.launch1D(num, WG_SIZE);
+ }
//fast prefix scan is not working properly on Mac OSX yet
#ifdef __APPLE__
- bool fastScan=false;
+ bool fastScan = false;
#else
- bool fastScan=!m_deviceCPU;
+ bool fastScan = !m_deviceCPU;
#endif
if (fastScan)
- {// prefix scan group histogram
- b3BufferInfoCL bInfo[] = { b3BufferInfoCL( srcHisto->getBufferCL() ) };
- b3LauncherCL launcher( m_commandQueue, m_prefixScanKernel,"m_prefixScanKernel" );
- launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
- launcher.setConst( cdata );
- launcher.launch1D( 128, 128 );
+ { // prefix scan group histogram
+ b3BufferInfoCL bInfo[] = {b3BufferInfoCL(srcHisto->getBufferCL())};
+ b3LauncherCL launcher(m_commandQueue, m_prefixScanKernel, "m_prefixScanKernel");
+ launcher.setBuffers(bInfo, sizeof(bInfo) / sizeof(b3BufferInfoCL));
+ launcher.setConst(cdata);
+ launcher.launch1D(128, 128);
destHisto = srcHisto;
- }else
+ }
+ else
{
//unsigned int sum; //for debugging
- m_scan->execute(*srcHisto,*destHisto,1920,0);//,&sum);
+ m_scan->execute(*srcHisto, *destHisto, 1920, 0); //,&sum);
}
if (src->size())
- {// local sort and distribute
- b3BufferInfoCL bInfo[] = { b3BufferInfoCL( src->getBufferCL(), true ), b3BufferInfoCL( destHisto->getBufferCL(), true ), b3BufferInfoCL( dst->getBufferCL() )};
- b3LauncherCL launcher( m_commandQueue, m_sortAndScatterKernel ,"m_sortAndScatterKernel");
- launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
- launcher.setConst( cdata );
- launcher.launch1D( nWGs*WG_SIZE, WG_SIZE );
-
+ { // local sort and distribute
+ b3BufferInfoCL bInfo[] = {b3BufferInfoCL(src->getBufferCL(), true), b3BufferInfoCL(destHisto->getBufferCL(), true), b3BufferInfoCL(dst->getBufferCL())};
+ b3LauncherCL launcher(m_commandQueue, m_sortAndScatterKernel, "m_sortAndScatterKernel");
+ launcher.setBuffers(bInfo, sizeof(bInfo) / sizeof(b3BufferInfoCL));
+ launcher.setConst(cdata);
+ launcher.launch1D(nWGs * WG_SIZE, WG_SIZE);
}
-
- b3Swap(src, dst );
- b3Swap(srcHisto,destHisto);
- count++;
+ b3Swap(src, dst);
+ b3Swap(srcHisto, destHisto);
+
+ count++;
}
-
- if (count&1)
+
+ if (count & 1)
{
- b3Assert(0);//need to copy from workbuffer to keyValuesInOut
+ b3Assert(0); //need to copy from workbuffer to keyValuesInOut
}
if (m_workBuffer4a->size())
@@ -699,12 +643,4 @@ void b3RadixSort32CL::execute(b3OpenCLArray<unsigned int>& keysInOut, int sortBi
m_workBuffer4a->resize(originalSize);
keysInOut.copyFromOpenCLArray(*m_workBuffer4a);
}
-
}
-
-
-
-
-
-
-
diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3RadixSort32CL.h b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3RadixSort32CL.h
index 975bd80e53..69caf182d7 100644
--- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3RadixSort32CL.h
+++ b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3RadixSort32CL.h
@@ -6,90 +6,79 @@
struct b3SortData
{
- union
- {
+ union {
unsigned int m_key;
unsigned int x;
};
- union
- {
+ union {
unsigned int m_value;
unsigned int y;
-
};
};
#include "b3BufferInfoCL.h"
-class b3RadixSort32CL
+class b3RadixSort32CL
{
+ b3OpenCLArray<unsigned int>* m_workBuffer1;
+ b3OpenCLArray<unsigned int>* m_workBuffer2;
- b3OpenCLArray<unsigned int>* m_workBuffer1;
- b3OpenCLArray<unsigned int>* m_workBuffer2;
-
- b3OpenCLArray<b3SortData>* m_workBuffer3;
- b3OpenCLArray<b3SortData>* m_workBuffer4;
-
- b3OpenCLArray<unsigned int>* m_workBuffer3a;
- b3OpenCLArray<unsigned int>* m_workBuffer4a;
+ b3OpenCLArray<b3SortData>* m_workBuffer3;
+ b3OpenCLArray<b3SortData>* m_workBuffer4;
- cl_command_queue m_commandQueue;
+ b3OpenCLArray<unsigned int>* m_workBuffer3a;
+ b3OpenCLArray<unsigned int>* m_workBuffer4a;
- cl_kernel m_streamCountSortDataKernel;
- cl_kernel m_streamCountKernel;
+ cl_command_queue m_commandQueue;
- cl_kernel m_prefixScanKernel;
- cl_kernel m_sortAndScatterSortDataKernel;
- cl_kernel m_sortAndScatterKernel;
+ cl_kernel m_streamCountSortDataKernel;
+ cl_kernel m_streamCountKernel;
+ cl_kernel m_prefixScanKernel;
+ cl_kernel m_sortAndScatterSortDataKernel;
+ cl_kernel m_sortAndScatterKernel;
- bool m_deviceCPU;
+ bool m_deviceCPU;
- class b3PrefixScanCL* m_scan;
- class b3FillCL* m_fill;
+ class b3PrefixScanCL* m_scan;
+ class b3FillCL* m_fill;
public:
struct b3ConstData
- {
- int m_n;
- int m_nWGs;
- int m_startBit;
- int m_nBlocksPerWG;
- };
+ {
+ int m_n;
+ int m_nWGs;
+ int m_startBit;
+ int m_nBlocksPerWG;
+ };
enum
- {
- DATA_ALIGNMENT = 256,
- WG_SIZE = 64,
- BLOCK_SIZE = 256,
- ELEMENTS_PER_WORK_ITEM = (BLOCK_SIZE/WG_SIZE),
- BITS_PER_PASS = 4,
- NUM_BUCKET=(1<<BITS_PER_PASS),
- // if you change this, change nPerWI in kernel as well
- NUM_WGS = 20*6, // cypress
-// NUM_WGS = 24*6, // cayman
-// NUM_WGS = 32*4, // nv
- };
-
+ {
+ DATA_ALIGNMENT = 256,
+ WG_SIZE = 64,
+ BLOCK_SIZE = 256,
+ ELEMENTS_PER_WORK_ITEM = (BLOCK_SIZE / WG_SIZE),
+ BITS_PER_PASS = 4,
+ NUM_BUCKET = (1 << BITS_PER_PASS),
+ // if you change this, change nPerWI in kernel as well
+ NUM_WGS = 20 * 6, // cypress
+ // NUM_WGS = 24*6, // cayman
+ // NUM_WGS = 32*4, // nv
+ };
private:
-
-
public:
+ b3RadixSort32CL(cl_context ctx, cl_device_id device, cl_command_queue queue, int initialCapacity = 0);
- b3RadixSort32CL(cl_context ctx, cl_device_id device, cl_command_queue queue, int initialCapacity =0);
+ virtual ~b3RadixSort32CL();
- virtual ~b3RadixSort32CL();
+ void execute(b3OpenCLArray<unsigned int>& keysIn, b3OpenCLArray<unsigned int>& keysOut, b3OpenCLArray<unsigned int>& valuesIn,
+ b3OpenCLArray<unsigned int>& valuesOut, int n, int sortBits = 32);
- void execute(b3OpenCLArray<unsigned int>& keysIn, b3OpenCLArray<unsigned int>& keysOut, b3OpenCLArray<unsigned int>& valuesIn,
- b3OpenCLArray<unsigned int>& valuesOut, int n, int sortBits = 32);
-
- ///keys only
- void execute(b3OpenCLArray<unsigned int>& keysInOut, int sortBits = 32 );
-
- void execute(b3OpenCLArray<b3SortData>& keyValuesInOut, int sortBits = 32 );
- void executeHost(b3OpenCLArray<b3SortData>& keyValuesInOut, int sortBits = 32);
- void executeHost(b3AlignedObjectArray<b3SortData>& keyValuesInOut, int sortBits = 32);
+ ///keys only
+ void execute(b3OpenCLArray<unsigned int>& keysInOut, int sortBits = 32);
+ void execute(b3OpenCLArray<b3SortData>& keyValuesInOut, int sortBits = 32);
+ void executeHost(b3OpenCLArray<b3SortData>& keyValuesInOut, int sortBits = 32);
+ void executeHost(b3AlignedObjectArray<b3SortData>& keyValuesInOut, int sortBits = 32);
};
-#endif //B3_RADIXSORT32_H
-
+#endif //B3_RADIXSORT32_H
diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/BoundSearchKernelsCL.h b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/BoundSearchKernelsCL.h
index 9c9e847138..1758dd41e3 100644
--- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/BoundSearchKernelsCL.h
+++ b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/BoundSearchKernelsCL.h
@@ -1,87 +1,86 @@
//this file is autogenerated using stringify.bat (premake --stringify) in the build folder of this project
-static const char* boundSearchKernelsCL= \
-"/*\n"
-"Copyright (c) 2012 Advanced Micro Devices, Inc. \n"
-"This software is provided 'as-is', without any express or implied warranty.\n"
-"In no event will the authors be held liable for any damages arising from the use of this software.\n"
-"Permission is granted to anyone to use this software for any purpose, \n"
-"including commercial applications, and to alter it and redistribute it freely, \n"
-"subject to the following restrictions:\n"
-"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.\n"
-"2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software.\n"
-"3. This notice may not be removed or altered from any source distribution.\n"
-"*/\n"
-"//Originally written by Takahiro Harada\n"
-"typedef unsigned int u32;\n"
-"#define GET_GROUP_IDX get_group_id(0)\n"
-"#define GET_LOCAL_IDX get_local_id(0)\n"
-"#define GET_GLOBAL_IDX get_global_id(0)\n"
-"#define GET_GROUP_SIZE get_local_size(0)\n"
-"#define GROUP_LDS_BARRIER barrier(CLK_LOCAL_MEM_FENCE)\n"
-"typedef struct\n"
-"{\n"
-" u32 m_key; \n"
-" u32 m_value;\n"
-"}SortData;\n"
-"typedef struct\n"
-"{\n"
-" u32 m_nSrc;\n"
-" u32 m_nDst;\n"
-" u32 m_padding[2];\n"
-"} ConstBuffer;\n"
-"__attribute__((reqd_work_group_size(64,1,1)))\n"
-"__kernel\n"
-"void SearchSortDataLowerKernel(__global SortData* src, __global u32 *dst, \n"
-" unsigned int nSrc, unsigned int nDst)\n"
-"{\n"
-" int gIdx = GET_GLOBAL_IDX;\n"
-" if( gIdx < nSrc )\n"
-" {\n"
-" SortData first; first.m_key = (u32)(-1); first.m_value = (u32)(-1);\n"
-" SortData end; end.m_key = nDst; end.m_value = nDst;\n"
-" SortData iData = (gIdx==0)? first: src[gIdx-1];\n"
-" SortData jData = (gIdx==nSrc)? end: src[gIdx];\n"
-" if( iData.m_key != jData.m_key )\n"
-" {\n"
-"// for(u32 k=iData.m_key+1; k<=min(jData.m_key, nDst-1); k++)\n"
-" u32 k = jData.m_key;\n"
-" {\n"
-" dst[k] = gIdx;\n"
-" }\n"
-" }\n"
-" }\n"
-"}\n"
-"__attribute__((reqd_work_group_size(64,1,1)))\n"
-"__kernel\n"
-"void SearchSortDataUpperKernel(__global SortData* src, __global u32 *dst, \n"
-" unsigned int nSrc, unsigned int nDst)\n"
-"{\n"
-" int gIdx = GET_GLOBAL_IDX+1;\n"
-" if( gIdx < nSrc+1 )\n"
-" {\n"
-" SortData first; first.m_key = 0; first.m_value = 0;\n"
-" SortData end; end.m_key = nDst; end.m_value = nDst;\n"
-" SortData iData = src[gIdx-1];\n"
-" SortData jData = (gIdx==nSrc)? end: src[gIdx];\n"
-" if( iData.m_key != jData.m_key )\n"
-" {\n"
-" u32 k = iData.m_key;\n"
-" {\n"
-" dst[k] = gIdx;\n"
-" }\n"
-" }\n"
-" }\n"
-"}\n"
-"__attribute__((reqd_work_group_size(64,1,1)))\n"
-"__kernel\n"
-"void SubtractKernel(__global u32* A, __global u32 *B, __global u32 *C, \n"
-" unsigned int nSrc, unsigned int nDst)\n"
-"{\n"
-" int gIdx = GET_GLOBAL_IDX;\n"
-" \n"
-" if( gIdx < nDst )\n"
-" {\n"
-" C[gIdx] = A[gIdx] - B[gIdx];\n"
-" }\n"
-"}\n"
-;
+static const char* boundSearchKernelsCL =
+ "/*\n"
+ "Copyright (c) 2012 Advanced Micro Devices, Inc. \n"
+ "This software is provided 'as-is', without any express or implied warranty.\n"
+ "In no event will the authors be held liable for any damages arising from the use of this software.\n"
+ "Permission is granted to anyone to use this software for any purpose, \n"
+ "including commercial applications, and to alter it and redistribute it freely, \n"
+ "subject to the following restrictions:\n"
+ "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.\n"
+ "2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software.\n"
+ "3. This notice may not be removed or altered from any source distribution.\n"
+ "*/\n"
+ "//Originally written by Takahiro Harada\n"
+ "typedef unsigned int u32;\n"
+ "#define GET_GROUP_IDX get_group_id(0)\n"
+ "#define GET_LOCAL_IDX get_local_id(0)\n"
+ "#define GET_GLOBAL_IDX get_global_id(0)\n"
+ "#define GET_GROUP_SIZE get_local_size(0)\n"
+ "#define GROUP_LDS_BARRIER barrier(CLK_LOCAL_MEM_FENCE)\n"
+ "typedef struct\n"
+ "{\n"
+ " u32 m_key; \n"
+ " u32 m_value;\n"
+ "}SortData;\n"
+ "typedef struct\n"
+ "{\n"
+ " u32 m_nSrc;\n"
+ " u32 m_nDst;\n"
+ " u32 m_padding[2];\n"
+ "} ConstBuffer;\n"
+ "__attribute__((reqd_work_group_size(64,1,1)))\n"
+ "__kernel\n"
+ "void SearchSortDataLowerKernel(__global SortData* src, __global u32 *dst, \n"
+ " unsigned int nSrc, unsigned int nDst)\n"
+ "{\n"
+ " int gIdx = GET_GLOBAL_IDX;\n"
+ " if( gIdx < nSrc )\n"
+ " {\n"
+ " SortData first; first.m_key = (u32)(-1); first.m_value = (u32)(-1);\n"
+ " SortData end; end.m_key = nDst; end.m_value = nDst;\n"
+ " SortData iData = (gIdx==0)? first: src[gIdx-1];\n"
+ " SortData jData = (gIdx==nSrc)? end: src[gIdx];\n"
+ " if( iData.m_key != jData.m_key )\n"
+ " {\n"
+ "// for(u32 k=iData.m_key+1; k<=min(jData.m_key, nDst-1); k++)\n"
+ " u32 k = jData.m_key;\n"
+ " {\n"
+ " dst[k] = gIdx;\n"
+ " }\n"
+ " }\n"
+ " }\n"
+ "}\n"
+ "__attribute__((reqd_work_group_size(64,1,1)))\n"
+ "__kernel\n"
+ "void SearchSortDataUpperKernel(__global SortData* src, __global u32 *dst, \n"
+ " unsigned int nSrc, unsigned int nDst)\n"
+ "{\n"
+ " int gIdx = GET_GLOBAL_IDX+1;\n"
+ " if( gIdx < nSrc+1 )\n"
+ " {\n"
+ " SortData first; first.m_key = 0; first.m_value = 0;\n"
+ " SortData end; end.m_key = nDst; end.m_value = nDst;\n"
+ " SortData iData = src[gIdx-1];\n"
+ " SortData jData = (gIdx==nSrc)? end: src[gIdx];\n"
+ " if( iData.m_key != jData.m_key )\n"
+ " {\n"
+ " u32 k = iData.m_key;\n"
+ " {\n"
+ " dst[k] = gIdx;\n"
+ " }\n"
+ " }\n"
+ " }\n"
+ "}\n"
+ "__attribute__((reqd_work_group_size(64,1,1)))\n"
+ "__kernel\n"
+ "void SubtractKernel(__global u32* A, __global u32 *B, __global u32 *C, \n"
+ " unsigned int nSrc, unsigned int nDst)\n"
+ "{\n"
+ " int gIdx = GET_GLOBAL_IDX;\n"
+ " \n"
+ " if( gIdx < nDst )\n"
+ " {\n"
+ " C[gIdx] = A[gIdx] - B[gIdx];\n"
+ " }\n"
+ "}\n";
diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/CopyKernelsCL.h b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/CopyKernelsCL.h
index e5670e3cd3..33c9279462 100644
--- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/CopyKernelsCL.h
+++ b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/CopyKernelsCL.h
@@ -1,132 +1,131 @@
//this file is autogenerated using stringify.bat (premake --stringify) in the build folder of this project
-static const char* copyKernelsCL= \
-"/*\n"
-"Copyright (c) 2012 Advanced Micro Devices, Inc. \n"
-"\n"
-"This software is provided 'as-is', without any express or implied warranty.\n"
-"In no event will the authors be held liable for any damages arising from the use of this software.\n"
-"Permission is granted to anyone to use this software for any purpose, \n"
-"including commercial applications, and to alter it and redistribute it freely, \n"
-"subject to the following restrictions:\n"
-"\n"
-"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.\n"
-"2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software.\n"
-"3. This notice may not be removed or altered from any source distribution.\n"
-"*/\n"
-"//Originally written by Takahiro Harada\n"
-"\n"
-"#pragma OPENCL EXTENSION cl_amd_printf : enable\n"
-"#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable\n"
-"\n"
-"typedef unsigned int u32;\n"
-"#define GET_GROUP_IDX get_group_id(0)\n"
-"#define GET_LOCAL_IDX get_local_id(0)\n"
-"#define GET_GLOBAL_IDX get_global_id(0)\n"
-"#define GET_GROUP_SIZE get_local_size(0)\n"
-"#define GROUP_LDS_BARRIER barrier(CLK_LOCAL_MEM_FENCE)\n"
-"#define GROUP_MEM_FENCE mem_fence(CLK_LOCAL_MEM_FENCE)\n"
-"#define AtomInc(x) atom_inc(&(x))\n"
-"#define AtomInc1(x, out) out = atom_inc(&(x))\n"
-"\n"
-"#define make_uint4 (uint4)\n"
-"#define make_uint2 (uint2)\n"
-"#define make_int2 (int2)\n"
-"\n"
-"typedef struct\n"
-"{\n"
-" int m_n;\n"
-" int m_padding[3];\n"
-"} ConstBuffer;\n"
-"\n"
-"\n"
-"\n"
-"__kernel\n"
-"__attribute__((reqd_work_group_size(64,1,1)))\n"
-"void Copy1F4Kernel(__global float4* dst, __global float4* src, \n"
-" ConstBuffer cb)\n"
-"{\n"
-" int gIdx = GET_GLOBAL_IDX;\n"
-"\n"
-" if( gIdx < cb.m_n )\n"
-" {\n"
-" float4 a0 = src[gIdx];\n"
-"\n"
-" dst[ gIdx ] = a0;\n"
-" }\n"
-"}\n"
-"\n"
-"__kernel\n"
-"__attribute__((reqd_work_group_size(64,1,1)))\n"
-"void Copy2F4Kernel(__global float4* dst, __global float4* src, \n"
-" ConstBuffer cb)\n"
-"{\n"
-" int gIdx = GET_GLOBAL_IDX;\n"
-"\n"
-" if( 2*gIdx <= cb.m_n )\n"
-" {\n"
-" float4 a0 = src[gIdx*2+0];\n"
-" float4 a1 = src[gIdx*2+1];\n"
-"\n"
-" dst[ gIdx*2+0 ] = a0;\n"
-" dst[ gIdx*2+1 ] = a1;\n"
-" }\n"
-"}\n"
-"\n"
-"__kernel\n"
-"__attribute__((reqd_work_group_size(64,1,1)))\n"
-"void Copy4F4Kernel(__global float4* dst, __global float4* src, \n"
-" ConstBuffer cb)\n"
-"{\n"
-" int gIdx = GET_GLOBAL_IDX;\n"
-"\n"
-" if( 4*gIdx <= cb.m_n )\n"
-" {\n"
-" int idx0 = gIdx*4+0;\n"
-" int idx1 = gIdx*4+1;\n"
-" int idx2 = gIdx*4+2;\n"
-" int idx3 = gIdx*4+3;\n"
-"\n"
-" float4 a0 = src[idx0];\n"
-" float4 a1 = src[idx1];\n"
-" float4 a2 = src[idx2];\n"
-" float4 a3 = src[idx3];\n"
-"\n"
-" dst[ idx0 ] = a0;\n"
-" dst[ idx1 ] = a1;\n"
-" dst[ idx2 ] = a2;\n"
-" dst[ idx3 ] = a3;\n"
-" }\n"
-"}\n"
-"\n"
-"__kernel\n"
-"__attribute__((reqd_work_group_size(64,1,1)))\n"
-"void CopyF1Kernel(__global float* dstF1, __global float* srcF1, \n"
-" ConstBuffer cb)\n"
-"{\n"
-" int gIdx = GET_GLOBAL_IDX;\n"
-"\n"
-" if( gIdx < cb.m_n )\n"
-" {\n"
-" float a0 = srcF1[gIdx];\n"
-"\n"
-" dstF1[ gIdx ] = a0;\n"
-" }\n"
-"}\n"
-"\n"
-"__kernel\n"
-"__attribute__((reqd_work_group_size(64,1,1)))\n"
-"void CopyF2Kernel(__global float2* dstF2, __global float2* srcF2, \n"
-" ConstBuffer cb)\n"
-"{\n"
-" int gIdx = GET_GLOBAL_IDX;\n"
-"\n"
-" if( gIdx < cb.m_n )\n"
-" {\n"
-" float2 a0 = srcF2[gIdx];\n"
-"\n"
-" dstF2[ gIdx ] = a0;\n"
-" }\n"
-"}\n"
-"\n"
-"\n"
-;
+static const char* copyKernelsCL =
+ "/*\n"
+ "Copyright (c) 2012 Advanced Micro Devices, Inc. \n"
+ "\n"
+ "This software is provided 'as-is', without any express or implied warranty.\n"
+ "In no event will the authors be held liable for any damages arising from the use of this software.\n"
+ "Permission is granted to anyone to use this software for any purpose, \n"
+ "including commercial applications, and to alter it and redistribute it freely, \n"
+ "subject to the following restrictions:\n"
+ "\n"
+ "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.\n"
+ "2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software.\n"
+ "3. This notice may not be removed or altered from any source distribution.\n"
+ "*/\n"
+ "//Originally written by Takahiro Harada\n"
+ "\n"
+ "#pragma OPENCL EXTENSION cl_amd_printf : enable\n"
+ "#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable\n"
+ "\n"
+ "typedef unsigned int u32;\n"
+ "#define GET_GROUP_IDX get_group_id(0)\n"
+ "#define GET_LOCAL_IDX get_local_id(0)\n"
+ "#define GET_GLOBAL_IDX get_global_id(0)\n"
+ "#define GET_GROUP_SIZE get_local_size(0)\n"
+ "#define GROUP_LDS_BARRIER barrier(CLK_LOCAL_MEM_FENCE)\n"
+ "#define GROUP_MEM_FENCE mem_fence(CLK_LOCAL_MEM_FENCE)\n"
+ "#define AtomInc(x) atom_inc(&(x))\n"
+ "#define AtomInc1(x, out) out = atom_inc(&(x))\n"
+ "\n"
+ "#define make_uint4 (uint4)\n"
+ "#define make_uint2 (uint2)\n"
+ "#define make_int2 (int2)\n"
+ "\n"
+ "typedef struct\n"
+ "{\n"
+ " int m_n;\n"
+ " int m_padding[3];\n"
+ "} ConstBuffer;\n"
+ "\n"
+ "\n"
+ "\n"
+ "__kernel\n"
+ "__attribute__((reqd_work_group_size(64,1,1)))\n"
+ "void Copy1F4Kernel(__global float4* dst, __global float4* src, \n"
+ " ConstBuffer cb)\n"
+ "{\n"
+ " int gIdx = GET_GLOBAL_IDX;\n"
+ "\n"
+ " if( gIdx < cb.m_n )\n"
+ " {\n"
+ " float4 a0 = src[gIdx];\n"
+ "\n"
+ " dst[ gIdx ] = a0;\n"
+ " }\n"
+ "}\n"
+ "\n"
+ "__kernel\n"
+ "__attribute__((reqd_work_group_size(64,1,1)))\n"
+ "void Copy2F4Kernel(__global float4* dst, __global float4* src, \n"
+ " ConstBuffer cb)\n"
+ "{\n"
+ " int gIdx = GET_GLOBAL_IDX;\n"
+ "\n"
+ " if( 2*gIdx <= cb.m_n )\n"
+ " {\n"
+ " float4 a0 = src[gIdx*2+0];\n"
+ " float4 a1 = src[gIdx*2+1];\n"
+ "\n"
+ " dst[ gIdx*2+0 ] = a0;\n"
+ " dst[ gIdx*2+1 ] = a1;\n"
+ " }\n"
+ "}\n"
+ "\n"
+ "__kernel\n"
+ "__attribute__((reqd_work_group_size(64,1,1)))\n"
+ "void Copy4F4Kernel(__global float4* dst, __global float4* src, \n"
+ " ConstBuffer cb)\n"
+ "{\n"
+ " int gIdx = GET_GLOBAL_IDX;\n"
+ "\n"
+ " if( 4*gIdx <= cb.m_n )\n"
+ " {\n"
+ " int idx0 = gIdx*4+0;\n"
+ " int idx1 = gIdx*4+1;\n"
+ " int idx2 = gIdx*4+2;\n"
+ " int idx3 = gIdx*4+3;\n"
+ "\n"
+ " float4 a0 = src[idx0];\n"
+ " float4 a1 = src[idx1];\n"
+ " float4 a2 = src[idx2];\n"
+ " float4 a3 = src[idx3];\n"
+ "\n"
+ " dst[ idx0 ] = a0;\n"
+ " dst[ idx1 ] = a1;\n"
+ " dst[ idx2 ] = a2;\n"
+ " dst[ idx3 ] = a3;\n"
+ " }\n"
+ "}\n"
+ "\n"
+ "__kernel\n"
+ "__attribute__((reqd_work_group_size(64,1,1)))\n"
+ "void CopyF1Kernel(__global float* dstF1, __global float* srcF1, \n"
+ " ConstBuffer cb)\n"
+ "{\n"
+ " int gIdx = GET_GLOBAL_IDX;\n"
+ "\n"
+ " if( gIdx < cb.m_n )\n"
+ " {\n"
+ " float a0 = srcF1[gIdx];\n"
+ "\n"
+ " dstF1[ gIdx ] = a0;\n"
+ " }\n"
+ "}\n"
+ "\n"
+ "__kernel\n"
+ "__attribute__((reqd_work_group_size(64,1,1)))\n"
+ "void CopyF2Kernel(__global float2* dstF2, __global float2* srcF2, \n"
+ " ConstBuffer cb)\n"
+ "{\n"
+ " int gIdx = GET_GLOBAL_IDX;\n"
+ "\n"
+ " if( gIdx < cb.m_n )\n"
+ " {\n"
+ " float2 a0 = srcF2[gIdx];\n"
+ "\n"
+ " dstF2[ gIdx ] = a0;\n"
+ " }\n"
+ "}\n"
+ "\n"
+ "\n";
diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/FillKernelsCL.h b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/FillKernelsCL.h
index 4f8b96e489..983e652270 100644
--- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/FillKernelsCL.h
+++ b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/FillKernelsCL.h
@@ -1,91 +1,90 @@
//this file is autogenerated using stringify.bat (premake --stringify) in the build folder of this project
-static const char* fillKernelsCL= \
-"/*\n"
-"Copyright (c) 2012 Advanced Micro Devices, Inc. \n"
-"This software is provided 'as-is', without any express or implied warranty.\n"
-"In no event will the authors be held liable for any damages arising from the use of this software.\n"
-"Permission is granted to anyone to use this software for any purpose, \n"
-"including commercial applications, and to alter it and redistribute it freely, \n"
-"subject to the following restrictions:\n"
-"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.\n"
-"2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software.\n"
-"3. This notice may not be removed or altered from any source distribution.\n"
-"*/\n"
-"//Originally written by Takahiro Harada\n"
-"#pragma OPENCL EXTENSION cl_amd_printf : enable\n"
-"#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable\n"
-"typedef unsigned int u32;\n"
-"#define GET_GROUP_IDX get_group_id(0)\n"
-"#define GET_LOCAL_IDX get_local_id(0)\n"
-"#define GET_GLOBAL_IDX get_global_id(0)\n"
-"#define GET_GROUP_SIZE get_local_size(0)\n"
-"#define GROUP_LDS_BARRIER barrier(CLK_LOCAL_MEM_FENCE)\n"
-"#define GROUP_MEM_FENCE mem_fence(CLK_LOCAL_MEM_FENCE)\n"
-"#define AtomInc(x) atom_inc(&(x))\n"
-"#define AtomInc1(x, out) out = atom_inc(&(x))\n"
-"#define make_uint4 (uint4)\n"
-"#define make_uint2 (uint2)\n"
-"#define make_int2 (int2)\n"
-"typedef struct\n"
-"{\n"
-" union\n"
-" {\n"
-" int4 m_data;\n"
-" uint4 m_unsignedData;\n"
-" float m_floatData;\n"
-" };\n"
-" int m_offset;\n"
-" int m_n;\n"
-" int m_padding[2];\n"
-"} ConstBuffer;\n"
-"__kernel\n"
-"__attribute__((reqd_work_group_size(64,1,1)))\n"
-"void FillIntKernel(__global int* dstInt, int num_elements, int value, const int offset)\n"
-"{\n"
-" int gIdx = GET_GLOBAL_IDX;\n"
-" if( gIdx < num_elements )\n"
-" {\n"
-" dstInt[ offset+gIdx ] = value;\n"
-" }\n"
-"}\n"
-"__kernel\n"
-"__attribute__((reqd_work_group_size(64,1,1)))\n"
-"void FillFloatKernel(__global float* dstFloat, int num_elements, float value, const int offset)\n"
-"{\n"
-" int gIdx = GET_GLOBAL_IDX;\n"
-" if( gIdx < num_elements )\n"
-" {\n"
-" dstFloat[ offset+gIdx ] = value;\n"
-" }\n"
-"}\n"
-"__kernel\n"
-"__attribute__((reqd_work_group_size(64,1,1)))\n"
-"void FillUnsignedIntKernel(__global unsigned int* dstInt, const int num, const unsigned int value, const int offset)\n"
-"{\n"
-" int gIdx = GET_GLOBAL_IDX;\n"
-" if( gIdx < num )\n"
-" {\n"
-" dstInt[ offset+gIdx ] = value;\n"
-" }\n"
-"}\n"
-"__kernel\n"
-"__attribute__((reqd_work_group_size(64,1,1)))\n"
-"void FillInt2Kernel(__global int2* dstInt2, const int num, const int2 value, const int offset)\n"
-"{\n"
-" int gIdx = GET_GLOBAL_IDX;\n"
-" if( gIdx < num )\n"
-" {\n"
-" dstInt2[ gIdx + offset] = make_int2( value.x, value.y );\n"
-" }\n"
-"}\n"
-"__kernel\n"
-"__attribute__((reqd_work_group_size(64,1,1)))\n"
-"void FillInt4Kernel(__global int4* dstInt4, const int num, const int4 value, const int offset)\n"
-"{\n"
-" int gIdx = GET_GLOBAL_IDX;\n"
-" if( gIdx < num )\n"
-" {\n"
-" dstInt4[ offset+gIdx ] = value;\n"
-" }\n"
-"}\n"
-;
+static const char* fillKernelsCL =
+ "/*\n"
+ "Copyright (c) 2012 Advanced Micro Devices, Inc. \n"
+ "This software is provided 'as-is', without any express or implied warranty.\n"
+ "In no event will the authors be held liable for any damages arising from the use of this software.\n"
+ "Permission is granted to anyone to use this software for any purpose, \n"
+ "including commercial applications, and to alter it and redistribute it freely, \n"
+ "subject to the following restrictions:\n"
+ "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.\n"
+ "2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software.\n"
+ "3. This notice may not be removed or altered from any source distribution.\n"
+ "*/\n"
+ "//Originally written by Takahiro Harada\n"
+ "#pragma OPENCL EXTENSION cl_amd_printf : enable\n"
+ "#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable\n"
+ "typedef unsigned int u32;\n"
+ "#define GET_GROUP_IDX get_group_id(0)\n"
+ "#define GET_LOCAL_IDX get_local_id(0)\n"
+ "#define GET_GLOBAL_IDX get_global_id(0)\n"
+ "#define GET_GROUP_SIZE get_local_size(0)\n"
+ "#define GROUP_LDS_BARRIER barrier(CLK_LOCAL_MEM_FENCE)\n"
+ "#define GROUP_MEM_FENCE mem_fence(CLK_LOCAL_MEM_FENCE)\n"
+ "#define AtomInc(x) atom_inc(&(x))\n"
+ "#define AtomInc1(x, out) out = atom_inc(&(x))\n"
+ "#define make_uint4 (uint4)\n"
+ "#define make_uint2 (uint2)\n"
+ "#define make_int2 (int2)\n"
+ "typedef struct\n"
+ "{\n"
+ " union\n"
+ " {\n"
+ " int4 m_data;\n"
+ " uint4 m_unsignedData;\n"
+ " float m_floatData;\n"
+ " };\n"
+ " int m_offset;\n"
+ " int m_n;\n"
+ " int m_padding[2];\n"
+ "} ConstBuffer;\n"
+ "__kernel\n"
+ "__attribute__((reqd_work_group_size(64,1,1)))\n"
+ "void FillIntKernel(__global int* dstInt, int num_elements, int value, const int offset)\n"
+ "{\n"
+ " int gIdx = GET_GLOBAL_IDX;\n"
+ " if( gIdx < num_elements )\n"
+ " {\n"
+ " dstInt[ offset+gIdx ] = value;\n"
+ " }\n"
+ "}\n"
+ "__kernel\n"
+ "__attribute__((reqd_work_group_size(64,1,1)))\n"
+ "void FillFloatKernel(__global float* dstFloat, int num_elements, float value, const int offset)\n"
+ "{\n"
+ " int gIdx = GET_GLOBAL_IDX;\n"
+ " if( gIdx < num_elements )\n"
+ " {\n"
+ " dstFloat[ offset+gIdx ] = value;\n"
+ " }\n"
+ "}\n"
+ "__kernel\n"
+ "__attribute__((reqd_work_group_size(64,1,1)))\n"
+ "void FillUnsignedIntKernel(__global unsigned int* dstInt, const int num, const unsigned int value, const int offset)\n"
+ "{\n"
+ " int gIdx = GET_GLOBAL_IDX;\n"
+ " if( gIdx < num )\n"
+ " {\n"
+ " dstInt[ offset+gIdx ] = value;\n"
+ " }\n"
+ "}\n"
+ "__kernel\n"
+ "__attribute__((reqd_work_group_size(64,1,1)))\n"
+ "void FillInt2Kernel(__global int2* dstInt2, const int num, const int2 value, const int offset)\n"
+ "{\n"
+ " int gIdx = GET_GLOBAL_IDX;\n"
+ " if( gIdx < num )\n"
+ " {\n"
+ " dstInt2[ gIdx + offset] = make_int2( value.x, value.y );\n"
+ " }\n"
+ "}\n"
+ "__kernel\n"
+ "__attribute__((reqd_work_group_size(64,1,1)))\n"
+ "void FillInt4Kernel(__global int4* dstInt4, const int num, const int4 value, const int offset)\n"
+ "{\n"
+ " int gIdx = GET_GLOBAL_IDX;\n"
+ " if( gIdx < num )\n"
+ " {\n"
+ " dstInt4[ offset+gIdx ] = value;\n"
+ " }\n"
+ "}\n";
diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/PrefixScanKernelsCL.h b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/PrefixScanKernelsCL.h
index 27baab8331..fc5e7b865c 100644
--- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/PrefixScanKernelsCL.h
+++ b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/PrefixScanKernelsCL.h
@@ -1,129 +1,128 @@
//this file is autogenerated using stringify.bat (premake --stringify) in the build folder of this project
-static const char* prefixScanKernelsCL= \
-"/*\n"
-"Copyright (c) 2012 Advanced Micro Devices, Inc. \n"
-"This software is provided 'as-is', without any express or implied warranty.\n"
-"In no event will the authors be held liable for any damages arising from the use of this software.\n"
-"Permission is granted to anyone to use this software for any purpose, \n"
-"including commercial applications, and to alter it and redistribute it freely, \n"
-"subject to the following restrictions:\n"
-"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.\n"
-"2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software.\n"
-"3. This notice may not be removed or altered from any source distribution.\n"
-"*/\n"
-"//Originally written by Takahiro Harada\n"
-"typedef unsigned int u32;\n"
-"#define GET_GROUP_IDX get_group_id(0)\n"
-"#define GET_LOCAL_IDX get_local_id(0)\n"
-"#define GET_GLOBAL_IDX get_global_id(0)\n"
-"#define GET_GROUP_SIZE get_local_size(0)\n"
-"#define GROUP_LDS_BARRIER barrier(CLK_LOCAL_MEM_FENCE)\n"
-"// takahiro end\n"
-"#define WG_SIZE 128 \n"
-"#define m_numElems x\n"
-"#define m_numBlocks y\n"
-"#define m_numScanBlocks z\n"
-"/*typedef struct\n"
-"{\n"
-" uint m_numElems;\n"
-" uint m_numBlocks;\n"
-" uint m_numScanBlocks;\n"
-" uint m_padding[1];\n"
-"} ConstBuffer;\n"
-"*/\n"
-"u32 ScanExclusive(__local u32* data, u32 n, int lIdx, int lSize)\n"
-"{\n"
-" u32 blocksum;\n"
-" int offset = 1;\n"
-" for(int nActive=n>>1; nActive>0; nActive>>=1, offset<<=1)\n"
-" {\n"
-" GROUP_LDS_BARRIER;\n"
-" for(int iIdx=lIdx; iIdx<nActive; iIdx+=lSize)\n"
-" {\n"
-" int ai = offset*(2*iIdx+1)-1;\n"
-" int bi = offset*(2*iIdx+2)-1;\n"
-" data[bi] += data[ai];\n"
-" }\n"
-" }\n"
-" GROUP_LDS_BARRIER;\n"
-" if( lIdx == 0 )\n"
-" {\n"
-" blocksum = data[ n-1 ];\n"
-" data[ n-1 ] = 0;\n"
-" }\n"
-" GROUP_LDS_BARRIER;\n"
-" offset >>= 1;\n"
-" for(int nActive=1; nActive<n; nActive<<=1, offset>>=1 )\n"
-" {\n"
-" GROUP_LDS_BARRIER;\n"
-" for( int iIdx = lIdx; iIdx<nActive; iIdx += lSize )\n"
-" {\n"
-" int ai = offset*(2*iIdx+1)-1;\n"
-" int bi = offset*(2*iIdx+2)-1;\n"
-" u32 temp = data[ai];\n"
-" data[ai] = data[bi];\n"
-" data[bi] += temp;\n"
-" }\n"
-" }\n"
-" GROUP_LDS_BARRIER;\n"
-" return blocksum;\n"
-"}\n"
-"__attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
-"__kernel\n"
-"void LocalScanKernel(__global u32* dst, __global u32 *src, __global u32 *sumBuffer,\n"
-" uint4 cb)\n"
-"{\n"
-" __local u32 ldsData[WG_SIZE*2];\n"
-" int gIdx = GET_GLOBAL_IDX;\n"
-" int lIdx = GET_LOCAL_IDX;\n"
-" ldsData[2*lIdx] = ( 2*gIdx < cb.m_numElems )? src[2*gIdx]: 0;\n"
-" ldsData[2*lIdx + 1] = ( 2*gIdx+1 < cb.m_numElems )? src[2*gIdx + 1]: 0;\n"
-" u32 sum = ScanExclusive(ldsData, WG_SIZE*2, GET_LOCAL_IDX, GET_GROUP_SIZE);\n"
-" if( lIdx == 0 ) sumBuffer[GET_GROUP_IDX] = sum;\n"
-" if( (2*gIdx) < cb.m_numElems )\n"
-" {\n"
-" dst[2*gIdx] = ldsData[2*lIdx];\n"
-" }\n"
-" if( (2*gIdx + 1) < cb.m_numElems )\n"
-" {\n"
-" dst[2*gIdx + 1] = ldsData[2*lIdx + 1];\n"
-" }\n"
-"}\n"
-"__attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
-"__kernel\n"
-"void AddOffsetKernel(__global u32 *dst, __global u32 *blockSum, uint4 cb)\n"
-"{\n"
-" const u32 blockSize = WG_SIZE*2;\n"
-" int myIdx = GET_GROUP_IDX+1;\n"
-" int lIdx = GET_LOCAL_IDX;\n"
-" u32 iBlockSum = blockSum[myIdx];\n"
-" int endValue = min((myIdx+1)*(blockSize), cb.m_numElems);\n"
-" for(int i=myIdx*blockSize+lIdx; i<endValue; i+=GET_GROUP_SIZE)\n"
-" {\n"
-" dst[i] += iBlockSum;\n"
-" }\n"
-"}\n"
-"__attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
-"__kernel\n"
-"void TopLevelScanKernel(__global u32* dst, uint4 cb)\n"
-"{\n"
-" __local u32 ldsData[2048];\n"
-" int gIdx = GET_GLOBAL_IDX;\n"
-" int lIdx = GET_LOCAL_IDX;\n"
-" int lSize = GET_GROUP_SIZE;\n"
-" for(int i=lIdx; i<cb.m_numScanBlocks; i+=lSize )\n"
-" {\n"
-" ldsData[i] = (i<cb.m_numBlocks)? dst[i]:0;\n"
-" }\n"
-" GROUP_LDS_BARRIER;\n"
-" u32 sum = ScanExclusive(ldsData, cb.m_numScanBlocks, GET_LOCAL_IDX, GET_GROUP_SIZE);\n"
-" for(int i=lIdx; i<cb.m_numBlocks; i+=lSize )\n"
-" {\n"
-" dst[i] = ldsData[i];\n"
-" }\n"
-" if( gIdx == 0 )\n"
-" {\n"
-" dst[cb.m_numBlocks] = sum;\n"
-" }\n"
-"}\n"
-;
+static const char* prefixScanKernelsCL =
+ "/*\n"
+ "Copyright (c) 2012 Advanced Micro Devices, Inc. \n"
+ "This software is provided 'as-is', without any express or implied warranty.\n"
+ "In no event will the authors be held liable for any damages arising from the use of this software.\n"
+ "Permission is granted to anyone to use this software for any purpose, \n"
+ "including commercial applications, and to alter it and redistribute it freely, \n"
+ "subject to the following restrictions:\n"
+ "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.\n"
+ "2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software.\n"
+ "3. This notice may not be removed or altered from any source distribution.\n"
+ "*/\n"
+ "//Originally written by Takahiro Harada\n"
+ "typedef unsigned int u32;\n"
+ "#define GET_GROUP_IDX get_group_id(0)\n"
+ "#define GET_LOCAL_IDX get_local_id(0)\n"
+ "#define GET_GLOBAL_IDX get_global_id(0)\n"
+ "#define GET_GROUP_SIZE get_local_size(0)\n"
+ "#define GROUP_LDS_BARRIER barrier(CLK_LOCAL_MEM_FENCE)\n"
+ "// takahiro end\n"
+ "#define WG_SIZE 128 \n"
+ "#define m_numElems x\n"
+ "#define m_numBlocks y\n"
+ "#define m_numScanBlocks z\n"
+ "/*typedef struct\n"
+ "{\n"
+ " uint m_numElems;\n"
+ " uint m_numBlocks;\n"
+ " uint m_numScanBlocks;\n"
+ " uint m_padding[1];\n"
+ "} ConstBuffer;\n"
+ "*/\n"
+ "u32 ScanExclusive(__local u32* data, u32 n, int lIdx, int lSize)\n"
+ "{\n"
+ " u32 blocksum;\n"
+ " int offset = 1;\n"
+ " for(int nActive=n>>1; nActive>0; nActive>>=1, offset<<=1)\n"
+ " {\n"
+ " GROUP_LDS_BARRIER;\n"
+ " for(int iIdx=lIdx; iIdx<nActive; iIdx+=lSize)\n"
+ " {\n"
+ " int ai = offset*(2*iIdx+1)-1;\n"
+ " int bi = offset*(2*iIdx+2)-1;\n"
+ " data[bi] += data[ai];\n"
+ " }\n"
+ " }\n"
+ " GROUP_LDS_BARRIER;\n"
+ " if( lIdx == 0 )\n"
+ " {\n"
+ " blocksum = data[ n-1 ];\n"
+ " data[ n-1 ] = 0;\n"
+ " }\n"
+ " GROUP_LDS_BARRIER;\n"
+ " offset >>= 1;\n"
+ " for(int nActive=1; nActive<n; nActive<<=1, offset>>=1 )\n"
+ " {\n"
+ " GROUP_LDS_BARRIER;\n"
+ " for( int iIdx = lIdx; iIdx<nActive; iIdx += lSize )\n"
+ " {\n"
+ " int ai = offset*(2*iIdx+1)-1;\n"
+ " int bi = offset*(2*iIdx+2)-1;\n"
+ " u32 temp = data[ai];\n"
+ " data[ai] = data[bi];\n"
+ " data[bi] += temp;\n"
+ " }\n"
+ " }\n"
+ " GROUP_LDS_BARRIER;\n"
+ " return blocksum;\n"
+ "}\n"
+ "__attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
+ "__kernel\n"
+ "void LocalScanKernel(__global u32* dst, __global u32 *src, __global u32 *sumBuffer,\n"
+ " uint4 cb)\n"
+ "{\n"
+ " __local u32 ldsData[WG_SIZE*2];\n"
+ " int gIdx = GET_GLOBAL_IDX;\n"
+ " int lIdx = GET_LOCAL_IDX;\n"
+ " ldsData[2*lIdx] = ( 2*gIdx < cb.m_numElems )? src[2*gIdx]: 0;\n"
+ " ldsData[2*lIdx + 1] = ( 2*gIdx+1 < cb.m_numElems )? src[2*gIdx + 1]: 0;\n"
+ " u32 sum = ScanExclusive(ldsData, WG_SIZE*2, GET_LOCAL_IDX, GET_GROUP_SIZE);\n"
+ " if( lIdx == 0 ) sumBuffer[GET_GROUP_IDX] = sum;\n"
+ " if( (2*gIdx) < cb.m_numElems )\n"
+ " {\n"
+ " dst[2*gIdx] = ldsData[2*lIdx];\n"
+ " }\n"
+ " if( (2*gIdx + 1) < cb.m_numElems )\n"
+ " {\n"
+ " dst[2*gIdx + 1] = ldsData[2*lIdx + 1];\n"
+ " }\n"
+ "}\n"
+ "__attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
+ "__kernel\n"
+ "void AddOffsetKernel(__global u32 *dst, __global u32 *blockSum, uint4 cb)\n"
+ "{\n"
+ " const u32 blockSize = WG_SIZE*2;\n"
+ " int myIdx = GET_GROUP_IDX+1;\n"
+ " int lIdx = GET_LOCAL_IDX;\n"
+ " u32 iBlockSum = blockSum[myIdx];\n"
+ " int endValue = min((myIdx+1)*(blockSize), cb.m_numElems);\n"
+ " for(int i=myIdx*blockSize+lIdx; i<endValue; i+=GET_GROUP_SIZE)\n"
+ " {\n"
+ " dst[i] += iBlockSum;\n"
+ " }\n"
+ "}\n"
+ "__attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
+ "__kernel\n"
+ "void TopLevelScanKernel(__global u32* dst, uint4 cb)\n"
+ "{\n"
+ " __local u32 ldsData[2048];\n"
+ " int gIdx = GET_GLOBAL_IDX;\n"
+ " int lIdx = GET_LOCAL_IDX;\n"
+ " int lSize = GET_GROUP_SIZE;\n"
+ " for(int i=lIdx; i<cb.m_numScanBlocks; i+=lSize )\n"
+ " {\n"
+ " ldsData[i] = (i<cb.m_numBlocks)? dst[i]:0;\n"
+ " }\n"
+ " GROUP_LDS_BARRIER;\n"
+ " u32 sum = ScanExclusive(ldsData, cb.m_numScanBlocks, GET_LOCAL_IDX, GET_GROUP_SIZE);\n"
+ " for(int i=lIdx; i<cb.m_numBlocks; i+=lSize )\n"
+ " {\n"
+ " dst[i] = ldsData[i];\n"
+ " }\n"
+ " if( gIdx == 0 )\n"
+ " {\n"
+ " dst[cb.m_numBlocks] = sum;\n"
+ " }\n"
+ "}\n";
diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/PrefixScanKernelsFloat4CL.h b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/PrefixScanKernelsFloat4CL.h
index 5b13254796..15d1bc5195 100644
--- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/PrefixScanKernelsFloat4CL.h
+++ b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/PrefixScanKernelsFloat4CL.h
@@ -1,129 +1,128 @@
//this file is autogenerated using stringify.bat (premake --stringify) in the build folder of this project
-static const char* prefixScanKernelsFloat4CL= \
-"/*\n"
-"Copyright (c) 2012 Advanced Micro Devices, Inc. \n"
-"This software is provided 'as-is', without any express or implied warranty.\n"
-"In no event will the authors be held liable for any damages arising from the use of this software.\n"
-"Permission is granted to anyone to use this software for any purpose, \n"
-"including commercial applications, and to alter it and redistribute it freely, \n"
-"subject to the following restrictions:\n"
-"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.\n"
-"2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software.\n"
-"3. This notice may not be removed or altered from any source distribution.\n"
-"*/\n"
-"//Originally written by Takahiro Harada\n"
-"typedef unsigned int u32;\n"
-"#define GET_GROUP_IDX get_group_id(0)\n"
-"#define GET_LOCAL_IDX get_local_id(0)\n"
-"#define GET_GLOBAL_IDX get_global_id(0)\n"
-"#define GET_GROUP_SIZE get_local_size(0)\n"
-"#define GROUP_LDS_BARRIER barrier(CLK_LOCAL_MEM_FENCE)\n"
-"// takahiro end\n"
-"#define WG_SIZE 128 \n"
-"#define m_numElems x\n"
-"#define m_numBlocks y\n"
-"#define m_numScanBlocks z\n"
-"/*typedef struct\n"
-"{\n"
-" uint m_numElems;\n"
-" uint m_numBlocks;\n"
-" uint m_numScanBlocks;\n"
-" uint m_padding[1];\n"
-"} ConstBuffer;\n"
-"*/\n"
-"float4 ScanExclusiveFloat4(__local float4* data, u32 n, int lIdx, int lSize)\n"
-"{\n"
-" float4 blocksum;\n"
-" int offset = 1;\n"
-" for(int nActive=n>>1; nActive>0; nActive>>=1, offset<<=1)\n"
-" {\n"
-" GROUP_LDS_BARRIER;\n"
-" for(int iIdx=lIdx; iIdx<nActive; iIdx+=lSize)\n"
-" {\n"
-" int ai = offset*(2*iIdx+1)-1;\n"
-" int bi = offset*(2*iIdx+2)-1;\n"
-" data[bi] += data[ai];\n"
-" }\n"
-" }\n"
-" GROUP_LDS_BARRIER;\n"
-" if( lIdx == 0 )\n"
-" {\n"
-" blocksum = data[ n-1 ];\n"
-" data[ n-1 ] = 0;\n"
-" }\n"
-" GROUP_LDS_BARRIER;\n"
-" offset >>= 1;\n"
-" for(int nActive=1; nActive<n; nActive<<=1, offset>>=1 )\n"
-" {\n"
-" GROUP_LDS_BARRIER;\n"
-" for( int iIdx = lIdx; iIdx<nActive; iIdx += lSize )\n"
-" {\n"
-" int ai = offset*(2*iIdx+1)-1;\n"
-" int bi = offset*(2*iIdx+2)-1;\n"
-" float4 temp = data[ai];\n"
-" data[ai] = data[bi];\n"
-" data[bi] += temp;\n"
-" }\n"
-" }\n"
-" GROUP_LDS_BARRIER;\n"
-" return blocksum;\n"
-"}\n"
-"__attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
-"__kernel\n"
-"void LocalScanKernel(__global float4* dst, __global float4* src, __global float4* sumBuffer, uint4 cb)\n"
-"{\n"
-" __local float4 ldsData[WG_SIZE*2];\n"
-" int gIdx = GET_GLOBAL_IDX;\n"
-" int lIdx = GET_LOCAL_IDX;\n"
-" ldsData[2*lIdx] = ( 2*gIdx < cb.m_numElems )? src[2*gIdx]: 0;\n"
-" ldsData[2*lIdx + 1] = ( 2*gIdx+1 < cb.m_numElems )? src[2*gIdx + 1]: 0;\n"
-" float4 sum = ScanExclusiveFloat4(ldsData, WG_SIZE*2, GET_LOCAL_IDX, GET_GROUP_SIZE);\n"
-" if( lIdx == 0 ) \n"
-" sumBuffer[GET_GROUP_IDX] = sum;\n"
-" if( (2*gIdx) < cb.m_numElems )\n"
-" {\n"
-" dst[2*gIdx] = ldsData[2*lIdx];\n"
-" }\n"
-" if( (2*gIdx + 1) < cb.m_numElems )\n"
-" {\n"
-" dst[2*gIdx + 1] = ldsData[2*lIdx + 1];\n"
-" }\n"
-"}\n"
-"__attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
-"__kernel\n"
-"void AddOffsetKernel(__global float4* dst, __global float4* blockSum, uint4 cb)\n"
-"{\n"
-" const u32 blockSize = WG_SIZE*2;\n"
-" int myIdx = GET_GROUP_IDX+1;\n"
-" int lIdx = GET_LOCAL_IDX;\n"
-" float4 iBlockSum = blockSum[myIdx];\n"
-" int endValue = min((myIdx+1)*(blockSize), cb.m_numElems);\n"
-" for(int i=myIdx*blockSize+lIdx; i<endValue; i+=GET_GROUP_SIZE)\n"
-" {\n"
-" dst[i] += iBlockSum;\n"
-" }\n"
-"}\n"
-"__attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
-"__kernel\n"
-"void TopLevelScanKernel(__global float4* dst, uint4 cb)\n"
-"{\n"
-" __local float4 ldsData[2048];\n"
-" int gIdx = GET_GLOBAL_IDX;\n"
-" int lIdx = GET_LOCAL_IDX;\n"
-" int lSize = GET_GROUP_SIZE;\n"
-" for(int i=lIdx; i<cb.m_numScanBlocks; i+=lSize )\n"
-" {\n"
-" ldsData[i] = (i<cb.m_numBlocks)? dst[i]:0;\n"
-" }\n"
-" GROUP_LDS_BARRIER;\n"
-" float4 sum = ScanExclusiveFloat4(ldsData, cb.m_numScanBlocks, GET_LOCAL_IDX, GET_GROUP_SIZE);\n"
-" for(int i=lIdx; i<cb.m_numBlocks; i+=lSize )\n"
-" {\n"
-" dst[i] = ldsData[i];\n"
-" }\n"
-" if( gIdx == 0 )\n"
-" {\n"
-" dst[cb.m_numBlocks] = sum;\n"
-" }\n"
-"}\n"
-;
+static const char* prefixScanKernelsFloat4CL =
+ "/*\n"
+ "Copyright (c) 2012 Advanced Micro Devices, Inc. \n"
+ "This software is provided 'as-is', without any express or implied warranty.\n"
+ "In no event will the authors be held liable for any damages arising from the use of this software.\n"
+ "Permission is granted to anyone to use this software for any purpose, \n"
+ "including commercial applications, and to alter it and redistribute it freely, \n"
+ "subject to the following restrictions:\n"
+ "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.\n"
+ "2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software.\n"
+ "3. This notice may not be removed or altered from any source distribution.\n"
+ "*/\n"
+ "//Originally written by Takahiro Harada\n"
+ "typedef unsigned int u32;\n"
+ "#define GET_GROUP_IDX get_group_id(0)\n"
+ "#define GET_LOCAL_IDX get_local_id(0)\n"
+ "#define GET_GLOBAL_IDX get_global_id(0)\n"
+ "#define GET_GROUP_SIZE get_local_size(0)\n"
+ "#define GROUP_LDS_BARRIER barrier(CLK_LOCAL_MEM_FENCE)\n"
+ "// takahiro end\n"
+ "#define WG_SIZE 128 \n"
+ "#define m_numElems x\n"
+ "#define m_numBlocks y\n"
+ "#define m_numScanBlocks z\n"
+ "/*typedef struct\n"
+ "{\n"
+ " uint m_numElems;\n"
+ " uint m_numBlocks;\n"
+ " uint m_numScanBlocks;\n"
+ " uint m_padding[1];\n"
+ "} ConstBuffer;\n"
+ "*/\n"
+ "float4 ScanExclusiveFloat4(__local float4* data, u32 n, int lIdx, int lSize)\n"
+ "{\n"
+ " float4 blocksum;\n"
+ " int offset = 1;\n"
+ " for(int nActive=n>>1; nActive>0; nActive>>=1, offset<<=1)\n"
+ " {\n"
+ " GROUP_LDS_BARRIER;\n"
+ " for(int iIdx=lIdx; iIdx<nActive; iIdx+=lSize)\n"
+ " {\n"
+ " int ai = offset*(2*iIdx+1)-1;\n"
+ " int bi = offset*(2*iIdx+2)-1;\n"
+ " data[bi] += data[ai];\n"
+ " }\n"
+ " }\n"
+ " GROUP_LDS_BARRIER;\n"
+ " if( lIdx == 0 )\n"
+ " {\n"
+ " blocksum = data[ n-1 ];\n"
+ " data[ n-1 ] = 0;\n"
+ " }\n"
+ " GROUP_LDS_BARRIER;\n"
+ " offset >>= 1;\n"
+ " for(int nActive=1; nActive<n; nActive<<=1, offset>>=1 )\n"
+ " {\n"
+ " GROUP_LDS_BARRIER;\n"
+ " for( int iIdx = lIdx; iIdx<nActive; iIdx += lSize )\n"
+ " {\n"
+ " int ai = offset*(2*iIdx+1)-1;\n"
+ " int bi = offset*(2*iIdx+2)-1;\n"
+ " float4 temp = data[ai];\n"
+ " data[ai] = data[bi];\n"
+ " data[bi] += temp;\n"
+ " }\n"
+ " }\n"
+ " GROUP_LDS_BARRIER;\n"
+ " return blocksum;\n"
+ "}\n"
+ "__attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
+ "__kernel\n"
+ "void LocalScanKernel(__global float4* dst, __global float4* src, __global float4* sumBuffer, uint4 cb)\n"
+ "{\n"
+ " __local float4 ldsData[WG_SIZE*2];\n"
+ " int gIdx = GET_GLOBAL_IDX;\n"
+ " int lIdx = GET_LOCAL_IDX;\n"
+ " ldsData[2*lIdx] = ( 2*gIdx < cb.m_numElems )? src[2*gIdx]: 0;\n"
+ " ldsData[2*lIdx + 1] = ( 2*gIdx+1 < cb.m_numElems )? src[2*gIdx + 1]: 0;\n"
+ " float4 sum = ScanExclusiveFloat4(ldsData, WG_SIZE*2, GET_LOCAL_IDX, GET_GROUP_SIZE);\n"
+ " if( lIdx == 0 ) \n"
+ " sumBuffer[GET_GROUP_IDX] = sum;\n"
+ " if( (2*gIdx) < cb.m_numElems )\n"
+ " {\n"
+ " dst[2*gIdx] = ldsData[2*lIdx];\n"
+ " }\n"
+ " if( (2*gIdx + 1) < cb.m_numElems )\n"
+ " {\n"
+ " dst[2*gIdx + 1] = ldsData[2*lIdx + 1];\n"
+ " }\n"
+ "}\n"
+ "__attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
+ "__kernel\n"
+ "void AddOffsetKernel(__global float4* dst, __global float4* blockSum, uint4 cb)\n"
+ "{\n"
+ " const u32 blockSize = WG_SIZE*2;\n"
+ " int myIdx = GET_GROUP_IDX+1;\n"
+ " int lIdx = GET_LOCAL_IDX;\n"
+ " float4 iBlockSum = blockSum[myIdx];\n"
+ " int endValue = min((myIdx+1)*(blockSize), cb.m_numElems);\n"
+ " for(int i=myIdx*blockSize+lIdx; i<endValue; i+=GET_GROUP_SIZE)\n"
+ " {\n"
+ " dst[i] += iBlockSum;\n"
+ " }\n"
+ "}\n"
+ "__attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
+ "__kernel\n"
+ "void TopLevelScanKernel(__global float4* dst, uint4 cb)\n"
+ "{\n"
+ " __local float4 ldsData[2048];\n"
+ " int gIdx = GET_GLOBAL_IDX;\n"
+ " int lIdx = GET_LOCAL_IDX;\n"
+ " int lSize = GET_GROUP_SIZE;\n"
+ " for(int i=lIdx; i<cb.m_numScanBlocks; i+=lSize )\n"
+ " {\n"
+ " ldsData[i] = (i<cb.m_numBlocks)? dst[i]:0;\n"
+ " }\n"
+ " GROUP_LDS_BARRIER;\n"
+ " float4 sum = ScanExclusiveFloat4(ldsData, cb.m_numScanBlocks, GET_LOCAL_IDX, GET_GROUP_SIZE);\n"
+ " for(int i=lIdx; i<cb.m_numBlocks; i+=lSize )\n"
+ " {\n"
+ " dst[i] = ldsData[i];\n"
+ " }\n"
+ " if( gIdx == 0 )\n"
+ " {\n"
+ " dst[cb.m_numBlocks] = sum;\n"
+ " }\n"
+ "}\n";
diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/RadixSort32KernelsCL.h b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/RadixSort32KernelsCL.h
index 8876c16aa6..fb4bdda303 100644
--- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/RadixSort32KernelsCL.h
+++ b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/RadixSort32KernelsCL.h
@@ -1,910 +1,909 @@
//this file is autogenerated using stringify.bat (premake --stringify) in the build folder of this project
-static const char* radixSort32KernelsCL= \
-"/*\n"
-"Bullet Continuous Collision Detection and Physics Library\n"
-"Copyright (c) 2011 Advanced Micro Devices, Inc. http://bulletphysics.org\n"
-"This software is provided 'as-is', without any express or implied warranty.\n"
-"In no event will the authors be held liable for any damages arising from the use of this software.\n"
-"Permission is granted to anyone to use this software for any purpose, \n"
-"including commercial applications, and to alter it and redistribute it freely, \n"
-"subject to the following restrictions:\n"
-"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.\n"
-"2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software.\n"
-"3. This notice may not be removed or altered from any source distribution.\n"
-"*/\n"
-"//Author Takahiro Harada\n"
-"//#pragma OPENCL EXTENSION cl_amd_printf : enable\n"
-"#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable\n"
-"#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable\n"
-"typedef unsigned int u32;\n"
-"#define GET_GROUP_IDX get_group_id(0)\n"
-"#define GET_LOCAL_IDX get_local_id(0)\n"
-"#define GET_GLOBAL_IDX get_global_id(0)\n"
-"#define GET_GROUP_SIZE get_local_size(0)\n"
-"#define GROUP_LDS_BARRIER barrier(CLK_LOCAL_MEM_FENCE)\n"
-"#define GROUP_MEM_FENCE mem_fence(CLK_LOCAL_MEM_FENCE)\n"
-"#define AtomInc(x) atom_inc(&(x))\n"
-"#define AtomInc1(x, out) out = atom_inc(&(x))\n"
-"#define AtomAdd(x, value) atom_add(&(x), value)\n"
-"#define SELECT_UINT4( b, a, condition ) select( b,a,condition )\n"
-"#define make_uint4 (uint4)\n"
-"#define make_uint2 (uint2)\n"
-"#define make_int2 (int2)\n"
-"#define WG_SIZE 64\n"
-"#define ELEMENTS_PER_WORK_ITEM (256/WG_SIZE)\n"
-"#define BITS_PER_PASS 4\n"
-"#define NUM_BUCKET (1<<BITS_PER_PASS)\n"
-"typedef uchar u8;\n"
-"// this isn't optimization for VLIW. But just reducing writes. \n"
-"#define USE_2LEVEL_REDUCE 1\n"
-"//#define CHECK_BOUNDARY 1\n"
-"//#define NV_GPU 1\n"
-"// Cypress\n"
-"#define nPerWI 16\n"
-"// Cayman\n"
-"//#define nPerWI 20\n"
-"#define m_n x\n"
-"#define m_nWGs y\n"
-"#define m_startBit z\n"
-"#define m_nBlocksPerWG w\n"
-"/*\n"
-"typedef struct\n"
-"{\n"
-" int m_n;\n"
-" int m_nWGs;\n"
-" int m_startBit;\n"
-" int m_nBlocksPerWG;\n"
-"} ConstBuffer;\n"
-"*/\n"
-"typedef struct\n"
-"{\n"
-" unsigned int m_key;\n"
-" unsigned int m_value;\n"
-"} SortDataCL;\n"
-"uint prefixScanVectorEx( uint4* data )\n"
-"{\n"
-" u32 sum = 0;\n"
-" u32 tmp = data[0].x;\n"
-" data[0].x = sum;\n"
-" sum += tmp;\n"
-" tmp = data[0].y;\n"
-" data[0].y = sum;\n"
-" sum += tmp;\n"
-" tmp = data[0].z;\n"
-" data[0].z = sum;\n"
-" sum += tmp;\n"
-" tmp = data[0].w;\n"
-" data[0].w = sum;\n"
-" sum += tmp;\n"
-" return sum;\n"
-"}\n"
-"u32 localPrefixSum( u32 pData, uint lIdx, uint* totalSum, __local u32* sorterSharedMemory, int wgSize /*64 or 128*/ )\n"
-"{\n"
-" { // Set data\n"
-" sorterSharedMemory[lIdx] = 0;\n"
-" sorterSharedMemory[lIdx+wgSize] = pData;\n"
-" }\n"
-" GROUP_LDS_BARRIER;\n"
-" { // Prefix sum\n"
-" int idx = 2*lIdx + (wgSize+1);\n"
-"#if defined(USE_2LEVEL_REDUCE)\n"
-" if( lIdx < 64 )\n"
-" {\n"
-" u32 u0, u1, u2;\n"
-" u0 = sorterSharedMemory[idx-3];\n"
-" u1 = sorterSharedMemory[idx-2];\n"
-" u2 = sorterSharedMemory[idx-1];\n"
-" AtomAdd( sorterSharedMemory[idx], u0+u1+u2 ); \n"
-" GROUP_MEM_FENCE;\n"
-" u0 = sorterSharedMemory[idx-12];\n"
-" u1 = sorterSharedMemory[idx-8];\n"
-" u2 = sorterSharedMemory[idx-4];\n"
-" AtomAdd( sorterSharedMemory[idx], u0+u1+u2 ); \n"
-" GROUP_MEM_FENCE;\n"
-" u0 = sorterSharedMemory[idx-48];\n"
-" u1 = sorterSharedMemory[idx-32];\n"
-" u2 = sorterSharedMemory[idx-16];\n"
-" AtomAdd( sorterSharedMemory[idx], u0+u1+u2 ); \n"
-" GROUP_MEM_FENCE;\n"
-" if( wgSize > 64 )\n"
-" {\n"
-" sorterSharedMemory[idx] += sorterSharedMemory[idx-64];\n"
-" GROUP_MEM_FENCE;\n"
-" }\n"
-" sorterSharedMemory[idx-1] += sorterSharedMemory[idx-2];\n"
-" GROUP_MEM_FENCE;\n"
-" }\n"
-"#else\n"
-" if( lIdx < 64 )\n"
-" {\n"
-" sorterSharedMemory[idx] += sorterSharedMemory[idx-1];\n"
-" GROUP_MEM_FENCE;\n"
-" sorterSharedMemory[idx] += sorterSharedMemory[idx-2]; \n"
-" GROUP_MEM_FENCE;\n"
-" sorterSharedMemory[idx] += sorterSharedMemory[idx-4];\n"
-" GROUP_MEM_FENCE;\n"
-" sorterSharedMemory[idx] += sorterSharedMemory[idx-8];\n"
-" GROUP_MEM_FENCE;\n"
-" sorterSharedMemory[idx] += sorterSharedMemory[idx-16];\n"
-" GROUP_MEM_FENCE;\n"
-" sorterSharedMemory[idx] += sorterSharedMemory[idx-32];\n"
-" GROUP_MEM_FENCE;\n"
-" if( wgSize > 64 )\n"
-" {\n"
-" sorterSharedMemory[idx] += sorterSharedMemory[idx-64];\n"
-" GROUP_MEM_FENCE;\n"
-" }\n"
-" sorterSharedMemory[idx-1] += sorterSharedMemory[idx-2];\n"
-" GROUP_MEM_FENCE;\n"
-" }\n"
-"#endif\n"
-" }\n"
-" GROUP_LDS_BARRIER;\n"
-" *totalSum = sorterSharedMemory[wgSize*2-1];\n"
-" u32 addValue = sorterSharedMemory[lIdx+wgSize-1];\n"
-" return addValue;\n"
-"}\n"
-"//__attribute__((reqd_work_group_size(128,1,1)))\n"
-"uint4 localPrefixSum128V( uint4 pData, uint lIdx, uint* totalSum, __local u32* sorterSharedMemory )\n"
-"{\n"
-" u32 s4 = prefixScanVectorEx( &pData );\n"
-" u32 rank = localPrefixSum( s4, lIdx, totalSum, sorterSharedMemory, 128 );\n"
-" return pData + make_uint4( rank, rank, rank, rank );\n"
-"}\n"
-"//__attribute__((reqd_work_group_size(64,1,1)))\n"
-"uint4 localPrefixSum64V( uint4 pData, uint lIdx, uint* totalSum, __local u32* sorterSharedMemory )\n"
-"{\n"
-" u32 s4 = prefixScanVectorEx( &pData );\n"
-" u32 rank = localPrefixSum( s4, lIdx, totalSum, sorterSharedMemory, 64 );\n"
-" return pData + make_uint4( rank, rank, rank, rank );\n"
-"}\n"
-"u32 unpack4Key( u32 key, int keyIdx ){ return (key>>(keyIdx*8)) & 0xff;}\n"
-"u32 bit8Scan(u32 v)\n"
-"{\n"
-" return (v<<8) + (v<<16) + (v<<24);\n"
-"}\n"
-"//===\n"
-"#define MY_HISTOGRAM(idx) localHistogramMat[(idx)*WG_SIZE+lIdx]\n"
-"__kernel\n"
-"__attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
-"void StreamCountKernel( __global u32* gSrc, __global u32* histogramOut, int4 cb )\n"
-"{\n"
-" __local u32 localHistogramMat[NUM_BUCKET*WG_SIZE];\n"
-" u32 gIdx = GET_GLOBAL_IDX;\n"
-" u32 lIdx = GET_LOCAL_IDX;\n"
-" u32 wgIdx = GET_GROUP_IDX;\n"
-" u32 wgSize = GET_GROUP_SIZE;\n"
-" const int startBit = cb.m_startBit;\n"
-" const int n = cb.m_n;\n"
-" const int nWGs = cb.m_nWGs;\n"
-" const int nBlocksPerWG = cb.m_nBlocksPerWG;\n"
-" for(int i=0; i<NUM_BUCKET; i++)\n"
-" {\n"
-" MY_HISTOGRAM(i) = 0;\n"
-" }\n"
-" GROUP_LDS_BARRIER;\n"
-" const int blockSize = ELEMENTS_PER_WORK_ITEM*WG_SIZE;\n"
-" u32 localKey;\n"
-" int nBlocks = (n)/blockSize - nBlocksPerWG*wgIdx;\n"
-" int addr = blockSize*nBlocksPerWG*wgIdx + ELEMENTS_PER_WORK_ITEM*lIdx;\n"
-" for(int iblock=0; iblock<min(nBlocksPerWG, nBlocks); iblock++, addr+=blockSize)\n"
-" {\n"
-" // MY_HISTOGRAM( localKeys.x ) ++ is much expensive than atomic add as it requires read and write while atomics can just add on AMD\n"
-" // Using registers didn't perform well. It seems like use localKeys to address requires a lot of alu ops\n"
-" // AMD: AtomInc performs better while NV prefers ++\n"
-" for(int i=0; i<ELEMENTS_PER_WORK_ITEM; i++)\n"
-" {\n"
-"#if defined(CHECK_BOUNDARY)\n"
-" if( addr+i < n )\n"
-"#endif\n"
-" {\n"
-" localKey = (gSrc[addr+i]>>startBit) & 0xf;\n"
-"#if defined(NV_GPU)\n"
-" MY_HISTOGRAM( localKey )++;\n"
-"#else\n"
-" AtomInc( MY_HISTOGRAM( localKey ) );\n"
-"#endif\n"
-" }\n"
-" }\n"
-" }\n"
-" GROUP_LDS_BARRIER;\n"
-" \n"
-" if( lIdx < NUM_BUCKET )\n"
-" {\n"
-" u32 sum = 0;\n"
-" for(int i=0; i<GET_GROUP_SIZE; i++)\n"
-" {\n"
-" sum += localHistogramMat[lIdx*WG_SIZE+(i+lIdx)%GET_GROUP_SIZE];\n"
-" }\n"
-" histogramOut[lIdx*nWGs+wgIdx] = sum;\n"
-" }\n"
-"}\n"
-"__kernel\n"
-"__attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
-"void StreamCountSortDataKernel( __global SortDataCL* gSrc, __global u32* histogramOut, int4 cb )\n"
-"{\n"
-" __local u32 localHistogramMat[NUM_BUCKET*WG_SIZE];\n"
-" u32 gIdx = GET_GLOBAL_IDX;\n"
-" u32 lIdx = GET_LOCAL_IDX;\n"
-" u32 wgIdx = GET_GROUP_IDX;\n"
-" u32 wgSize = GET_GROUP_SIZE;\n"
-" const int startBit = cb.m_startBit;\n"
-" const int n = cb.m_n;\n"
-" const int nWGs = cb.m_nWGs;\n"
-" const int nBlocksPerWG = cb.m_nBlocksPerWG;\n"
-" for(int i=0; i<NUM_BUCKET; i++)\n"
-" {\n"
-" MY_HISTOGRAM(i) = 0;\n"
-" }\n"
-" GROUP_LDS_BARRIER;\n"
-" const int blockSize = ELEMENTS_PER_WORK_ITEM*WG_SIZE;\n"
-" u32 localKey;\n"
-" int nBlocks = (n)/blockSize - nBlocksPerWG*wgIdx;\n"
-" int addr = blockSize*nBlocksPerWG*wgIdx + ELEMENTS_PER_WORK_ITEM*lIdx;\n"
-" for(int iblock=0; iblock<min(nBlocksPerWG, nBlocks); iblock++, addr+=blockSize)\n"
-" {\n"
-" // MY_HISTOGRAM( localKeys.x ) ++ is much expensive than atomic add as it requires read and write while atomics can just add on AMD\n"
-" // Using registers didn't perform well. It seems like use localKeys to address requires a lot of alu ops\n"
-" // AMD: AtomInc performs better while NV prefers ++\n"
-" for(int i=0; i<ELEMENTS_PER_WORK_ITEM; i++)\n"
-" {\n"
-"#if defined(CHECK_BOUNDARY)\n"
-" if( addr+i < n )\n"
-"#endif\n"
-" {\n"
-" localKey = (gSrc[addr+i].m_key>>startBit) & 0xf;\n"
-"#if defined(NV_GPU)\n"
-" MY_HISTOGRAM( localKey )++;\n"
-"#else\n"
-" AtomInc( MY_HISTOGRAM( localKey ) );\n"
-"#endif\n"
-" }\n"
-" }\n"
-" }\n"
-" GROUP_LDS_BARRIER;\n"
-" \n"
-" if( lIdx < NUM_BUCKET )\n"
-" {\n"
-" u32 sum = 0;\n"
-" for(int i=0; i<GET_GROUP_SIZE; i++)\n"
-" {\n"
-" sum += localHistogramMat[lIdx*WG_SIZE+(i+lIdx)%GET_GROUP_SIZE];\n"
-" }\n"
-" histogramOut[lIdx*nWGs+wgIdx] = sum;\n"
-" }\n"
-"}\n"
-"#define nPerLane (nPerWI/4)\n"
-"// NUM_BUCKET*nWGs < 128*nPerWI\n"
-"__kernel\n"
-"__attribute__((reqd_work_group_size(128,1,1)))\n"
-"void PrefixScanKernel( __global u32* wHistogram1, int4 cb )\n"
-"{\n"
-" __local u32 ldsTopScanData[128*2];\n"
-" u32 lIdx = GET_LOCAL_IDX;\n"
-" u32 wgIdx = GET_GROUP_IDX;\n"
-" const int nWGs = cb.m_nWGs;\n"
-" u32 data[nPerWI];\n"
-" for(int i=0; i<nPerWI; i++)\n"
-" {\n"
-" data[i] = 0;\n"
-" if( (nPerWI*lIdx+i) < NUM_BUCKET*nWGs )\n"
-" data[i] = wHistogram1[nPerWI*lIdx+i];\n"
-" }\n"
-" uint4 myData = make_uint4(0,0,0,0);\n"
-" for(int i=0; i<nPerLane; i++)\n"
-" {\n"
-" myData.x += data[nPerLane*0+i];\n"
-" myData.y += data[nPerLane*1+i];\n"
-" myData.z += data[nPerLane*2+i];\n"
-" myData.w += data[nPerLane*3+i];\n"
-" }\n"
-" uint totalSum;\n"
-" uint4 scanned = localPrefixSum128V( myData, lIdx, &totalSum, ldsTopScanData );\n"
-"// for(int j=0; j<4; j++) // somehow it introduces a lot of branches\n"
-" { int j = 0;\n"
-" u32 sum = 0;\n"
-" for(int i=0; i<nPerLane; i++)\n"
-" {\n"
-" u32 tmp = data[nPerLane*j+i];\n"
-" data[nPerLane*j+i] = sum;\n"
-" sum += tmp;\n"
-" }\n"
-" }\n"
-" { int j = 1;\n"
-" u32 sum = 0;\n"
-" for(int i=0; i<nPerLane; i++)\n"
-" {\n"
-" u32 tmp = data[nPerLane*j+i];\n"
-" data[nPerLane*j+i] = sum;\n"
-" sum += tmp;\n"
-" }\n"
-" }\n"
-" { int j = 2;\n"
-" u32 sum = 0;\n"
-" for(int i=0; i<nPerLane; i++)\n"
-" {\n"
-" u32 tmp = data[nPerLane*j+i];\n"
-" data[nPerLane*j+i] = sum;\n"
-" sum += tmp;\n"
-" }\n"
-" }\n"
-" { int j = 3;\n"
-" u32 sum = 0;\n"
-" for(int i=0; i<nPerLane; i++)\n"
-" {\n"
-" u32 tmp = data[nPerLane*j+i];\n"
-" data[nPerLane*j+i] = sum;\n"
-" sum += tmp;\n"
-" }\n"
-" }\n"
-" for(int i=0; i<nPerLane; i++)\n"
-" {\n"
-" data[nPerLane*0+i] += scanned.x;\n"
-" data[nPerLane*1+i] += scanned.y;\n"
-" data[nPerLane*2+i] += scanned.z;\n"
-" data[nPerLane*3+i] += scanned.w;\n"
-" }\n"
-" for(int i=0; i<nPerWI; i++)\n"
-" {\n"
-" int index = nPerWI*lIdx+i;\n"
-" if (index < NUM_BUCKET*nWGs)\n"
-" wHistogram1[nPerWI*lIdx+i] = data[i];\n"
-" }\n"
-"}\n"
-"// 4 scan, 4 exchange\n"
-"void sort4Bits(u32 sortData[4], int startBit, int lIdx, __local u32* ldsSortData)\n"
-"{\n"
-" for(int bitIdx=0; bitIdx<BITS_PER_PASS; bitIdx++)\n"
-" {\n"
-" u32 mask = (1<<bitIdx);\n"
-" uint4 cmpResult = make_uint4( (sortData[0]>>startBit) & mask, (sortData[1]>>startBit) & mask, (sortData[2]>>startBit) & mask, (sortData[3]>>startBit) & mask );\n"
-" uint4 prefixSum = SELECT_UINT4( make_uint4(1,1,1,1), make_uint4(0,0,0,0), cmpResult != make_uint4(0,0,0,0) );\n"
-" u32 total;\n"
-" prefixSum = localPrefixSum64V( prefixSum, lIdx, &total, ldsSortData );\n"
-" {\n"
-" uint4 localAddr = make_uint4(lIdx*4+0,lIdx*4+1,lIdx*4+2,lIdx*4+3);\n"
-" uint4 dstAddr = localAddr - prefixSum + make_uint4( total, total, total, total );\n"
-" dstAddr = SELECT_UINT4( prefixSum, dstAddr, cmpResult != make_uint4(0, 0, 0, 0) );\n"
-" GROUP_LDS_BARRIER;\n"
-" ldsSortData[dstAddr.x] = sortData[0];\n"
-" ldsSortData[dstAddr.y] = sortData[1];\n"
-" ldsSortData[dstAddr.z] = sortData[2];\n"
-" ldsSortData[dstAddr.w] = sortData[3];\n"
-" GROUP_LDS_BARRIER;\n"
-" sortData[0] = ldsSortData[localAddr.x];\n"
-" sortData[1] = ldsSortData[localAddr.y];\n"
-" sortData[2] = ldsSortData[localAddr.z];\n"
-" sortData[3] = ldsSortData[localAddr.w];\n"
-" GROUP_LDS_BARRIER;\n"
-" }\n"
-" }\n"
-"}\n"
-"// 2 scan, 2 exchange\n"
-"void sort4Bits1(u32 sortData[4], int startBit, int lIdx, __local u32* ldsSortData)\n"
-"{\n"
-" for(uint ibit=0; ibit<BITS_PER_PASS; ibit+=2)\n"
-" {\n"
-" uint4 b = make_uint4((sortData[0]>>(startBit+ibit)) & 0x3, \n"
-" (sortData[1]>>(startBit+ibit)) & 0x3, \n"
-" (sortData[2]>>(startBit+ibit)) & 0x3, \n"
-" (sortData[3]>>(startBit+ibit)) & 0x3);\n"
-" u32 key4;\n"
-" u32 sKeyPacked[4] = { 0, 0, 0, 0 };\n"
-" {\n"
-" sKeyPacked[0] |= 1<<(8*b.x);\n"
-" sKeyPacked[1] |= 1<<(8*b.y);\n"
-" sKeyPacked[2] |= 1<<(8*b.z);\n"
-" sKeyPacked[3] |= 1<<(8*b.w);\n"
-" key4 = sKeyPacked[0] + sKeyPacked[1] + sKeyPacked[2] + sKeyPacked[3];\n"
-" }\n"
-" u32 rankPacked;\n"
-" u32 sumPacked;\n"
-" {\n"
-" rankPacked = localPrefixSum( key4, lIdx, &sumPacked, ldsSortData, WG_SIZE );\n"
-" }\n"
-" GROUP_LDS_BARRIER;\n"
-" u32 newOffset[4] = { 0,0,0,0 };\n"
-" {\n"
-" u32 sumScanned = bit8Scan( sumPacked );\n"
-" u32 scannedKeys[4];\n"
-" scannedKeys[0] = 1<<(8*b.x);\n"
-" scannedKeys[1] = 1<<(8*b.y);\n"
-" scannedKeys[2] = 1<<(8*b.z);\n"
-" scannedKeys[3] = 1<<(8*b.w);\n"
-" { // 4 scans at once\n"
-" u32 sum4 = 0;\n"
-" for(int ie=0; ie<4; ie++)\n"
-" {\n"
-" u32 tmp = scannedKeys[ie];\n"
-" scannedKeys[ie] = sum4;\n"
-" sum4 += tmp;\n"
-" }\n"
-" }\n"
-" {\n"
-" u32 sumPlusRank = sumScanned + rankPacked;\n"
-" { u32 ie = b.x;\n"
-" scannedKeys[0] += sumPlusRank;\n"
-" newOffset[0] = unpack4Key( scannedKeys[0], ie );\n"
-" }\n"
-" { u32 ie = b.y;\n"
-" scannedKeys[1] += sumPlusRank;\n"
-" newOffset[1] = unpack4Key( scannedKeys[1], ie );\n"
-" }\n"
-" { u32 ie = b.z;\n"
-" scannedKeys[2] += sumPlusRank;\n"
-" newOffset[2] = unpack4Key( scannedKeys[2], ie );\n"
-" }\n"
-" { u32 ie = b.w;\n"
-" scannedKeys[3] += sumPlusRank;\n"
-" newOffset[3] = unpack4Key( scannedKeys[3], ie );\n"
-" }\n"
-" }\n"
-" }\n"
-" GROUP_LDS_BARRIER;\n"
-" {\n"
-" ldsSortData[newOffset[0]] = sortData[0];\n"
-" ldsSortData[newOffset[1]] = sortData[1];\n"
-" ldsSortData[newOffset[2]] = sortData[2];\n"
-" ldsSortData[newOffset[3]] = sortData[3];\n"
-" GROUP_LDS_BARRIER;\n"
-" u32 dstAddr = 4*lIdx;\n"
-" sortData[0] = ldsSortData[dstAddr+0];\n"
-" sortData[1] = ldsSortData[dstAddr+1];\n"
-" sortData[2] = ldsSortData[dstAddr+2];\n"
-" sortData[3] = ldsSortData[dstAddr+3];\n"
-" GROUP_LDS_BARRIER;\n"
-" }\n"
-" }\n"
-"}\n"
-"#define SET_HISTOGRAM(setIdx, key) ldsSortData[(setIdx)*NUM_BUCKET+key]\n"
-"__kernel\n"
-"__attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
-"void SortAndScatterKernel( __global const u32* restrict gSrc, __global const u32* rHistogram, __global u32* restrict gDst, int4 cb )\n"
-"{\n"
-" __local u32 ldsSortData[WG_SIZE*ELEMENTS_PER_WORK_ITEM+16];\n"
-" __local u32 localHistogramToCarry[NUM_BUCKET];\n"
-" __local u32 localHistogram[NUM_BUCKET*2];\n"
-" u32 gIdx = GET_GLOBAL_IDX;\n"
-" u32 lIdx = GET_LOCAL_IDX;\n"
-" u32 wgIdx = GET_GROUP_IDX;\n"
-" u32 wgSize = GET_GROUP_SIZE;\n"
-" const int n = cb.m_n;\n"
-" const int nWGs = cb.m_nWGs;\n"
-" const int startBit = cb.m_startBit;\n"
-" const int nBlocksPerWG = cb.m_nBlocksPerWG;\n"
-" if( lIdx < (NUM_BUCKET) )\n"
-" {\n"
-" localHistogramToCarry[lIdx] = rHistogram[lIdx*nWGs + wgIdx];\n"
-" }\n"
-" GROUP_LDS_BARRIER;\n"
-" const int blockSize = ELEMENTS_PER_WORK_ITEM*WG_SIZE;\n"
-" int nBlocks = n/blockSize - nBlocksPerWG*wgIdx;\n"
-" int addr = blockSize*nBlocksPerWG*wgIdx + ELEMENTS_PER_WORK_ITEM*lIdx;\n"
-" for(int iblock=0; iblock<min(nBlocksPerWG, nBlocks); iblock++, addr+=blockSize)\n"
-" {\n"
-" u32 myHistogram = 0;\n"
-" u32 sortData[ELEMENTS_PER_WORK_ITEM];\n"
-" for(int i=0; i<ELEMENTS_PER_WORK_ITEM; i++)\n"
-"#if defined(CHECK_BOUNDARY)\n"
-" sortData[i] = ( addr+i < n )? gSrc[ addr+i ] : 0xffffffff;\n"
-"#else\n"
-" sortData[i] = gSrc[ addr+i ];\n"
-"#endif\n"
-" sort4Bits(sortData, startBit, lIdx, ldsSortData);\n"
-" u32 keys[ELEMENTS_PER_WORK_ITEM];\n"
-" for(int i=0; i<ELEMENTS_PER_WORK_ITEM; i++)\n"
-" keys[i] = (sortData[i]>>startBit) & 0xf;\n"
-" { // create histogram\n"
-" u32 setIdx = lIdx/16;\n"
-" if( lIdx < NUM_BUCKET )\n"
-" {\n"
-" localHistogram[lIdx] = 0;\n"
-" }\n"
-" ldsSortData[lIdx] = 0;\n"
-" GROUP_LDS_BARRIER;\n"
-" for(int i=0; i<ELEMENTS_PER_WORK_ITEM; i++)\n"
-"#if defined(CHECK_BOUNDARY)\n"
-" if( addr+i < n )\n"
-"#endif\n"
-"#if defined(NV_GPU)\n"
-" SET_HISTOGRAM( setIdx, keys[i] )++;\n"
-"#else\n"
-" AtomInc( SET_HISTOGRAM( setIdx, keys[i] ) );\n"
-"#endif\n"
-" \n"
-" GROUP_LDS_BARRIER;\n"
-" \n"
-" uint hIdx = NUM_BUCKET+lIdx;\n"
-" if( lIdx < NUM_BUCKET )\n"
-" {\n"
-" u32 sum = 0;\n"
-" for(int i=0; i<WG_SIZE/16; i++)\n"
-" {\n"
-" sum += SET_HISTOGRAM( i, lIdx );\n"
-" }\n"
-" myHistogram = sum;\n"
-" localHistogram[hIdx] = sum;\n"
-" }\n"
-" GROUP_LDS_BARRIER;\n"
-"#if defined(USE_2LEVEL_REDUCE)\n"
-" if( lIdx < NUM_BUCKET )\n"
-" {\n"
-" localHistogram[hIdx] = localHistogram[hIdx-1];\n"
-" GROUP_MEM_FENCE;\n"
-" u32 u0, u1, u2;\n"
-" u0 = localHistogram[hIdx-3];\n"
-" u1 = localHistogram[hIdx-2];\n"
-" u2 = localHistogram[hIdx-1];\n"
-" AtomAdd( localHistogram[hIdx], u0 + u1 + u2 );\n"
-" GROUP_MEM_FENCE;\n"
-" u0 = localHistogram[hIdx-12];\n"
-" u1 = localHistogram[hIdx-8];\n"
-" u2 = localHistogram[hIdx-4];\n"
-" AtomAdd( localHistogram[hIdx], u0 + u1 + u2 );\n"
-" GROUP_MEM_FENCE;\n"
-" }\n"
-"#else\n"
-" if( lIdx < NUM_BUCKET )\n"
-" {\n"
-" localHistogram[hIdx] = localHistogram[hIdx-1];\n"
-" GROUP_MEM_FENCE;\n"
-" localHistogram[hIdx] += localHistogram[hIdx-1];\n"
-" GROUP_MEM_FENCE;\n"
-" localHistogram[hIdx] += localHistogram[hIdx-2];\n"
-" GROUP_MEM_FENCE;\n"
-" localHistogram[hIdx] += localHistogram[hIdx-4];\n"
-" GROUP_MEM_FENCE;\n"
-" localHistogram[hIdx] += localHistogram[hIdx-8];\n"
-" GROUP_MEM_FENCE;\n"
-" }\n"
-"#endif\n"
-" GROUP_LDS_BARRIER;\n"
-" }\n"
-" {\n"
-" for(int ie=0; ie<ELEMENTS_PER_WORK_ITEM; ie++)\n"
-" {\n"
-" int dataIdx = ELEMENTS_PER_WORK_ITEM*lIdx+ie;\n"
-" int binIdx = keys[ie];\n"
-" int groupOffset = localHistogramToCarry[binIdx];\n"
-" int myIdx = dataIdx - localHistogram[NUM_BUCKET+binIdx];\n"
-"#if defined(CHECK_BOUNDARY)\n"
-" if( addr+ie < n )\n"
-"#endif\n"
-" gDst[ groupOffset + myIdx ] = sortData[ie];\n"
-" }\n"
-" }\n"
-" GROUP_LDS_BARRIER;\n"
-" if( lIdx < NUM_BUCKET )\n"
-" {\n"
-" localHistogramToCarry[lIdx] += myHistogram;\n"
-" }\n"
-" GROUP_LDS_BARRIER;\n"
-" }\n"
-"}\n"
-"// 2 scan, 2 exchange\n"
-"void sort4Bits1KeyValue(u32 sortData[4], int sortVal[4], int startBit, int lIdx, __local u32* ldsSortData, __local int *ldsSortVal)\n"
-"{\n"
-" for(uint ibit=0; ibit<BITS_PER_PASS; ibit+=2)\n"
-" {\n"
-" uint4 b = make_uint4((sortData[0]>>(startBit+ibit)) & 0x3, \n"
-" (sortData[1]>>(startBit+ibit)) & 0x3, \n"
-" (sortData[2]>>(startBit+ibit)) & 0x3, \n"
-" (sortData[3]>>(startBit+ibit)) & 0x3);\n"
-" u32 key4;\n"
-" u32 sKeyPacked[4] = { 0, 0, 0, 0 };\n"
-" {\n"
-" sKeyPacked[0] |= 1<<(8*b.x);\n"
-" sKeyPacked[1] |= 1<<(8*b.y);\n"
-" sKeyPacked[2] |= 1<<(8*b.z);\n"
-" sKeyPacked[3] |= 1<<(8*b.w);\n"
-" key4 = sKeyPacked[0] + sKeyPacked[1] + sKeyPacked[2] + sKeyPacked[3];\n"
-" }\n"
-" u32 rankPacked;\n"
-" u32 sumPacked;\n"
-" {\n"
-" rankPacked = localPrefixSum( key4, lIdx, &sumPacked, ldsSortData, WG_SIZE );\n"
-" }\n"
-" GROUP_LDS_BARRIER;\n"
-" u32 newOffset[4] = { 0,0,0,0 };\n"
-" {\n"
-" u32 sumScanned = bit8Scan( sumPacked );\n"
-" u32 scannedKeys[4];\n"
-" scannedKeys[0] = 1<<(8*b.x);\n"
-" scannedKeys[1] = 1<<(8*b.y);\n"
-" scannedKeys[2] = 1<<(8*b.z);\n"
-" scannedKeys[3] = 1<<(8*b.w);\n"
-" { // 4 scans at once\n"
-" u32 sum4 = 0;\n"
-" for(int ie=0; ie<4; ie++)\n"
-" {\n"
-" u32 tmp = scannedKeys[ie];\n"
-" scannedKeys[ie] = sum4;\n"
-" sum4 += tmp;\n"
-" }\n"
-" }\n"
-" {\n"
-" u32 sumPlusRank = sumScanned + rankPacked;\n"
-" { u32 ie = b.x;\n"
-" scannedKeys[0] += sumPlusRank;\n"
-" newOffset[0] = unpack4Key( scannedKeys[0], ie );\n"
-" }\n"
-" { u32 ie = b.y;\n"
-" scannedKeys[1] += sumPlusRank;\n"
-" newOffset[1] = unpack4Key( scannedKeys[1], ie );\n"
-" }\n"
-" { u32 ie = b.z;\n"
-" scannedKeys[2] += sumPlusRank;\n"
-" newOffset[2] = unpack4Key( scannedKeys[2], ie );\n"
-" }\n"
-" { u32 ie = b.w;\n"
-" scannedKeys[3] += sumPlusRank;\n"
-" newOffset[3] = unpack4Key( scannedKeys[3], ie );\n"
-" }\n"
-" }\n"
-" }\n"
-" GROUP_LDS_BARRIER;\n"
-" {\n"
-" ldsSortData[newOffset[0]] = sortData[0];\n"
-" ldsSortData[newOffset[1]] = sortData[1];\n"
-" ldsSortData[newOffset[2]] = sortData[2];\n"
-" ldsSortData[newOffset[3]] = sortData[3];\n"
-" ldsSortVal[newOffset[0]] = sortVal[0];\n"
-" ldsSortVal[newOffset[1]] = sortVal[1];\n"
-" ldsSortVal[newOffset[2]] = sortVal[2];\n"
-" ldsSortVal[newOffset[3]] = sortVal[3];\n"
-" GROUP_LDS_BARRIER;\n"
-" u32 dstAddr = 4*lIdx;\n"
-" sortData[0] = ldsSortData[dstAddr+0];\n"
-" sortData[1] = ldsSortData[dstAddr+1];\n"
-" sortData[2] = ldsSortData[dstAddr+2];\n"
-" sortData[3] = ldsSortData[dstAddr+3];\n"
-" sortVal[0] = ldsSortVal[dstAddr+0];\n"
-" sortVal[1] = ldsSortVal[dstAddr+1];\n"
-" sortVal[2] = ldsSortVal[dstAddr+2];\n"
-" sortVal[3] = ldsSortVal[dstAddr+3];\n"
-" GROUP_LDS_BARRIER;\n"
-" }\n"
-" }\n"
-"}\n"
-"__kernel\n"
-"__attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
-"void SortAndScatterSortDataKernel( __global const SortDataCL* restrict gSrc, __global const u32* rHistogram, __global SortDataCL* restrict gDst, int4 cb)\n"
-"{\n"
-" __local int ldsSortData[WG_SIZE*ELEMENTS_PER_WORK_ITEM+16];\n"
-" __local int ldsSortVal[WG_SIZE*ELEMENTS_PER_WORK_ITEM+16];\n"
-" __local u32 localHistogramToCarry[NUM_BUCKET];\n"
-" __local u32 localHistogram[NUM_BUCKET*2];\n"
-" u32 gIdx = GET_GLOBAL_IDX;\n"
-" u32 lIdx = GET_LOCAL_IDX;\n"
-" u32 wgIdx = GET_GROUP_IDX;\n"
-" u32 wgSize = GET_GROUP_SIZE;\n"
-" const int n = cb.m_n;\n"
-" const int nWGs = cb.m_nWGs;\n"
-" const int startBit = cb.m_startBit;\n"
-" const int nBlocksPerWG = cb.m_nBlocksPerWG;\n"
-" if( lIdx < (NUM_BUCKET) )\n"
-" {\n"
-" localHistogramToCarry[lIdx] = rHistogram[lIdx*nWGs + wgIdx];\n"
-" }\n"
-" GROUP_LDS_BARRIER;\n"
-" \n"
-" const int blockSize = ELEMENTS_PER_WORK_ITEM*WG_SIZE;\n"
-" int nBlocks = n/blockSize - nBlocksPerWG*wgIdx;\n"
-" int addr = blockSize*nBlocksPerWG*wgIdx + ELEMENTS_PER_WORK_ITEM*lIdx;\n"
-" for(int iblock=0; iblock<min(nBlocksPerWG, nBlocks); iblock++, addr+=blockSize)\n"
-" {\n"
-" u32 myHistogram = 0;\n"
-" int sortData[ELEMENTS_PER_WORK_ITEM];\n"
-" int sortVal[ELEMENTS_PER_WORK_ITEM];\n"
-" for(int i=0; i<ELEMENTS_PER_WORK_ITEM; i++)\n"
-"#if defined(CHECK_BOUNDARY)\n"
-" {\n"
-" sortData[i] = ( addr+i < n )? gSrc[ addr+i ].m_key : 0xffffffff;\n"
-" sortVal[i] = ( addr+i < n )? gSrc[ addr+i ].m_value : 0xffffffff;\n"
-" }\n"
-"#else\n"
-" {\n"
-" sortData[i] = gSrc[ addr+i ].m_key;\n"
-" sortVal[i] = gSrc[ addr+i ].m_value;\n"
-" }\n"
-"#endif\n"
-" sort4Bits1KeyValue(sortData, sortVal, startBit, lIdx, ldsSortData, ldsSortVal);\n"
-" u32 keys[ELEMENTS_PER_WORK_ITEM];\n"
-" for(int i=0; i<ELEMENTS_PER_WORK_ITEM; i++)\n"
-" keys[i] = (sortData[i]>>startBit) & 0xf;\n"
-" { // create histogram\n"
-" u32 setIdx = lIdx/16;\n"
-" if( lIdx < NUM_BUCKET )\n"
-" {\n"
-" localHistogram[lIdx] = 0;\n"
-" }\n"
-" ldsSortData[lIdx] = 0;\n"
-" GROUP_LDS_BARRIER;\n"
-" for(int i=0; i<ELEMENTS_PER_WORK_ITEM; i++)\n"
-"#if defined(CHECK_BOUNDARY)\n"
-" if( addr+i < n )\n"
-"#endif\n"
-"#if defined(NV_GPU)\n"
-" SET_HISTOGRAM( setIdx, keys[i] )++;\n"
-"#else\n"
-" AtomInc( SET_HISTOGRAM( setIdx, keys[i] ) );\n"
-"#endif\n"
-" \n"
-" GROUP_LDS_BARRIER;\n"
-" \n"
-" uint hIdx = NUM_BUCKET+lIdx;\n"
-" if( lIdx < NUM_BUCKET )\n"
-" {\n"
-" u32 sum = 0;\n"
-" for(int i=0; i<WG_SIZE/16; i++)\n"
-" {\n"
-" sum += SET_HISTOGRAM( i, lIdx );\n"
-" }\n"
-" myHistogram = sum;\n"
-" localHistogram[hIdx] = sum;\n"
-" }\n"
-" GROUP_LDS_BARRIER;\n"
-"#if defined(USE_2LEVEL_REDUCE)\n"
-" if( lIdx < NUM_BUCKET )\n"
-" {\n"
-" localHistogram[hIdx] = localHistogram[hIdx-1];\n"
-" GROUP_MEM_FENCE;\n"
-" u32 u0, u1, u2;\n"
-" u0 = localHistogram[hIdx-3];\n"
-" u1 = localHistogram[hIdx-2];\n"
-" u2 = localHistogram[hIdx-1];\n"
-" AtomAdd( localHistogram[hIdx], u0 + u1 + u2 );\n"
-" GROUP_MEM_FENCE;\n"
-" u0 = localHistogram[hIdx-12];\n"
-" u1 = localHistogram[hIdx-8];\n"
-" u2 = localHistogram[hIdx-4];\n"
-" AtomAdd( localHistogram[hIdx], u0 + u1 + u2 );\n"
-" GROUP_MEM_FENCE;\n"
-" }\n"
-"#else\n"
-" if( lIdx < NUM_BUCKET )\n"
-" {\n"
-" localHistogram[hIdx] = localHistogram[hIdx-1];\n"
-" GROUP_MEM_FENCE;\n"
-" localHistogram[hIdx] += localHistogram[hIdx-1];\n"
-" GROUP_MEM_FENCE;\n"
-" localHistogram[hIdx] += localHistogram[hIdx-2];\n"
-" GROUP_MEM_FENCE;\n"
-" localHistogram[hIdx] += localHistogram[hIdx-4];\n"
-" GROUP_MEM_FENCE;\n"
-" localHistogram[hIdx] += localHistogram[hIdx-8];\n"
-" GROUP_MEM_FENCE;\n"
-" }\n"
-"#endif\n"
-" GROUP_LDS_BARRIER;\n"
-" }\n"
-" {\n"
-" for(int ie=0; ie<ELEMENTS_PER_WORK_ITEM; ie++)\n"
-" {\n"
-" int dataIdx = ELEMENTS_PER_WORK_ITEM*lIdx+ie;\n"
-" int binIdx = keys[ie];\n"
-" int groupOffset = localHistogramToCarry[binIdx];\n"
-" int myIdx = dataIdx - localHistogram[NUM_BUCKET+binIdx];\n"
-"#if defined(CHECK_BOUNDARY)\n"
-" if( addr+ie < n )\n"
-" {\n"
-" if ((groupOffset + myIdx)<n)\n"
-" {\n"
-" if (sortData[ie]==sortVal[ie])\n"
-" {\n"
-" \n"
-" SortDataCL tmp;\n"
-" tmp.m_key = sortData[ie];\n"
-" tmp.m_value = sortVal[ie];\n"
-" if (tmp.m_key == tmp.m_value)\n"
-" gDst[groupOffset + myIdx ] = tmp;\n"
-" }\n"
-" \n"
-" }\n"
-" }\n"
-"#else\n"
-" if ((groupOffset + myIdx)<n)\n"
-" {\n"
-" gDst[ groupOffset + myIdx ].m_key = sortData[ie];\n"
-" gDst[ groupOffset + myIdx ].m_value = sortVal[ie];\n"
-" }\n"
-"#endif\n"
-" }\n"
-" }\n"
-" GROUP_LDS_BARRIER;\n"
-" if( lIdx < NUM_BUCKET )\n"
-" {\n"
-" localHistogramToCarry[lIdx] += myHistogram;\n"
-" }\n"
-" GROUP_LDS_BARRIER;\n"
-" }\n"
-"}\n"
-"__kernel\n"
-"__attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
-"void SortAndScatterSortDataKernelSerial( __global const SortDataCL* restrict gSrc, __global const u32* rHistogram, __global SortDataCL* restrict gDst, int4 cb)\n"
-"{\n"
-" \n"
-" u32 gIdx = GET_GLOBAL_IDX;\n"
-" u32 realLocalIdx = GET_LOCAL_IDX;\n"
-" u32 wgIdx = GET_GROUP_IDX;\n"
-" u32 wgSize = GET_GROUP_SIZE;\n"
-" const int startBit = cb.m_startBit;\n"
-" const int n = cb.m_n;\n"
-" const int nWGs = cb.m_nWGs;\n"
-" const int nBlocksPerWG = cb.m_nBlocksPerWG;\n"
-" int counter[NUM_BUCKET];\n"
-" \n"
-" if (realLocalIdx>0)\n"
-" return;\n"
-" \n"
-" for (int c=0;c<NUM_BUCKET;c++)\n"
-" counter[c]=0;\n"
-" const int blockSize = ELEMENTS_PER_WORK_ITEM*WG_SIZE;\n"
-" \n"
-" int nBlocks = (n)/blockSize - nBlocksPerWG*wgIdx;\n"
-" for(int iblock=0; iblock<min(nBlocksPerWG, nBlocks); iblock++)\n"
-" {\n"
-" for (int lIdx=0;lIdx<WG_SIZE;lIdx++)\n"
-" {\n"
-" int addr2 = iblock*blockSize + blockSize*nBlocksPerWG*wgIdx + ELEMENTS_PER_WORK_ITEM*lIdx;\n"
-" \n"
-" for(int j=0; j<ELEMENTS_PER_WORK_ITEM; j++)\n"
-" {\n"
-" int i = addr2+j;\n"
-" if( i < n )\n"
-" {\n"
-" int tableIdx;\n"
-" tableIdx = (gSrc[i].m_key>>startBit) & 0xf;//0xf = NUM_TABLES-1\n"
-" gDst[rHistogram[tableIdx*nWGs+wgIdx] + counter[tableIdx]] = gSrc[i];\n"
-" counter[tableIdx] ++;\n"
-" }\n"
-" }\n"
-" }\n"
-" }\n"
-" \n"
-"}\n"
-"__kernel\n"
-"__attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
-"void SortAndScatterKernelSerial( __global const u32* restrict gSrc, __global const u32* rHistogram, __global u32* restrict gDst, int4 cb )\n"
-"{\n"
-" \n"
-" u32 gIdx = GET_GLOBAL_IDX;\n"
-" u32 realLocalIdx = GET_LOCAL_IDX;\n"
-" u32 wgIdx = GET_GROUP_IDX;\n"
-" u32 wgSize = GET_GROUP_SIZE;\n"
-" const int startBit = cb.m_startBit;\n"
-" const int n = cb.m_n;\n"
-" const int nWGs = cb.m_nWGs;\n"
-" const int nBlocksPerWG = cb.m_nBlocksPerWG;\n"
-" int counter[NUM_BUCKET];\n"
-" \n"
-" if (realLocalIdx>0)\n"
-" return;\n"
-" \n"
-" for (int c=0;c<NUM_BUCKET;c++)\n"
-" counter[c]=0;\n"
-" const int blockSize = ELEMENTS_PER_WORK_ITEM*WG_SIZE;\n"
-" \n"
-" int nBlocks = (n)/blockSize - nBlocksPerWG*wgIdx;\n"
-" for(int iblock=0; iblock<min(nBlocksPerWG, nBlocks); iblock++)\n"
-" {\n"
-" for (int lIdx=0;lIdx<WG_SIZE;lIdx++)\n"
-" {\n"
-" int addr2 = iblock*blockSize + blockSize*nBlocksPerWG*wgIdx + ELEMENTS_PER_WORK_ITEM*lIdx;\n"
-" \n"
-" for(int j=0; j<ELEMENTS_PER_WORK_ITEM; j++)\n"
-" {\n"
-" int i = addr2+j;\n"
-" if( i < n )\n"
-" {\n"
-" int tableIdx;\n"
-" tableIdx = (gSrc[i]>>startBit) & 0xf;//0xf = NUM_TABLES-1\n"
-" gDst[rHistogram[tableIdx*nWGs+wgIdx] + counter[tableIdx]] = gSrc[i];\n"
-" counter[tableIdx] ++;\n"
-" }\n"
-" }\n"
-" }\n"
-" }\n"
-" \n"
-"}\n"
-;
+static const char* radixSort32KernelsCL =
+ "/*\n"
+ "Bullet Continuous Collision Detection and Physics Library\n"
+ "Copyright (c) 2011 Advanced Micro Devices, Inc. http://bulletphysics.org\n"
+ "This software is provided 'as-is', without any express or implied warranty.\n"
+ "In no event will the authors be held liable for any damages arising from the use of this software.\n"
+ "Permission is granted to anyone to use this software for any purpose, \n"
+ "including commercial applications, and to alter it and redistribute it freely, \n"
+ "subject to the following restrictions:\n"
+ "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.\n"
+ "2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software.\n"
+ "3. This notice may not be removed or altered from any source distribution.\n"
+ "*/\n"
+ "//Author Takahiro Harada\n"
+ "//#pragma OPENCL EXTENSION cl_amd_printf : enable\n"
+ "#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable\n"
+ "#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable\n"
+ "typedef unsigned int u32;\n"
+ "#define GET_GROUP_IDX get_group_id(0)\n"
+ "#define GET_LOCAL_IDX get_local_id(0)\n"
+ "#define GET_GLOBAL_IDX get_global_id(0)\n"
+ "#define GET_GROUP_SIZE get_local_size(0)\n"
+ "#define GROUP_LDS_BARRIER barrier(CLK_LOCAL_MEM_FENCE)\n"
+ "#define GROUP_MEM_FENCE mem_fence(CLK_LOCAL_MEM_FENCE)\n"
+ "#define AtomInc(x) atom_inc(&(x))\n"
+ "#define AtomInc1(x, out) out = atom_inc(&(x))\n"
+ "#define AtomAdd(x, value) atom_add(&(x), value)\n"
+ "#define SELECT_UINT4( b, a, condition ) select( b,a,condition )\n"
+ "#define make_uint4 (uint4)\n"
+ "#define make_uint2 (uint2)\n"
+ "#define make_int2 (int2)\n"
+ "#define WG_SIZE 64\n"
+ "#define ELEMENTS_PER_WORK_ITEM (256/WG_SIZE)\n"
+ "#define BITS_PER_PASS 4\n"
+ "#define NUM_BUCKET (1<<BITS_PER_PASS)\n"
+ "typedef uchar u8;\n"
+ "// this isn't optimization for VLIW. But just reducing writes. \n"
+ "#define USE_2LEVEL_REDUCE 1\n"
+ "//#define CHECK_BOUNDARY 1\n"
+ "//#define NV_GPU 1\n"
+ "// Cypress\n"
+ "#define nPerWI 16\n"
+ "// Cayman\n"
+ "//#define nPerWI 20\n"
+ "#define m_n x\n"
+ "#define m_nWGs y\n"
+ "#define m_startBit z\n"
+ "#define m_nBlocksPerWG w\n"
+ "/*\n"
+ "typedef struct\n"
+ "{\n"
+ " int m_n;\n"
+ " int m_nWGs;\n"
+ " int m_startBit;\n"
+ " int m_nBlocksPerWG;\n"
+ "} ConstBuffer;\n"
+ "*/\n"
+ "typedef struct\n"
+ "{\n"
+ " unsigned int m_key;\n"
+ " unsigned int m_value;\n"
+ "} SortDataCL;\n"
+ "uint prefixScanVectorEx( uint4* data )\n"
+ "{\n"
+ " u32 sum = 0;\n"
+ " u32 tmp = data[0].x;\n"
+ " data[0].x = sum;\n"
+ " sum += tmp;\n"
+ " tmp = data[0].y;\n"
+ " data[0].y = sum;\n"
+ " sum += tmp;\n"
+ " tmp = data[0].z;\n"
+ " data[0].z = sum;\n"
+ " sum += tmp;\n"
+ " tmp = data[0].w;\n"
+ " data[0].w = sum;\n"
+ " sum += tmp;\n"
+ " return sum;\n"
+ "}\n"
+ "u32 localPrefixSum( u32 pData, uint lIdx, uint* totalSum, __local u32* sorterSharedMemory, int wgSize /*64 or 128*/ )\n"
+ "{\n"
+ " { // Set data\n"
+ " sorterSharedMemory[lIdx] = 0;\n"
+ " sorterSharedMemory[lIdx+wgSize] = pData;\n"
+ " }\n"
+ " GROUP_LDS_BARRIER;\n"
+ " { // Prefix sum\n"
+ " int idx = 2*lIdx + (wgSize+1);\n"
+ "#if defined(USE_2LEVEL_REDUCE)\n"
+ " if( lIdx < 64 )\n"
+ " {\n"
+ " u32 u0, u1, u2;\n"
+ " u0 = sorterSharedMemory[idx-3];\n"
+ " u1 = sorterSharedMemory[idx-2];\n"
+ " u2 = sorterSharedMemory[idx-1];\n"
+ " AtomAdd( sorterSharedMemory[idx], u0+u1+u2 ); \n"
+ " GROUP_MEM_FENCE;\n"
+ " u0 = sorterSharedMemory[idx-12];\n"
+ " u1 = sorterSharedMemory[idx-8];\n"
+ " u2 = sorterSharedMemory[idx-4];\n"
+ " AtomAdd( sorterSharedMemory[idx], u0+u1+u2 ); \n"
+ " GROUP_MEM_FENCE;\n"
+ " u0 = sorterSharedMemory[idx-48];\n"
+ " u1 = sorterSharedMemory[idx-32];\n"
+ " u2 = sorterSharedMemory[idx-16];\n"
+ " AtomAdd( sorterSharedMemory[idx], u0+u1+u2 ); \n"
+ " GROUP_MEM_FENCE;\n"
+ " if( wgSize > 64 )\n"
+ " {\n"
+ " sorterSharedMemory[idx] += sorterSharedMemory[idx-64];\n"
+ " GROUP_MEM_FENCE;\n"
+ " }\n"
+ " sorterSharedMemory[idx-1] += sorterSharedMemory[idx-2];\n"
+ " GROUP_MEM_FENCE;\n"
+ " }\n"
+ "#else\n"
+ " if( lIdx < 64 )\n"
+ " {\n"
+ " sorterSharedMemory[idx] += sorterSharedMemory[idx-1];\n"
+ " GROUP_MEM_FENCE;\n"
+ " sorterSharedMemory[idx] += sorterSharedMemory[idx-2]; \n"
+ " GROUP_MEM_FENCE;\n"
+ " sorterSharedMemory[idx] += sorterSharedMemory[idx-4];\n"
+ " GROUP_MEM_FENCE;\n"
+ " sorterSharedMemory[idx] += sorterSharedMemory[idx-8];\n"
+ " GROUP_MEM_FENCE;\n"
+ " sorterSharedMemory[idx] += sorterSharedMemory[idx-16];\n"
+ " GROUP_MEM_FENCE;\n"
+ " sorterSharedMemory[idx] += sorterSharedMemory[idx-32];\n"
+ " GROUP_MEM_FENCE;\n"
+ " if( wgSize > 64 )\n"
+ " {\n"
+ " sorterSharedMemory[idx] += sorterSharedMemory[idx-64];\n"
+ " GROUP_MEM_FENCE;\n"
+ " }\n"
+ " sorterSharedMemory[idx-1] += sorterSharedMemory[idx-2];\n"
+ " GROUP_MEM_FENCE;\n"
+ " }\n"
+ "#endif\n"
+ " }\n"
+ " GROUP_LDS_BARRIER;\n"
+ " *totalSum = sorterSharedMemory[wgSize*2-1];\n"
+ " u32 addValue = sorterSharedMemory[lIdx+wgSize-1];\n"
+ " return addValue;\n"
+ "}\n"
+ "//__attribute__((reqd_work_group_size(128,1,1)))\n"
+ "uint4 localPrefixSum128V( uint4 pData, uint lIdx, uint* totalSum, __local u32* sorterSharedMemory )\n"
+ "{\n"
+ " u32 s4 = prefixScanVectorEx( &pData );\n"
+ " u32 rank = localPrefixSum( s4, lIdx, totalSum, sorterSharedMemory, 128 );\n"
+ " return pData + make_uint4( rank, rank, rank, rank );\n"
+ "}\n"
+ "//__attribute__((reqd_work_group_size(64,1,1)))\n"
+ "uint4 localPrefixSum64V( uint4 pData, uint lIdx, uint* totalSum, __local u32* sorterSharedMemory )\n"
+ "{\n"
+ " u32 s4 = prefixScanVectorEx( &pData );\n"
+ " u32 rank = localPrefixSum( s4, lIdx, totalSum, sorterSharedMemory, 64 );\n"
+ " return pData + make_uint4( rank, rank, rank, rank );\n"
+ "}\n"
+ "u32 unpack4Key( u32 key, int keyIdx ){ return (key>>(keyIdx*8)) & 0xff;}\n"
+ "u32 bit8Scan(u32 v)\n"
+ "{\n"
+ " return (v<<8) + (v<<16) + (v<<24);\n"
+ "}\n"
+ "//===\n"
+ "#define MY_HISTOGRAM(idx) localHistogramMat[(idx)*WG_SIZE+lIdx]\n"
+ "__kernel\n"
+ "__attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
+ "void StreamCountKernel( __global u32* gSrc, __global u32* histogramOut, int4 cb )\n"
+ "{\n"
+ " __local u32 localHistogramMat[NUM_BUCKET*WG_SIZE];\n"
+ " u32 gIdx = GET_GLOBAL_IDX;\n"
+ " u32 lIdx = GET_LOCAL_IDX;\n"
+ " u32 wgIdx = GET_GROUP_IDX;\n"
+ " u32 wgSize = GET_GROUP_SIZE;\n"
+ " const int startBit = cb.m_startBit;\n"
+ " const int n = cb.m_n;\n"
+ " const int nWGs = cb.m_nWGs;\n"
+ " const int nBlocksPerWG = cb.m_nBlocksPerWG;\n"
+ " for(int i=0; i<NUM_BUCKET; i++)\n"
+ " {\n"
+ " MY_HISTOGRAM(i) = 0;\n"
+ " }\n"
+ " GROUP_LDS_BARRIER;\n"
+ " const int blockSize = ELEMENTS_PER_WORK_ITEM*WG_SIZE;\n"
+ " u32 localKey;\n"
+ " int nBlocks = (n)/blockSize - nBlocksPerWG*wgIdx;\n"
+ " int addr = blockSize*nBlocksPerWG*wgIdx + ELEMENTS_PER_WORK_ITEM*lIdx;\n"
+ " for(int iblock=0; iblock<min(nBlocksPerWG, nBlocks); iblock++, addr+=blockSize)\n"
+ " {\n"
+ " // MY_HISTOGRAM( localKeys.x ) ++ is much expensive than atomic add as it requires read and write while atomics can just add on AMD\n"
+ " // Using registers didn't perform well. It seems like use localKeys to address requires a lot of alu ops\n"
+ " // AMD: AtomInc performs better while NV prefers ++\n"
+ " for(int i=0; i<ELEMENTS_PER_WORK_ITEM; i++)\n"
+ " {\n"
+ "#if defined(CHECK_BOUNDARY)\n"
+ " if( addr+i < n )\n"
+ "#endif\n"
+ " {\n"
+ " localKey = (gSrc[addr+i]>>startBit) & 0xf;\n"
+ "#if defined(NV_GPU)\n"
+ " MY_HISTOGRAM( localKey )++;\n"
+ "#else\n"
+ " AtomInc( MY_HISTOGRAM( localKey ) );\n"
+ "#endif\n"
+ " }\n"
+ " }\n"
+ " }\n"
+ " GROUP_LDS_BARRIER;\n"
+ " \n"
+ " if( lIdx < NUM_BUCKET )\n"
+ " {\n"
+ " u32 sum = 0;\n"
+ " for(int i=0; i<GET_GROUP_SIZE; i++)\n"
+ " {\n"
+ " sum += localHistogramMat[lIdx*WG_SIZE+(i+lIdx)%GET_GROUP_SIZE];\n"
+ " }\n"
+ " histogramOut[lIdx*nWGs+wgIdx] = sum;\n"
+ " }\n"
+ "}\n"
+ "__kernel\n"
+ "__attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
+ "void StreamCountSortDataKernel( __global SortDataCL* gSrc, __global u32* histogramOut, int4 cb )\n"
+ "{\n"
+ " __local u32 localHistogramMat[NUM_BUCKET*WG_SIZE];\n"
+ " u32 gIdx = GET_GLOBAL_IDX;\n"
+ " u32 lIdx = GET_LOCAL_IDX;\n"
+ " u32 wgIdx = GET_GROUP_IDX;\n"
+ " u32 wgSize = GET_GROUP_SIZE;\n"
+ " const int startBit = cb.m_startBit;\n"
+ " const int n = cb.m_n;\n"
+ " const int nWGs = cb.m_nWGs;\n"
+ " const int nBlocksPerWG = cb.m_nBlocksPerWG;\n"
+ " for(int i=0; i<NUM_BUCKET; i++)\n"
+ " {\n"
+ " MY_HISTOGRAM(i) = 0;\n"
+ " }\n"
+ " GROUP_LDS_BARRIER;\n"
+ " const int blockSize = ELEMENTS_PER_WORK_ITEM*WG_SIZE;\n"
+ " u32 localKey;\n"
+ " int nBlocks = (n)/blockSize - nBlocksPerWG*wgIdx;\n"
+ " int addr = blockSize*nBlocksPerWG*wgIdx + ELEMENTS_PER_WORK_ITEM*lIdx;\n"
+ " for(int iblock=0; iblock<min(nBlocksPerWG, nBlocks); iblock++, addr+=blockSize)\n"
+ " {\n"
+ " // MY_HISTOGRAM( localKeys.x ) ++ is much expensive than atomic add as it requires read and write while atomics can just add on AMD\n"
+ " // Using registers didn't perform well. It seems like use localKeys to address requires a lot of alu ops\n"
+ " // AMD: AtomInc performs better while NV prefers ++\n"
+ " for(int i=0; i<ELEMENTS_PER_WORK_ITEM; i++)\n"
+ " {\n"
+ "#if defined(CHECK_BOUNDARY)\n"
+ " if( addr+i < n )\n"
+ "#endif\n"
+ " {\n"
+ " localKey = (gSrc[addr+i].m_key>>startBit) & 0xf;\n"
+ "#if defined(NV_GPU)\n"
+ " MY_HISTOGRAM( localKey )++;\n"
+ "#else\n"
+ " AtomInc( MY_HISTOGRAM( localKey ) );\n"
+ "#endif\n"
+ " }\n"
+ " }\n"
+ " }\n"
+ " GROUP_LDS_BARRIER;\n"
+ " \n"
+ " if( lIdx < NUM_BUCKET )\n"
+ " {\n"
+ " u32 sum = 0;\n"
+ " for(int i=0; i<GET_GROUP_SIZE; i++)\n"
+ " {\n"
+ " sum += localHistogramMat[lIdx*WG_SIZE+(i+lIdx)%GET_GROUP_SIZE];\n"
+ " }\n"
+ " histogramOut[lIdx*nWGs+wgIdx] = sum;\n"
+ " }\n"
+ "}\n"
+ "#define nPerLane (nPerWI/4)\n"
+ "// NUM_BUCKET*nWGs < 128*nPerWI\n"
+ "__kernel\n"
+ "__attribute__((reqd_work_group_size(128,1,1)))\n"
+ "void PrefixScanKernel( __global u32* wHistogram1, int4 cb )\n"
+ "{\n"
+ " __local u32 ldsTopScanData[128*2];\n"
+ " u32 lIdx = GET_LOCAL_IDX;\n"
+ " u32 wgIdx = GET_GROUP_IDX;\n"
+ " const int nWGs = cb.m_nWGs;\n"
+ " u32 data[nPerWI];\n"
+ " for(int i=0; i<nPerWI; i++)\n"
+ " {\n"
+ " data[i] = 0;\n"
+ " if( (nPerWI*lIdx+i) < NUM_BUCKET*nWGs )\n"
+ " data[i] = wHistogram1[nPerWI*lIdx+i];\n"
+ " }\n"
+ " uint4 myData = make_uint4(0,0,0,0);\n"
+ " for(int i=0; i<nPerLane; i++)\n"
+ " {\n"
+ " myData.x += data[nPerLane*0+i];\n"
+ " myData.y += data[nPerLane*1+i];\n"
+ " myData.z += data[nPerLane*2+i];\n"
+ " myData.w += data[nPerLane*3+i];\n"
+ " }\n"
+ " uint totalSum;\n"
+ " uint4 scanned = localPrefixSum128V( myData, lIdx, &totalSum, ldsTopScanData );\n"
+ "// for(int j=0; j<4; j++) // somehow it introduces a lot of branches\n"
+ " { int j = 0;\n"
+ " u32 sum = 0;\n"
+ " for(int i=0; i<nPerLane; i++)\n"
+ " {\n"
+ " u32 tmp = data[nPerLane*j+i];\n"
+ " data[nPerLane*j+i] = sum;\n"
+ " sum += tmp;\n"
+ " }\n"
+ " }\n"
+ " { int j = 1;\n"
+ " u32 sum = 0;\n"
+ " for(int i=0; i<nPerLane; i++)\n"
+ " {\n"
+ " u32 tmp = data[nPerLane*j+i];\n"
+ " data[nPerLane*j+i] = sum;\n"
+ " sum += tmp;\n"
+ " }\n"
+ " }\n"
+ " { int j = 2;\n"
+ " u32 sum = 0;\n"
+ " for(int i=0; i<nPerLane; i++)\n"
+ " {\n"
+ " u32 tmp = data[nPerLane*j+i];\n"
+ " data[nPerLane*j+i] = sum;\n"
+ " sum += tmp;\n"
+ " }\n"
+ " }\n"
+ " { int j = 3;\n"
+ " u32 sum = 0;\n"
+ " for(int i=0; i<nPerLane; i++)\n"
+ " {\n"
+ " u32 tmp = data[nPerLane*j+i];\n"
+ " data[nPerLane*j+i] = sum;\n"
+ " sum += tmp;\n"
+ " }\n"
+ " }\n"
+ " for(int i=0; i<nPerLane; i++)\n"
+ " {\n"
+ " data[nPerLane*0+i] += scanned.x;\n"
+ " data[nPerLane*1+i] += scanned.y;\n"
+ " data[nPerLane*2+i] += scanned.z;\n"
+ " data[nPerLane*3+i] += scanned.w;\n"
+ " }\n"
+ " for(int i=0; i<nPerWI; i++)\n"
+ " {\n"
+ " int index = nPerWI*lIdx+i;\n"
+ " if (index < NUM_BUCKET*nWGs)\n"
+ " wHistogram1[nPerWI*lIdx+i] = data[i];\n"
+ " }\n"
+ "}\n"
+ "// 4 scan, 4 exchange\n"
+ "void sort4Bits(u32 sortData[4], int startBit, int lIdx, __local u32* ldsSortData)\n"
+ "{\n"
+ " for(int bitIdx=0; bitIdx<BITS_PER_PASS; bitIdx++)\n"
+ " {\n"
+ " u32 mask = (1<<bitIdx);\n"
+ " uint4 cmpResult = make_uint4( (sortData[0]>>startBit) & mask, (sortData[1]>>startBit) & mask, (sortData[2]>>startBit) & mask, (sortData[3]>>startBit) & mask );\n"
+ " uint4 prefixSum = SELECT_UINT4( make_uint4(1,1,1,1), make_uint4(0,0,0,0), cmpResult != make_uint4(0,0,0,0) );\n"
+ " u32 total;\n"
+ " prefixSum = localPrefixSum64V( prefixSum, lIdx, &total, ldsSortData );\n"
+ " {\n"
+ " uint4 localAddr = make_uint4(lIdx*4+0,lIdx*4+1,lIdx*4+2,lIdx*4+3);\n"
+ " uint4 dstAddr = localAddr - prefixSum + make_uint4( total, total, total, total );\n"
+ " dstAddr = SELECT_UINT4( prefixSum, dstAddr, cmpResult != make_uint4(0, 0, 0, 0) );\n"
+ " GROUP_LDS_BARRIER;\n"
+ " ldsSortData[dstAddr.x] = sortData[0];\n"
+ " ldsSortData[dstAddr.y] = sortData[1];\n"
+ " ldsSortData[dstAddr.z] = sortData[2];\n"
+ " ldsSortData[dstAddr.w] = sortData[3];\n"
+ " GROUP_LDS_BARRIER;\n"
+ " sortData[0] = ldsSortData[localAddr.x];\n"
+ " sortData[1] = ldsSortData[localAddr.y];\n"
+ " sortData[2] = ldsSortData[localAddr.z];\n"
+ " sortData[3] = ldsSortData[localAddr.w];\n"
+ " GROUP_LDS_BARRIER;\n"
+ " }\n"
+ " }\n"
+ "}\n"
+ "// 2 scan, 2 exchange\n"
+ "void sort4Bits1(u32 sortData[4], int startBit, int lIdx, __local u32* ldsSortData)\n"
+ "{\n"
+ " for(uint ibit=0; ibit<BITS_PER_PASS; ibit+=2)\n"
+ " {\n"
+ " uint4 b = make_uint4((sortData[0]>>(startBit+ibit)) & 0x3, \n"
+ " (sortData[1]>>(startBit+ibit)) & 0x3, \n"
+ " (sortData[2]>>(startBit+ibit)) & 0x3, \n"
+ " (sortData[3]>>(startBit+ibit)) & 0x3);\n"
+ " u32 key4;\n"
+ " u32 sKeyPacked[4] = { 0, 0, 0, 0 };\n"
+ " {\n"
+ " sKeyPacked[0] |= 1<<(8*b.x);\n"
+ " sKeyPacked[1] |= 1<<(8*b.y);\n"
+ " sKeyPacked[2] |= 1<<(8*b.z);\n"
+ " sKeyPacked[3] |= 1<<(8*b.w);\n"
+ " key4 = sKeyPacked[0] + sKeyPacked[1] + sKeyPacked[2] + sKeyPacked[3];\n"
+ " }\n"
+ " u32 rankPacked;\n"
+ " u32 sumPacked;\n"
+ " {\n"
+ " rankPacked = localPrefixSum( key4, lIdx, &sumPacked, ldsSortData, WG_SIZE );\n"
+ " }\n"
+ " GROUP_LDS_BARRIER;\n"
+ " u32 newOffset[4] = { 0,0,0,0 };\n"
+ " {\n"
+ " u32 sumScanned = bit8Scan( sumPacked );\n"
+ " u32 scannedKeys[4];\n"
+ " scannedKeys[0] = 1<<(8*b.x);\n"
+ " scannedKeys[1] = 1<<(8*b.y);\n"
+ " scannedKeys[2] = 1<<(8*b.z);\n"
+ " scannedKeys[3] = 1<<(8*b.w);\n"
+ " { // 4 scans at once\n"
+ " u32 sum4 = 0;\n"
+ " for(int ie=0; ie<4; ie++)\n"
+ " {\n"
+ " u32 tmp = scannedKeys[ie];\n"
+ " scannedKeys[ie] = sum4;\n"
+ " sum4 += tmp;\n"
+ " }\n"
+ " }\n"
+ " {\n"
+ " u32 sumPlusRank = sumScanned + rankPacked;\n"
+ " { u32 ie = b.x;\n"
+ " scannedKeys[0] += sumPlusRank;\n"
+ " newOffset[0] = unpack4Key( scannedKeys[0], ie );\n"
+ " }\n"
+ " { u32 ie = b.y;\n"
+ " scannedKeys[1] += sumPlusRank;\n"
+ " newOffset[1] = unpack4Key( scannedKeys[1], ie );\n"
+ " }\n"
+ " { u32 ie = b.z;\n"
+ " scannedKeys[2] += sumPlusRank;\n"
+ " newOffset[2] = unpack4Key( scannedKeys[2], ie );\n"
+ " }\n"
+ " { u32 ie = b.w;\n"
+ " scannedKeys[3] += sumPlusRank;\n"
+ " newOffset[3] = unpack4Key( scannedKeys[3], ie );\n"
+ " }\n"
+ " }\n"
+ " }\n"
+ " GROUP_LDS_BARRIER;\n"
+ " {\n"
+ " ldsSortData[newOffset[0]] = sortData[0];\n"
+ " ldsSortData[newOffset[1]] = sortData[1];\n"
+ " ldsSortData[newOffset[2]] = sortData[2];\n"
+ " ldsSortData[newOffset[3]] = sortData[3];\n"
+ " GROUP_LDS_BARRIER;\n"
+ " u32 dstAddr = 4*lIdx;\n"
+ " sortData[0] = ldsSortData[dstAddr+0];\n"
+ " sortData[1] = ldsSortData[dstAddr+1];\n"
+ " sortData[2] = ldsSortData[dstAddr+2];\n"
+ " sortData[3] = ldsSortData[dstAddr+3];\n"
+ " GROUP_LDS_BARRIER;\n"
+ " }\n"
+ " }\n"
+ "}\n"
+ "#define SET_HISTOGRAM(setIdx, key) ldsSortData[(setIdx)*NUM_BUCKET+key]\n"
+ "__kernel\n"
+ "__attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
+ "void SortAndScatterKernel( __global const u32* restrict gSrc, __global const u32* rHistogram, __global u32* restrict gDst, int4 cb )\n"
+ "{\n"
+ " __local u32 ldsSortData[WG_SIZE*ELEMENTS_PER_WORK_ITEM+16];\n"
+ " __local u32 localHistogramToCarry[NUM_BUCKET];\n"
+ " __local u32 localHistogram[NUM_BUCKET*2];\n"
+ " u32 gIdx = GET_GLOBAL_IDX;\n"
+ " u32 lIdx = GET_LOCAL_IDX;\n"
+ " u32 wgIdx = GET_GROUP_IDX;\n"
+ " u32 wgSize = GET_GROUP_SIZE;\n"
+ " const int n = cb.m_n;\n"
+ " const int nWGs = cb.m_nWGs;\n"
+ " const int startBit = cb.m_startBit;\n"
+ " const int nBlocksPerWG = cb.m_nBlocksPerWG;\n"
+ " if( lIdx < (NUM_BUCKET) )\n"
+ " {\n"
+ " localHistogramToCarry[lIdx] = rHistogram[lIdx*nWGs + wgIdx];\n"
+ " }\n"
+ " GROUP_LDS_BARRIER;\n"
+ " const int blockSize = ELEMENTS_PER_WORK_ITEM*WG_SIZE;\n"
+ " int nBlocks = n/blockSize - nBlocksPerWG*wgIdx;\n"
+ " int addr = blockSize*nBlocksPerWG*wgIdx + ELEMENTS_PER_WORK_ITEM*lIdx;\n"
+ " for(int iblock=0; iblock<min(nBlocksPerWG, nBlocks); iblock++, addr+=blockSize)\n"
+ " {\n"
+ " u32 myHistogram = 0;\n"
+ " u32 sortData[ELEMENTS_PER_WORK_ITEM];\n"
+ " for(int i=0; i<ELEMENTS_PER_WORK_ITEM; i++)\n"
+ "#if defined(CHECK_BOUNDARY)\n"
+ " sortData[i] = ( addr+i < n )? gSrc[ addr+i ] : 0xffffffff;\n"
+ "#else\n"
+ " sortData[i] = gSrc[ addr+i ];\n"
+ "#endif\n"
+ " sort4Bits(sortData, startBit, lIdx, ldsSortData);\n"
+ " u32 keys[ELEMENTS_PER_WORK_ITEM];\n"
+ " for(int i=0; i<ELEMENTS_PER_WORK_ITEM; i++)\n"
+ " keys[i] = (sortData[i]>>startBit) & 0xf;\n"
+ " { // create histogram\n"
+ " u32 setIdx = lIdx/16;\n"
+ " if( lIdx < NUM_BUCKET )\n"
+ " {\n"
+ " localHistogram[lIdx] = 0;\n"
+ " }\n"
+ " ldsSortData[lIdx] = 0;\n"
+ " GROUP_LDS_BARRIER;\n"
+ " for(int i=0; i<ELEMENTS_PER_WORK_ITEM; i++)\n"
+ "#if defined(CHECK_BOUNDARY)\n"
+ " if( addr+i < n )\n"
+ "#endif\n"
+ "#if defined(NV_GPU)\n"
+ " SET_HISTOGRAM( setIdx, keys[i] )++;\n"
+ "#else\n"
+ " AtomInc( SET_HISTOGRAM( setIdx, keys[i] ) );\n"
+ "#endif\n"
+ " \n"
+ " GROUP_LDS_BARRIER;\n"
+ " \n"
+ " uint hIdx = NUM_BUCKET+lIdx;\n"
+ " if( lIdx < NUM_BUCKET )\n"
+ " {\n"
+ " u32 sum = 0;\n"
+ " for(int i=0; i<WG_SIZE/16; i++)\n"
+ " {\n"
+ " sum += SET_HISTOGRAM( i, lIdx );\n"
+ " }\n"
+ " myHistogram = sum;\n"
+ " localHistogram[hIdx] = sum;\n"
+ " }\n"
+ " GROUP_LDS_BARRIER;\n"
+ "#if defined(USE_2LEVEL_REDUCE)\n"
+ " if( lIdx < NUM_BUCKET )\n"
+ " {\n"
+ " localHistogram[hIdx] = localHistogram[hIdx-1];\n"
+ " GROUP_MEM_FENCE;\n"
+ " u32 u0, u1, u2;\n"
+ " u0 = localHistogram[hIdx-3];\n"
+ " u1 = localHistogram[hIdx-2];\n"
+ " u2 = localHistogram[hIdx-1];\n"
+ " AtomAdd( localHistogram[hIdx], u0 + u1 + u2 );\n"
+ " GROUP_MEM_FENCE;\n"
+ " u0 = localHistogram[hIdx-12];\n"
+ " u1 = localHistogram[hIdx-8];\n"
+ " u2 = localHistogram[hIdx-4];\n"
+ " AtomAdd( localHistogram[hIdx], u0 + u1 + u2 );\n"
+ " GROUP_MEM_FENCE;\n"
+ " }\n"
+ "#else\n"
+ " if( lIdx < NUM_BUCKET )\n"
+ " {\n"
+ " localHistogram[hIdx] = localHistogram[hIdx-1];\n"
+ " GROUP_MEM_FENCE;\n"
+ " localHistogram[hIdx] += localHistogram[hIdx-1];\n"
+ " GROUP_MEM_FENCE;\n"
+ " localHistogram[hIdx] += localHistogram[hIdx-2];\n"
+ " GROUP_MEM_FENCE;\n"
+ " localHistogram[hIdx] += localHistogram[hIdx-4];\n"
+ " GROUP_MEM_FENCE;\n"
+ " localHistogram[hIdx] += localHistogram[hIdx-8];\n"
+ " GROUP_MEM_FENCE;\n"
+ " }\n"
+ "#endif\n"
+ " GROUP_LDS_BARRIER;\n"
+ " }\n"
+ " {\n"
+ " for(int ie=0; ie<ELEMENTS_PER_WORK_ITEM; ie++)\n"
+ " {\n"
+ " int dataIdx = ELEMENTS_PER_WORK_ITEM*lIdx+ie;\n"
+ " int binIdx = keys[ie];\n"
+ " int groupOffset = localHistogramToCarry[binIdx];\n"
+ " int myIdx = dataIdx - localHistogram[NUM_BUCKET+binIdx];\n"
+ "#if defined(CHECK_BOUNDARY)\n"
+ " if( addr+ie < n )\n"
+ "#endif\n"
+ " gDst[ groupOffset + myIdx ] = sortData[ie];\n"
+ " }\n"
+ " }\n"
+ " GROUP_LDS_BARRIER;\n"
+ " if( lIdx < NUM_BUCKET )\n"
+ " {\n"
+ " localHistogramToCarry[lIdx] += myHistogram;\n"
+ " }\n"
+ " GROUP_LDS_BARRIER;\n"
+ " }\n"
+ "}\n"
+ "// 2 scan, 2 exchange\n"
+ "void sort4Bits1KeyValue(u32 sortData[4], int sortVal[4], int startBit, int lIdx, __local u32* ldsSortData, __local int *ldsSortVal)\n"
+ "{\n"
+ " for(uint ibit=0; ibit<BITS_PER_PASS; ibit+=2)\n"
+ " {\n"
+ " uint4 b = make_uint4((sortData[0]>>(startBit+ibit)) & 0x3, \n"
+ " (sortData[1]>>(startBit+ibit)) & 0x3, \n"
+ " (sortData[2]>>(startBit+ibit)) & 0x3, \n"
+ " (sortData[3]>>(startBit+ibit)) & 0x3);\n"
+ " u32 key4;\n"
+ " u32 sKeyPacked[4] = { 0, 0, 0, 0 };\n"
+ " {\n"
+ " sKeyPacked[0] |= 1<<(8*b.x);\n"
+ " sKeyPacked[1] |= 1<<(8*b.y);\n"
+ " sKeyPacked[2] |= 1<<(8*b.z);\n"
+ " sKeyPacked[3] |= 1<<(8*b.w);\n"
+ " key4 = sKeyPacked[0] + sKeyPacked[1] + sKeyPacked[2] + sKeyPacked[3];\n"
+ " }\n"
+ " u32 rankPacked;\n"
+ " u32 sumPacked;\n"
+ " {\n"
+ " rankPacked = localPrefixSum( key4, lIdx, &sumPacked, ldsSortData, WG_SIZE );\n"
+ " }\n"
+ " GROUP_LDS_BARRIER;\n"
+ " u32 newOffset[4] = { 0,0,0,0 };\n"
+ " {\n"
+ " u32 sumScanned = bit8Scan( sumPacked );\n"
+ " u32 scannedKeys[4];\n"
+ " scannedKeys[0] = 1<<(8*b.x);\n"
+ " scannedKeys[1] = 1<<(8*b.y);\n"
+ " scannedKeys[2] = 1<<(8*b.z);\n"
+ " scannedKeys[3] = 1<<(8*b.w);\n"
+ " { // 4 scans at once\n"
+ " u32 sum4 = 0;\n"
+ " for(int ie=0; ie<4; ie++)\n"
+ " {\n"
+ " u32 tmp = scannedKeys[ie];\n"
+ " scannedKeys[ie] = sum4;\n"
+ " sum4 += tmp;\n"
+ " }\n"
+ " }\n"
+ " {\n"
+ " u32 sumPlusRank = sumScanned + rankPacked;\n"
+ " { u32 ie = b.x;\n"
+ " scannedKeys[0] += sumPlusRank;\n"
+ " newOffset[0] = unpack4Key( scannedKeys[0], ie );\n"
+ " }\n"
+ " { u32 ie = b.y;\n"
+ " scannedKeys[1] += sumPlusRank;\n"
+ " newOffset[1] = unpack4Key( scannedKeys[1], ie );\n"
+ " }\n"
+ " { u32 ie = b.z;\n"
+ " scannedKeys[2] += sumPlusRank;\n"
+ " newOffset[2] = unpack4Key( scannedKeys[2], ie );\n"
+ " }\n"
+ " { u32 ie = b.w;\n"
+ " scannedKeys[3] += sumPlusRank;\n"
+ " newOffset[3] = unpack4Key( scannedKeys[3], ie );\n"
+ " }\n"
+ " }\n"
+ " }\n"
+ " GROUP_LDS_BARRIER;\n"
+ " {\n"
+ " ldsSortData[newOffset[0]] = sortData[0];\n"
+ " ldsSortData[newOffset[1]] = sortData[1];\n"
+ " ldsSortData[newOffset[2]] = sortData[2];\n"
+ " ldsSortData[newOffset[3]] = sortData[3];\n"
+ " ldsSortVal[newOffset[0]] = sortVal[0];\n"
+ " ldsSortVal[newOffset[1]] = sortVal[1];\n"
+ " ldsSortVal[newOffset[2]] = sortVal[2];\n"
+ " ldsSortVal[newOffset[3]] = sortVal[3];\n"
+ " GROUP_LDS_BARRIER;\n"
+ " u32 dstAddr = 4*lIdx;\n"
+ " sortData[0] = ldsSortData[dstAddr+0];\n"
+ " sortData[1] = ldsSortData[dstAddr+1];\n"
+ " sortData[2] = ldsSortData[dstAddr+2];\n"
+ " sortData[3] = ldsSortData[dstAddr+3];\n"
+ " sortVal[0] = ldsSortVal[dstAddr+0];\n"
+ " sortVal[1] = ldsSortVal[dstAddr+1];\n"
+ " sortVal[2] = ldsSortVal[dstAddr+2];\n"
+ " sortVal[3] = ldsSortVal[dstAddr+3];\n"
+ " GROUP_LDS_BARRIER;\n"
+ " }\n"
+ " }\n"
+ "}\n"
+ "__kernel\n"
+ "__attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
+ "void SortAndScatterSortDataKernel( __global const SortDataCL* restrict gSrc, __global const u32* rHistogram, __global SortDataCL* restrict gDst, int4 cb)\n"
+ "{\n"
+ " __local int ldsSortData[WG_SIZE*ELEMENTS_PER_WORK_ITEM+16];\n"
+ " __local int ldsSortVal[WG_SIZE*ELEMENTS_PER_WORK_ITEM+16];\n"
+ " __local u32 localHistogramToCarry[NUM_BUCKET];\n"
+ " __local u32 localHistogram[NUM_BUCKET*2];\n"
+ " u32 gIdx = GET_GLOBAL_IDX;\n"
+ " u32 lIdx = GET_LOCAL_IDX;\n"
+ " u32 wgIdx = GET_GROUP_IDX;\n"
+ " u32 wgSize = GET_GROUP_SIZE;\n"
+ " const int n = cb.m_n;\n"
+ " const int nWGs = cb.m_nWGs;\n"
+ " const int startBit = cb.m_startBit;\n"
+ " const int nBlocksPerWG = cb.m_nBlocksPerWG;\n"
+ " if( lIdx < (NUM_BUCKET) )\n"
+ " {\n"
+ " localHistogramToCarry[lIdx] = rHistogram[lIdx*nWGs + wgIdx];\n"
+ " }\n"
+ " GROUP_LDS_BARRIER;\n"
+ " \n"
+ " const int blockSize = ELEMENTS_PER_WORK_ITEM*WG_SIZE;\n"
+ " int nBlocks = n/blockSize - nBlocksPerWG*wgIdx;\n"
+ " int addr = blockSize*nBlocksPerWG*wgIdx + ELEMENTS_PER_WORK_ITEM*lIdx;\n"
+ " for(int iblock=0; iblock<min(nBlocksPerWG, nBlocks); iblock++, addr+=blockSize)\n"
+ " {\n"
+ " u32 myHistogram = 0;\n"
+ " int sortData[ELEMENTS_PER_WORK_ITEM];\n"
+ " int sortVal[ELEMENTS_PER_WORK_ITEM];\n"
+ " for(int i=0; i<ELEMENTS_PER_WORK_ITEM; i++)\n"
+ "#if defined(CHECK_BOUNDARY)\n"
+ " {\n"
+ " sortData[i] = ( addr+i < n )? gSrc[ addr+i ].m_key : 0xffffffff;\n"
+ " sortVal[i] = ( addr+i < n )? gSrc[ addr+i ].m_value : 0xffffffff;\n"
+ " }\n"
+ "#else\n"
+ " {\n"
+ " sortData[i] = gSrc[ addr+i ].m_key;\n"
+ " sortVal[i] = gSrc[ addr+i ].m_value;\n"
+ " }\n"
+ "#endif\n"
+ " sort4Bits1KeyValue(sortData, sortVal, startBit, lIdx, ldsSortData, ldsSortVal);\n"
+ " u32 keys[ELEMENTS_PER_WORK_ITEM];\n"
+ " for(int i=0; i<ELEMENTS_PER_WORK_ITEM; i++)\n"
+ " keys[i] = (sortData[i]>>startBit) & 0xf;\n"
+ " { // create histogram\n"
+ " u32 setIdx = lIdx/16;\n"
+ " if( lIdx < NUM_BUCKET )\n"
+ " {\n"
+ " localHistogram[lIdx] = 0;\n"
+ " }\n"
+ " ldsSortData[lIdx] = 0;\n"
+ " GROUP_LDS_BARRIER;\n"
+ " for(int i=0; i<ELEMENTS_PER_WORK_ITEM; i++)\n"
+ "#if defined(CHECK_BOUNDARY)\n"
+ " if( addr+i < n )\n"
+ "#endif\n"
+ "#if defined(NV_GPU)\n"
+ " SET_HISTOGRAM( setIdx, keys[i] )++;\n"
+ "#else\n"
+ " AtomInc( SET_HISTOGRAM( setIdx, keys[i] ) );\n"
+ "#endif\n"
+ " \n"
+ " GROUP_LDS_BARRIER;\n"
+ " \n"
+ " uint hIdx = NUM_BUCKET+lIdx;\n"
+ " if( lIdx < NUM_BUCKET )\n"
+ " {\n"
+ " u32 sum = 0;\n"
+ " for(int i=0; i<WG_SIZE/16; i++)\n"
+ " {\n"
+ " sum += SET_HISTOGRAM( i, lIdx );\n"
+ " }\n"
+ " myHistogram = sum;\n"
+ " localHistogram[hIdx] = sum;\n"
+ " }\n"
+ " GROUP_LDS_BARRIER;\n"
+ "#if defined(USE_2LEVEL_REDUCE)\n"
+ " if( lIdx < NUM_BUCKET )\n"
+ " {\n"
+ " localHistogram[hIdx] = localHistogram[hIdx-1];\n"
+ " GROUP_MEM_FENCE;\n"
+ " u32 u0, u1, u2;\n"
+ " u0 = localHistogram[hIdx-3];\n"
+ " u1 = localHistogram[hIdx-2];\n"
+ " u2 = localHistogram[hIdx-1];\n"
+ " AtomAdd( localHistogram[hIdx], u0 + u1 + u2 );\n"
+ " GROUP_MEM_FENCE;\n"
+ " u0 = localHistogram[hIdx-12];\n"
+ " u1 = localHistogram[hIdx-8];\n"
+ " u2 = localHistogram[hIdx-4];\n"
+ " AtomAdd( localHistogram[hIdx], u0 + u1 + u2 );\n"
+ " GROUP_MEM_FENCE;\n"
+ " }\n"
+ "#else\n"
+ " if( lIdx < NUM_BUCKET )\n"
+ " {\n"
+ " localHistogram[hIdx] = localHistogram[hIdx-1];\n"
+ " GROUP_MEM_FENCE;\n"
+ " localHistogram[hIdx] += localHistogram[hIdx-1];\n"
+ " GROUP_MEM_FENCE;\n"
+ " localHistogram[hIdx] += localHistogram[hIdx-2];\n"
+ " GROUP_MEM_FENCE;\n"
+ " localHistogram[hIdx] += localHistogram[hIdx-4];\n"
+ " GROUP_MEM_FENCE;\n"
+ " localHistogram[hIdx] += localHistogram[hIdx-8];\n"
+ " GROUP_MEM_FENCE;\n"
+ " }\n"
+ "#endif\n"
+ " GROUP_LDS_BARRIER;\n"
+ " }\n"
+ " {\n"
+ " for(int ie=0; ie<ELEMENTS_PER_WORK_ITEM; ie++)\n"
+ " {\n"
+ " int dataIdx = ELEMENTS_PER_WORK_ITEM*lIdx+ie;\n"
+ " int binIdx = keys[ie];\n"
+ " int groupOffset = localHistogramToCarry[binIdx];\n"
+ " int myIdx = dataIdx - localHistogram[NUM_BUCKET+binIdx];\n"
+ "#if defined(CHECK_BOUNDARY)\n"
+ " if( addr+ie < n )\n"
+ " {\n"
+ " if ((groupOffset + myIdx)<n)\n"
+ " {\n"
+ " if (sortData[ie]==sortVal[ie])\n"
+ " {\n"
+ " \n"
+ " SortDataCL tmp;\n"
+ " tmp.m_key = sortData[ie];\n"
+ " tmp.m_value = sortVal[ie];\n"
+ " if (tmp.m_key == tmp.m_value)\n"
+ " gDst[groupOffset + myIdx ] = tmp;\n"
+ " }\n"
+ " \n"
+ " }\n"
+ " }\n"
+ "#else\n"
+ " if ((groupOffset + myIdx)<n)\n"
+ " {\n"
+ " gDst[ groupOffset + myIdx ].m_key = sortData[ie];\n"
+ " gDst[ groupOffset + myIdx ].m_value = sortVal[ie];\n"
+ " }\n"
+ "#endif\n"
+ " }\n"
+ " }\n"
+ " GROUP_LDS_BARRIER;\n"
+ " if( lIdx < NUM_BUCKET )\n"
+ " {\n"
+ " localHistogramToCarry[lIdx] += myHistogram;\n"
+ " }\n"
+ " GROUP_LDS_BARRIER;\n"
+ " }\n"
+ "}\n"
+ "__kernel\n"
+ "__attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
+ "void SortAndScatterSortDataKernelSerial( __global const SortDataCL* restrict gSrc, __global const u32* rHistogram, __global SortDataCL* restrict gDst, int4 cb)\n"
+ "{\n"
+ " \n"
+ " u32 gIdx = GET_GLOBAL_IDX;\n"
+ " u32 realLocalIdx = GET_LOCAL_IDX;\n"
+ " u32 wgIdx = GET_GROUP_IDX;\n"
+ " u32 wgSize = GET_GROUP_SIZE;\n"
+ " const int startBit = cb.m_startBit;\n"
+ " const int n = cb.m_n;\n"
+ " const int nWGs = cb.m_nWGs;\n"
+ " const int nBlocksPerWG = cb.m_nBlocksPerWG;\n"
+ " int counter[NUM_BUCKET];\n"
+ " \n"
+ " if (realLocalIdx>0)\n"
+ " return;\n"
+ " \n"
+ " for (int c=0;c<NUM_BUCKET;c++)\n"
+ " counter[c]=0;\n"
+ " const int blockSize = ELEMENTS_PER_WORK_ITEM*WG_SIZE;\n"
+ " \n"
+ " int nBlocks = (n)/blockSize - nBlocksPerWG*wgIdx;\n"
+ " for(int iblock=0; iblock<min(nBlocksPerWG, nBlocks); iblock++)\n"
+ " {\n"
+ " for (int lIdx=0;lIdx<WG_SIZE;lIdx++)\n"
+ " {\n"
+ " int addr2 = iblock*blockSize + blockSize*nBlocksPerWG*wgIdx + ELEMENTS_PER_WORK_ITEM*lIdx;\n"
+ " \n"
+ " for(int j=0; j<ELEMENTS_PER_WORK_ITEM; j++)\n"
+ " {\n"
+ " int i = addr2+j;\n"
+ " if( i < n )\n"
+ " {\n"
+ " int tableIdx;\n"
+ " tableIdx = (gSrc[i].m_key>>startBit) & 0xf;//0xf = NUM_TABLES-1\n"
+ " gDst[rHistogram[tableIdx*nWGs+wgIdx] + counter[tableIdx]] = gSrc[i];\n"
+ " counter[tableIdx] ++;\n"
+ " }\n"
+ " }\n"
+ " }\n"
+ " }\n"
+ " \n"
+ "}\n"
+ "__kernel\n"
+ "__attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
+ "void SortAndScatterKernelSerial( __global const u32* restrict gSrc, __global const u32* rHistogram, __global u32* restrict gDst, int4 cb )\n"
+ "{\n"
+ " \n"
+ " u32 gIdx = GET_GLOBAL_IDX;\n"
+ " u32 realLocalIdx = GET_LOCAL_IDX;\n"
+ " u32 wgIdx = GET_GROUP_IDX;\n"
+ " u32 wgSize = GET_GROUP_SIZE;\n"
+ " const int startBit = cb.m_startBit;\n"
+ " const int n = cb.m_n;\n"
+ " const int nWGs = cb.m_nWGs;\n"
+ " const int nBlocksPerWG = cb.m_nBlocksPerWG;\n"
+ " int counter[NUM_BUCKET];\n"
+ " \n"
+ " if (realLocalIdx>0)\n"
+ " return;\n"
+ " \n"
+ " for (int c=0;c<NUM_BUCKET;c++)\n"
+ " counter[c]=0;\n"
+ " const int blockSize = ELEMENTS_PER_WORK_ITEM*WG_SIZE;\n"
+ " \n"
+ " int nBlocks = (n)/blockSize - nBlocksPerWG*wgIdx;\n"
+ " for(int iblock=0; iblock<min(nBlocksPerWG, nBlocks); iblock++)\n"
+ " {\n"
+ " for (int lIdx=0;lIdx<WG_SIZE;lIdx++)\n"
+ " {\n"
+ " int addr2 = iblock*blockSize + blockSize*nBlocksPerWG*wgIdx + ELEMENTS_PER_WORK_ITEM*lIdx;\n"
+ " \n"
+ " for(int j=0; j<ELEMENTS_PER_WORK_ITEM; j++)\n"
+ " {\n"
+ " int i = addr2+j;\n"
+ " if( i < n )\n"
+ " {\n"
+ " int tableIdx;\n"
+ " tableIdx = (gSrc[i]>>startBit) & 0xf;//0xf = NUM_TABLES-1\n"
+ " gDst[rHistogram[tableIdx*nWGs+wgIdx] + counter[tableIdx]] = gSrc[i];\n"
+ " counter[tableIdx] ++;\n"
+ " }\n"
+ " }\n"
+ " }\n"
+ " }\n"
+ " \n"
+ "}\n";