diff options
Diffstat (limited to 'thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives')
26 files changed, 0 insertions, 5413 deletions
diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3BoundSearchCL.cpp b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3BoundSearchCL.cpp deleted file mode 100644 index c0e11bfb26..0000000000 --- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3BoundSearchCL.cpp +++ /dev/null @@ -1,203 +0,0 @@ -/* -Copyright (c) 2012 Advanced Micro Devices, Inc. - -This software is provided 'as-is', without any express or implied warranty. -In no event will the authors be held liable for any damages arising from the use of this software. -Permission is granted to anyone to use this software for any purpose, -including commercial applications, and to alter it and redistribute it freely, -subject to the following restrictions: - -1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required. -2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software. -3. This notice may not be removed or altered from any source distribution. -*/ -//Originally written by Takahiro Harada -//Host-code rewritten by Erwin Coumans - -#define BOUNDSEARCH_PATH "src/Bullet3OpenCL/ParallelPrimitives/kernels/BoundSearchKernels.cl" -#define KERNEL0 "SearchSortDataLowerKernel" -#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) -{ - 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); - b3Assert(boundSearchProg); - - 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); - b3Assert(m_upperSortDataKernel); - - m_subtractKernel = 0; - - if (maxSize) - { - 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); -} - -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) -{ - b3Int4 constBuffer; - constBuffer.x = nSrc; - constBuffer.y = nDst; - - 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); - } - else if (option == BOUND_UPPER) - { - 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); - - launcher.launch1D(nSrc, 64); - } - else if (option == COUNT) - { - 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); - - 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())}; - - 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); - } - } - else - { - b3Assert(0); - } -} - -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); - - b3SortData minData, zeroData, maxData; - minData.m_key = -1; - minData.m_value = -1; - zeroData.m_key = 0; - zeroData.m_value = 0; - maxData.m_key = nDst; - maxData.m_value = nDst; - - if (option == BOUND_LOWER) - { - for (int i = 0; i < nSrc; i++) - { - b3SortData& iData = (i == 0) ? minData : src[i - 1]; - b3SortData& jData = (i == nSrc) ? maxData : src[i]; - - if (iData.m_key != jData.m_key) - { - int k = jData.m_key; - { - dst[k] = i; - } - } - } - } - else if (option == BOUND_UPPER) - { - for (int i = 1; i < nSrc + 1; i++) - { - b3SortData& iData = src[i - 1]; - b3SortData& jData = (i == nSrc) ? maxData : src[i]; - - if (iData.m_key != jData.m_key) - { - int k = iData.m_key; - { - dst[k] = i; - } - } - } - } - else if (option == COUNT) - { - b3AlignedObjectArray<unsigned int> lower; - lower.resize(nDst); - b3AlignedObjectArray<unsigned int> upper; - upper.resize(nDst); - - 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); - - for (int i = 0; i < nDst; i++) - { - dst[i] = upper[i] - lower[i]; - } - } - else - { - b3Assert(0); - } -} diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3BoundSearchCL.h b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3BoundSearchCL.h deleted file mode 100644 index 0d633e3d23..0000000000 --- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3BoundSearchCL.h +++ /dev/null @@ -1,64 +0,0 @@ -/* -Copyright (c) 2012 Advanced Micro Devices, Inc. - -This software is provided 'as-is', without any express or implied warranty. -In no event will the authors be held liable for any damages arising from the use of this software. -Permission is granted to anyone to use this software for any purpose, -including commercial applications, and to alter it and redistribute it freely, -subject to the following restrictions: - -1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required. -2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software. -3. This notice may not be removed or altered from any source distribution. -*/ -//Originally written by Takahiro Harada - -#ifndef B3_BOUNDSEARCH_H -#define B3_BOUNDSEARCH_H - -#pragma once - -/*#include <Adl/Adl.h> -#include <AdlPrimitives/Math/Math.h> -#include <AdlPrimitives/Sort/SortData.h> -#include <AdlPrimitives/Fill/Fill.h> -*/ - -#include "b3OpenCLArray.h" -#include "b3FillCL.h" -#include "b3RadixSort32CL.h" //for b3SortData (perhaps move it?) -class b3BoundSearchCL -{ -public: - enum Option - { - BOUND_LOWER, - BOUND_UPPER, - COUNT, - }; - - 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; - - 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); - - 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 diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3BufferInfoCL.h b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3BufferInfoCL.h deleted file mode 100644 index 35fc467b20..0000000000 --- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3BufferInfoCL.h +++ /dev/null @@ -1,18 +0,0 @@ - -#ifndef B3_BUFFER_INFO_CL_H -#define B3_BUFFER_INFO_CL_H - -#include "b3OpenCLArray.h" - -struct b3BufferInfoCL -{ - //b3BufferInfoCL(){} - - // 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 diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3FillCL.cpp b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3FillCL.cpp deleted file mode 100644 index bd25bb2101..0000000000 --- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3FillCL.cpp +++ /dev/null @@ -1,119 +0,0 @@ -#include "b3FillCL.h" -#include "Bullet3OpenCL/Initialize/b3OpenCLUtils.h" -#include "b3BufferInfoCL.h" -#include "b3LauncherCL.h" - -#define FILL_CL_PROGRAM_PATH "src/Bullet3OpenCL/ParallelPrimitives/kernels/FillKernels.cl" - -#include "kernels/FillKernelsCL.h" - -b3FillCL::b3FillCL(cl_context ctx, cl_device_id device, cl_command_queue 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); - b3Assert(fillProg); - - 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); - b3Assert(m_fillIntKernel); - - 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); - b3Assert(m_fillKernelInt2); -} - -b3FillCL::~b3FillCL() -{ - clReleaseKernel(m_fillKernelInt2); - 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); - - { - b3LauncherCL launcher(m_commandQueue, m_fillFloatKernel, "m_fillFloatKernel"); - launcher.setBuffer(src.getBufferCL()); - launcher.setConst(n); - launcher.setConst(value); - launcher.setConst(offset); - - launcher.launch1D(n); - } -} - -void b3FillCL::execute(b3OpenCLArray<int>& src, const int value, int n, int offset) -{ - b3Assert(n > 0); - - { - b3LauncherCL launcher(m_commandQueue, m_fillIntKernel, "m_fillIntKernel"); - launcher.setBuffer(src.getBufferCL()); - 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); - - { - 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); - launcher.setConst(offset); - - launcher.launch1D(n); - } -} - -void b3FillCL::executeHost(b3AlignedObjectArray<b3Int2>& src, const b3Int2& value, int n, int offset) -{ - for (int i = 0; i < n; i++) - { - src[i + offset] = value; - } -} - -void b3FillCL::executeHost(b3AlignedObjectArray<int>& src, const int value, int n, int offset) -{ - for (int i = 0; i < n; i++) - { - src[i + offset] = value; - } -} - -void b3FillCL::execute(b3OpenCLArray<b3Int2>& src, const b3Int2& value, int n, int offset) -{ - b3Assert(n > 0); - - { - b3BufferInfoCL bInfo[] = {b3BufferInfoCL(src.getBufferCL())}; - - 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); - } -} diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3FillCL.h b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3FillCL.h deleted file mode 100644 index c92c3e5119..0000000000 --- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3FillCL.h +++ /dev/null @@ -1,52 +0,0 @@ -#ifndef B3_FILL_CL_H -#define B3_FILL_CL_H - -#include "b3OpenCLArray.h" -#include "Bullet3Common/b3Scalar.h" - -#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]; - }; - -protected: -public: - b3FillCL(cl_context ctx, cl_device_id device, cl_command_queue queue); - - virtual ~b3FillCL(); - - 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<float>& src, const float 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<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 diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3LauncherCL.cpp b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3LauncherCL.cpp deleted file mode 100644 index c97d02eb45..0000000000 --- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3LauncherCL.cpp +++ /dev/null @@ -1,296 +0,0 @@ -#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) -{ - if (gDebugLauncherCL) - { - static int counter = 0; - printf("[%d] Prepare to launch OpenCL kernel %s\n", counter++, name); - } - - m_serializationSizeInBytes = sizeof(int); -} - -b3LauncherCL::~b3LauncherCL() -{ - 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); - } -} - -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, - ¶m_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) -{ - for (int i = 0; i < n; i++) - { - if (m_enableSerialization) - { - b3KernelArgData kernelArg; - 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, - ¶m_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); - } -} - -struct b3KernelArgDataUnaligned -{ - 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]; - }; -}; -#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); - } - b3KernelArgData b; - 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); - - if (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]; - - 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); - return -2; - } - - { - int expected = argGold->m_isBuffer; - int found = m_kernelArguments[ii].m_isBuffer; - - if (expected != found) - { - printf("failed validation: argument %d isBuffer expected: %d, found %d\n", ii, expected, found); - return -3; - } - } - index += sizeof(b3KernelArgData); - - if (argGold->m_isBuffer) - { - 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++) - { - 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); - - 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); - return -4; - } - } - - index += argGold->m_argSizeInBytes; - } - else - { - //compare content - 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]; - if (expected != found) - { - printf("failed validation: argument %d const data at byte position %d expected: %d, found %d\n", - ii, b, expected, found); - return -5; - } - } - } - } - 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; -} - -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* 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* ptr = (int*)&buf[buffSize]; - - *ptr = num; - - 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 deleted file mode 100644 index 18e9c1db2b..0000000000 --- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3LauncherCL.h +++ /dev/null @@ -1,128 +0,0 @@ - -#ifndef B3_LAUNCHER_CL_H -#define B3_LAUNCHER_CL_H - -#include "b3BufferInfoCL.h" -#include "Bullet3Common/b3MinMax.h" -#include "b3OpenCLArray.h" -#include <stdio.h> - -#define B3_DEBUG_SERIALIZE_CL - -#ifdef _WIN32 -#pragma warning(disable : 4996) -#endif -#define B3_CL_MAX_ARG_SIZE 16 -B3_ATTRIBUTE_ALIGNED16(struct) -b3KernelArgData -{ - 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]; - }; -}; - -class b3LauncherCL -{ - cl_command_queue m_commandQueue; - cl_kernel m_kernel; - int m_idx; - - 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); - - inline int validateResults(unsigned char* goldBuffer, int goldBufferCapacity, cl_context ctx); - - int serializeArguments(unsigned char* destBuffer, int destBufferCapacity); - - int getNumArguments() const - { - return m_kernelArguments.size(); - } - - b3KernelArgData getArgument(int index) - { - return m_kernelArguments[index]; - } - - 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); - } - - 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) - { - printf("Error: OpenCL status = %d\n", status); - } - b3Assert(status == CL_SUCCESS); - } - - void enableSerialization(bool serialize) - { - m_enableSerialization = serialize; - } -}; - -#endif //B3_LAUNCHER_CL_H diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3OpenCLArray.h b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3OpenCLArray.h deleted file mode 100644 index e837cceb66..0000000000 --- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3OpenCLArray.h +++ /dev/null @@ -1,300 +0,0 @@ -#ifndef B3_OPENCL_ARRAY_H -#define B3_OPENCL_ARRAY_H - -#include "Bullet3Common/b3AlignedObjectArray.h" -#include "Bullet3OpenCL/Initialize/b3OpenCLInclude.h" - -template <typename T> -class b3OpenCLArray -{ - size_t m_size; - size_t m_capacity; - cl_mem m_clBuffer; - - cl_context m_clContext; - cl_command_queue m_commandQueue; - - bool m_ownsMemory; - - bool m_allowGrowingCapacity; - - void deallocate() - { - if (m_clBuffer && m_ownsMemory) - { - clReleaseMemObject(m_clBuffer); - } - m_clBuffer = 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); - } - -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) - { - if (initialCapacity) - { - reserve(initialCapacity); - } - m_allowGrowingCapacity = allowGrowingCapacity; - } - - ///this is an error-prone method with no error checking, be careful! - void setFromOpenCLBuffer(cl_mem buffer, size_t sizeInElements) - { - deallocate(); - m_ownsMemory = false; - m_allowGrowingCapacity = false; - m_clBuffer = buffer; - m_size = sizeInElements; - 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; - // } - - cl_mem getBufferCL() const - { - return m_clBuffer; - } - - virtual ~b3OpenCLArray() - { - deallocate(); - m_size = 0; - m_capacity = 0; - } - - B3_FORCE_INLINE bool push_back(const T& _Val, bool waitForCompletion = true) - { - bool result = true; - size_t sz = size(); - if (sz == capacity()) - { - result = reserve(allocSize(size())); - } - copyFromHostPointer(&_Val, 1, sz, waitForCompletion); - m_size++; - return result; - } - - B3_FORCE_INLINE T forcedAt(size_t n) const - { - b3Assert(n >= 0); - b3Assert(n < capacity()); - T elem; - copyToHostPointer(&elem, 1, n, true); - return elem; - } - - B3_FORCE_INLINE T at(size_t n) const - { - b3Assert(n >= 0); - b3Assert(n < size()); - T elem; - copyToHostPointer(&elem, 1, n, true); - return elem; - } - - B3_FORCE_INLINE bool resize(size_t newsize, bool copyOldContents = true) - { - bool result = true; - size_t curSize = size(); - - if (newsize < curSize) - { - //leave the OpenCL memory for now - } - else - { - if (newsize > size()) - { - result = reserve(newsize, copyOldContents); - } - - //leave new data uninitialized (init in debug mode?) - //for (size_t i=curSize;i<newsize;i++) ... - } - - if (result) - { - m_size = newsize; - } - else - { - m_size = 0; - } - return result; - } - - B3_FORCE_INLINE size_t size() const - { - return m_size; - } - - B3_FORCE_INLINE size_t capacity() const - { - return m_capacity; - } - - B3_FORCE_INLINE bool reserve(size_t _Count, bool copyOldContents = true) - { - bool result = true; - // determine new minimum length of allocated storage - if (capacity() < _Count) - { // not enough room, reallocate - - if (m_allowGrowingCapacity) - { - cl_int ciErrNum; - //create a new OpenCL buffer - size_t memSizeInBytes = sizeof(T) * _Count; - cl_mem buf = clCreateBuffer(m_clContext, CL_MEM_READ_WRITE, memSizeInBytes, NULL, &ciErrNum); - if (ciErrNum != CL_SUCCESS) - { - b3Error("OpenCL out-of-memory\n"); - _Count = 0; - result = false; - } -//#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++) - src[i] = 0xbb; - 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 - - if (result) - { - if (copyOldContents) - copyToCL(buf, size()); - } - - //deallocate the old buffer - deallocate(); - - m_clBuffer = buf; - - m_capacity = _Count; - } - else - { - //fail: assert and - b3Assert(0); - deallocate(); - result = false; - } - } - return result; - } - - void copyToCL(cl_mem destination, size_t numElements, size_t firstElem = 0, size_t dstOffsetInElems = 0) const - { - if (numElements <= 0) - return; - - b3Assert(m_clBuffer); - b3Assert(destination); - - //likely some error, destination is same as source - b3Assert(m_clBuffer != destination); - - b3Assert((firstElem + numElements) <= m_size); - - cl_int status = 0; - - b3Assert(numElements > 0); - b3Assert(numElements <= m_size); - - 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); - - b3Assert(status == CL_SUCCESS); - } - - void copyFromHost(const b3AlignedObjectArray<T>& srcArray, bool waitForCompletion = true) - { - size_t newSize = srcArray.size(); - - bool copyOldContents = false; - resize(newSize, copyOldContents); - if (newSize) - copyFromHostPointer(&srcArray[0], newSize, 0, waitForCompletion); - } - - void copyFromHostPointer(const T* src, size_t numElems, size_t destFirstElem = 0, bool waitForCompletion = true) - { - b3Assert(numElems + destFirstElem <= capacity()); - - 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); - if (waitForCompletion) - clFinish(m_commandQueue); - } - else - { - b3Error("copyFromHostPointer invalid range\n"); - } - } - - void copyToHost(b3AlignedObjectArray<T>& destArray, bool waitForCompletion = true) const - { - destArray.resize(this->size()); - if (size()) - copyToHostPointer(&destArray[0], size(), 0, waitForCompletion); - } - - void copyToHostPointer(T* destPtr, size_t numElem, size_t srcFirstElem = 0, bool waitForCompletion = true) const - { - b3Assert(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); - - if (waitForCompletion) - clFinish(m_commandQueue); - } - else - { - b3Error("copyToHostPointer invalid range\n"); - } - } - - void copyFromOpenCLArray(const b3OpenCLArray& src) - { - size_t newSize = src.size(); - resize(newSize); - if (size()) - { - src.copyToCL(m_clBuffer, size()); - } - } -}; - -#endif //B3_OPENCL_ARRAY_H diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3PrefixScanCL.cpp b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3PrefixScanCL.cpp deleted file mode 100644 index 822b511633..0000000000 --- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3PrefixScanCL.cpp +++ /dev/null @@ -1,120 +0,0 @@ -#include "b3PrefixScanCL.h" -#include "b3FillCL.h" -#define B3_PREFIXSCAN_PROG_PATH "src/Bullet3OpenCL/ParallelPrimitives/kernels/PrefixScanKernels.cl" - -#include "b3LauncherCL.h" -#include "Bullet3OpenCL/Initialize/b3OpenCLUtils.h" -#include "kernels/PrefixScanKernelsCL.h" - -b3PrefixScanCL::b3PrefixScanCL(cl_context ctx, cl_device_id device, cl_command_queue queue, int size) - : m_commandQueue(queue) -{ - const char* scanKernelSource = prefixScanKernelsCL; - cl_int pErrNum; - 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); - 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); -} - -b3PrefixScanCL::~b3PrefixScanCL() -{ - delete m_workBuffer; - clReleaseKernel(m_localScanKernel); - clReleaseKernel(m_blockSumKernel); - clReleaseKernel(m_propagationKernel); -} - -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; -} - -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)); - - dst.resize(src.size()); - m_workBuffer->resize(src.size()); - - b3Int4 constBuffer; - constBuffer.x = n; - constBuffer.y = 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())}; - - 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())}; - - 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) - { - 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) - { - clFinish(m_commandQueue); - 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++) - { - dst[i] = s; - s += src[i]; - } - } - /*else - { - for(int i=0; i<n; i++) - { - s += hSrc[i]; - hDst[i] = s; - } - } - */ - - if (sum) - { - *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 deleted file mode 100644 index 346efa0c73..0000000000 --- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3PrefixScanCL.h +++ /dev/null @@ -1,35 +0,0 @@ - -#ifndef B3_PREFIX_SCAN_CL_H -#define B3_PREFIX_SCAN_CL_H - -#include "b3OpenCLArray.h" -#include "b3BufferInfoCL.h" -#include "Bullet3Common/b3AlignedObjectArray.h" - -class b3PrefixScanCL -{ - enum - { - BLOCK_SIZE = 128 - }; - - // Option m_option; - - cl_command_queue m_commandQueue; - - cl_kernel m_localScanKernel; - cl_kernel m_blockSumKernel; - cl_kernel m_propagationKernel; - - b3OpenCLArray<unsigned int>* m_workBuffer; - -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); -}; - -#endif //B3_PREFIX_SCAN_CL_H diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3PrefixScanFloat4CL.cpp b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3PrefixScanFloat4CL.cpp deleted file mode 100644 index 1cac97c988..0000000000 --- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3PrefixScanFloat4CL.cpp +++ /dev/null @@ -1,120 +0,0 @@ -#include "b3PrefixScanFloat4CL.h" -#include "b3FillCL.h" -#define B3_PREFIXSCAN_FLOAT4_PROG_PATH "src/Bullet3OpenCL/ParallelPrimitives/kernels/PrefixScanFloat4Kernels.cl" - -#include "b3LauncherCL.h" -#include "Bullet3OpenCL/Initialize/b3OpenCLUtils.h" -#include "kernels/PrefixScanKernelsFloat4CL.h" - -b3PrefixScanFloat4CL::b3PrefixScanFloat4CL(cl_context ctx, cl_device_id device, cl_command_queue queue, int size) - : m_commandQueue(queue) -{ - const char* scanKernelSource = prefixScanKernelsFloat4CL; - cl_int pErrNum; - 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); - 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); -} - -b3PrefixScanFloat4CL::~b3PrefixScanFloat4CL() -{ - delete m_workBuffer; - clReleaseKernel(m_localScanKernel); - clReleaseKernel(m_blockSumKernel); - clReleaseKernel(m_propagationKernel); -} - -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; -} - -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)); - - dst.resize(src.size()); - m_workBuffer->resize(src.size()); - - b3Int4 constBuffer; - constBuffer.x = n; - constBuffer.y = 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())}; - - 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())}; - - 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) - { - 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) - { - clFinish(m_commandQueue); - 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); - //if( data->m_option == EXCLUSIVE ) - { - for (int i = 0; i < n; i++) - { - dst[i] = s; - s += src[i]; - } - } - /*else - { - for(int i=0; i<n; i++) - { - s += hSrc[i]; - hDst[i] = s; - } - } - */ - - if (sum) - { - *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 deleted file mode 100644 index 122b0bfd68..0000000000 --- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3PrefixScanFloat4CL.h +++ /dev/null @@ -1,36 +0,0 @@ - -#ifndef B3_PREFIX_SCAN_CL_H -#define B3_PREFIX_SCAN_CL_H - -#include "b3OpenCLArray.h" -#include "b3BufferInfoCL.h" -#include "Bullet3Common/b3AlignedObjectArray.h" -#include "Bullet3Common/b3Vector3.h" - -class b3PrefixScanFloat4CL -{ - enum - { - BLOCK_SIZE = 128 - }; - - // Option m_option; - - cl_command_queue m_commandQueue; - - cl_kernel m_localScanKernel; - cl_kernel m_blockSumKernel; - cl_kernel m_propagationKernel; - - b3OpenCLArray<b3Vector3>* m_workBuffer; - -public: - b3PrefixScanFloat4CL(cl_context ctx, cl_device_id device, cl_command_queue queue, int size = 0); - - virtual ~b3PrefixScanFloat4CL(); - - void execute(b3OpenCLArray<b3Vector3>& src, b3OpenCLArray<b3Vector3>& dst, int n, b3Vector3* sum = 0); - void executeHost(b3AlignedObjectArray<b3Vector3>& src, b3AlignedObjectArray<b3Vector3>& dst, int n, b3Vector3* sum); -}; - -#endif //B3_PREFIX_SCAN_CL_H diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3RadixSort32CL.cpp b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3RadixSort32CL.cpp deleted file mode 100644 index e86af6583f..0000000000 --- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3RadixSort32CL.cpp +++ /dev/null @@ -1,646 +0,0 @@ - -#include "b3RadixSort32CL.h" -#include "b3LauncherCL.h" -#include "Bullet3OpenCL/Initialize/b3OpenCLUtils.h" -#include "b3PrefixScanCL.h" -#include "b3FillCL.h" - -#define RADIXSORT32_PATH "src/Bullet3OpenCL/ParallelPrimitives/kernels/RadixSort32Kernels.cl" - -#include "kernels/RadixSort32KernelsCL.h" - -b3RadixSort32CL::b3RadixSort32CL(cl_context ctx, cl_device_id device, cl_command_queue queue, int initialCapacity) - : m_commandQueue(queue) -{ - b3OpenCLDeviceInfo info; - 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); - - if (initialCapacity > 0) - { - m_workBuffer1->resize(initialCapacity); - m_workBuffer3->resize(initialCapacity); - m_workBuffer3a->resize(initialCapacity); - m_workBuffer4->resize(initialCapacity); - m_workBuffer4a->resize(initialCapacity); - } - - 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); - - 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); - b3Assert(m_sortAndScatterSortDataKernel); - m_sortAndScatterKernel = b3OpenCLUtils::compileCLKernelFromString(ctx, device, kernelSource, "SortAndScatterKernelSerial", &pErrNum, sortProg, additionalMacros); - b3Assert(m_sortAndScatterKernel); - } - else - { - 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); - b3Assert(m_sortAndScatterKernel); - } - - m_prefixScanKernel = b3OpenCLUtils::compileCLKernelFromString(ctx, device, kernelSource, "PrefixScanKernel", &pErrNum, sortProg, additionalMacros); - b3Assert(m_prefixScanKernel); -} - -b3RadixSort32CL::~b3RadixSort32CL() -{ - delete m_scan; - delete m_fill; - delete m_workBuffer1; - delete m_workBuffer2; - delete m_workBuffer3; - delete m_workBuffer3a; - delete m_workBuffer4; - delete m_workBuffer4a; - - clReleaseKernel(m_streamCountSortDataKernel); - clReleaseKernel(m_streamCountKernel); - clReleaseKernel(m_sortAndScatterSortDataKernel); - clReleaseKernel(m_sortAndScatterKernel); - clReleaseKernel(m_prefixScanKernel); -} - -void b3RadixSort32CL::executeHost(b3AlignedObjectArray<b3SortData>& inout, int sortBits /* = 32 */) -{ - int n = inout.size(); - const int BITS_PER_PASS = 8; - const int NUM_TABLES = (1 << BITS_PER_PASS); - - int tables[NUM_TABLES]; - int counter[NUM_TABLES]; - - b3SortData* src = &inout[0]; - b3AlignedObjectArray<b3SortData> workbuffer; - workbuffer.resize(inout.size()); - b3SortData* dst = &workbuffer[0]; - - int count = 0; - for (int startBit = 0; startBit < sortBits; startBit += BITS_PER_PASS) - { - for (int i = 0; i < NUM_TABLES; i++) - { - tables[i] = 0; - } - - for (int i = 0; i < n; i++) - { - 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++) - { - if (tables[i] != 0) - { - printf("tables[%d]=%d]\n", i, tables[i]); - } - } -#endif //TEST \ - // prefix scan - int sum = 0; - for (int i = 0; i < NUM_TABLES; i++) - { - int iData = tables[i]; - tables[i] = sum; - sum += iData; - counter[i] = 0; - } - - // distribute - for (int i = 0; i < n; i++) - { - int tableIdx = (src[i].m_key >> startBit) & (NUM_TABLES - 1); - - dst[tables[tableIdx] + counter[tableIdx]] = src[i]; - counter[tableIdx]++; - } - - b3Swap(src, dst); - count++; - } - - if (count & 1) - { - b3Assert(0); //need to copy - } -} - -void b3RadixSort32CL::executeHost(b3OpenCLArray<b3SortData>& keyValuesInOut, int sortBits /* = 32 */) -{ - b3AlignedObjectArray<b3SortData> inout; - keyValuesInOut.copyToHost(inout); - - 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) -{ -} - -//#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 - - b3OpenCLArray<b3SortData>* src = 0; - - if (workingSize % dataAlignment) - { - workingSize += dataAlignment - (workingSize % dataAlignment); - m_workBuffer4->copyFromOpenCLArray(keyValuesInOut); - m_workBuffer4->resize(workingSize); - b3SortData fillValue; - fillValue.m_key = 0xffffffff; - fillValue.m_value = 0xffffffff; - -#define USE_BTFILL -#ifdef USE_BTFILL - 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++) - { - m_workBuffer4->copyFromHostPointer(&fillValue, 1, i); - } -#endif //USE_BTFILL - - src = m_workBuffer4; - } - else - { - src = &keyValuesInOut; - m_workBuffer4->resize(0); - } - - 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); - - 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); - 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 = 1; - nWGs = nBlocks; - } - } - - 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 - - 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); - } - -#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++) - { - if (testHist[i] != 0) - printf("testHist[%d]=%d\n", i, testHist[i]); - } -#endif //DEBUG_RADIXSORT - -//fast prefix scan is not working properly on Mac OSX yet -#ifdef __APPLE__ - bool fastScan = false; -#else - 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); - destHisto = srcHisto; - } - else - { - //unsigned int sum; //for debugging - 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++) - { - if (testHist[i] != 0) - printf("testHist[%d]=%d\n", i, testHist[i]); - } - - for (int i = 0; i < testHist.size(); i += NUM_WGS) - { - printf("testHist[%d]=%d\n", i / NUM_WGS, testHist[i]); - } - -#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); - } -#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]++; - } - -#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 - -#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++) - { - if (testHist[i] != 0) - printf("testHist[%d]=%d\n", i, testHist[i]); - } -#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++; - } - - if (count & 1) - { - b3Assert(0); //need to copy from workbuffer to keyValuesInOut - } - - if (m_workBuffer4->size()) - { - m_workBuffer4->resize(originalSize); - 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 -} - -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) - { - 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); - - src = m_workBuffer4a; - } - else - { - src = &keysInOut; - m_workBuffer4a->resize(0); - } - - 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); - - 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); - 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 = 1; - nWGs = nBlocks; - } - } - - 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); - } - -//fast prefix scan is not working properly on Mac OSX yet -#ifdef __APPLE__ - bool fastScan = false; -#else - 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); - destHisto = srcHisto; - } - else - { - //unsigned int sum; //for debugging - 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); - } - - b3Swap(src, dst); - b3Swap(srcHisto, destHisto); - - count++; - } - - if (count & 1) - { - b3Assert(0); //need to copy from workbuffer to keyValuesInOut - } - - if (m_workBuffer4a->size()) - { - m_workBuffer4a->resize(originalSize); - keysInOut.copyFromOpenCLArray(*m_workBuffer4a); - } -} diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3RadixSort32CL.h b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3RadixSort32CL.h deleted file mode 100644 index 69caf182d7..0000000000 --- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/b3RadixSort32CL.h +++ /dev/null @@ -1,84 +0,0 @@ - -#ifndef B3_RADIXSORT32_H -#define B3_RADIXSORT32_H - -#include "b3OpenCLArray.h" - -struct b3SortData -{ - union { - unsigned int m_key; - unsigned int x; - }; - - union { - unsigned int m_value; - unsigned int y; - }; -}; -#include "b3BufferInfoCL.h" - -class b3RadixSort32CL -{ - 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; - - cl_command_queue m_commandQueue; - - cl_kernel m_streamCountSortDataKernel; - cl_kernel m_streamCountKernel; - - cl_kernel m_prefixScanKernel; - cl_kernel m_sortAndScatterSortDataKernel; - cl_kernel m_sortAndScatterKernel; - - bool m_deviceCPU; - - class b3PrefixScanCL* m_scan; - class b3FillCL* m_fill; - -public: - struct b3ConstData - { - 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 - }; - -private: -public: - b3RadixSort32CL(cl_context ctx, cl_device_id device, cl_command_queue queue, int initialCapacity = 0); - - 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); - - ///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 diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/BoundSearchKernels.cl b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/BoundSearchKernels.cl deleted file mode 100644 index f3b4a1e8a7..0000000000 --- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/BoundSearchKernels.cl +++ /dev/null @@ -1,106 +0,0 @@ -/* -Copyright (c) 2012 Advanced Micro Devices, Inc. - -This software is provided 'as-is', without any express or implied warranty. -In no event will the authors be held liable for any damages arising from the use of this software. -Permission is granted to anyone to use this software for any purpose, -including commercial applications, and to alter it and redistribute it freely, -subject to the following restrictions: - -1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required. -2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software. -3. This notice may not be removed or altered from any source distribution. -*/ -//Originally written by Takahiro Harada - - -typedef unsigned int u32; -#define GET_GROUP_IDX get_group_id(0) -#define GET_LOCAL_IDX get_local_id(0) -#define GET_GLOBAL_IDX get_global_id(0) -#define GET_GROUP_SIZE get_local_size(0) -#define GROUP_LDS_BARRIER barrier(CLK_LOCAL_MEM_FENCE) - -typedef struct -{ - u32 m_key; - u32 m_value; -}SortData; - - - -typedef struct -{ - u32 m_nSrc; - u32 m_nDst; - u32 m_padding[2]; -} ConstBuffer; - - - -__attribute__((reqd_work_group_size(64,1,1))) -__kernel -void SearchSortDataLowerKernel(__global SortData* src, __global u32 *dst, - unsigned int nSrc, unsigned int nDst) -{ - int gIdx = GET_GLOBAL_IDX; - - if( gIdx < nSrc ) - { - SortData first; first.m_key = (u32)(-1); first.m_value = (u32)(-1); - SortData end; end.m_key = nDst; end.m_value = nDst; - - SortData iData = (gIdx==0)? first: src[gIdx-1]; - SortData jData = (gIdx==nSrc)? end: src[gIdx]; - - if( iData.m_key != jData.m_key ) - { -// for(u32 k=iData.m_key+1; k<=min(jData.m_key, nDst-1); k++) - u32 k = jData.m_key; - { - dst[k] = gIdx; - } - } - } -} - - -__attribute__((reqd_work_group_size(64,1,1))) -__kernel -void SearchSortDataUpperKernel(__global SortData* src, __global u32 *dst, - unsigned int nSrc, unsigned int nDst) -{ - int gIdx = GET_GLOBAL_IDX+1; - - if( gIdx < nSrc+1 ) - { - SortData first; first.m_key = 0; first.m_value = 0; - SortData end; end.m_key = nDst; end.m_value = nDst; - - SortData iData = src[gIdx-1]; - SortData jData = (gIdx==nSrc)? end: src[gIdx]; - - if( iData.m_key != jData.m_key ) - { - u32 k = iData.m_key; - { - dst[k] = gIdx; - } - } - } -} - -__attribute__((reqd_work_group_size(64,1,1))) -__kernel -void SubtractKernel(__global u32* A, __global u32 *B, __global u32 *C, - unsigned int nSrc, unsigned int nDst) -{ - int gIdx = GET_GLOBAL_IDX; - - - if( gIdx < nDst ) - { - C[gIdx] = A[gIdx] - B[gIdx]; - } -} - diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/BoundSearchKernelsCL.h b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/BoundSearchKernelsCL.h deleted file mode 100644 index 1758dd41e3..0000000000 --- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/BoundSearchKernelsCL.h +++ /dev/null @@ -1,86 +0,0 @@ -//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"; diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/CopyKernels.cl b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/CopyKernels.cl deleted file mode 100644 index 2eee5752ec..0000000000 --- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/CopyKernels.cl +++ /dev/null @@ -1,128 +0,0 @@ -/* -Copyright (c) 2012 Advanced Micro Devices, Inc. - -This software is provided 'as-is', without any express or implied warranty. -In no event will the authors be held liable for any damages arising from the use of this software. -Permission is granted to anyone to use this software for any purpose, -including commercial applications, and to alter it and redistribute it freely, -subject to the following restrictions: - -1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required. -2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software. -3. This notice may not be removed or altered from any source distribution. -*/ -//Originally written by Takahiro Harada - -#pragma OPENCL EXTENSION cl_amd_printf : enable -#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable - -typedef unsigned int u32; -#define GET_GROUP_IDX get_group_id(0) -#define GET_LOCAL_IDX get_local_id(0) -#define GET_GLOBAL_IDX get_global_id(0) -#define GET_GROUP_SIZE get_local_size(0) -#define GROUP_LDS_BARRIER barrier(CLK_LOCAL_MEM_FENCE) -#define GROUP_MEM_FENCE mem_fence(CLK_LOCAL_MEM_FENCE) -#define AtomInc(x) atom_inc(&(x)) -#define AtomInc1(x, out) out = atom_inc(&(x)) - -#define make_uint4 (uint4) -#define make_uint2 (uint2) -#define make_int2 (int2) - -typedef struct -{ - int m_n; - int m_padding[3]; -} ConstBuffer; - - - -__kernel -__attribute__((reqd_work_group_size(64,1,1))) -void Copy1F4Kernel(__global float4* dst, __global float4* src, - ConstBuffer cb) -{ - int gIdx = GET_GLOBAL_IDX; - - if( gIdx < cb.m_n ) - { - float4 a0 = src[gIdx]; - - dst[ gIdx ] = a0; - } -} - -__kernel -__attribute__((reqd_work_group_size(64,1,1))) -void Copy2F4Kernel(__global float4* dst, __global float4* src, - ConstBuffer cb) -{ - int gIdx = GET_GLOBAL_IDX; - - if( 2*gIdx <= cb.m_n ) - { - float4 a0 = src[gIdx*2+0]; - float4 a1 = src[gIdx*2+1]; - - dst[ gIdx*2+0 ] = a0; - dst[ gIdx*2+1 ] = a1; - } -} - -__kernel -__attribute__((reqd_work_group_size(64,1,1))) -void Copy4F4Kernel(__global float4* dst, __global float4* src, - ConstBuffer cb) -{ - int gIdx = GET_GLOBAL_IDX; - - if( 4*gIdx <= cb.m_n ) - { - int idx0 = gIdx*4+0; - int idx1 = gIdx*4+1; - int idx2 = gIdx*4+2; - int idx3 = gIdx*4+3; - - float4 a0 = src[idx0]; - float4 a1 = src[idx1]; - float4 a2 = src[idx2]; - float4 a3 = src[idx3]; - - dst[ idx0 ] = a0; - dst[ idx1 ] = a1; - dst[ idx2 ] = a2; - dst[ idx3 ] = a3; - } -} - -__kernel -__attribute__((reqd_work_group_size(64,1,1))) -void CopyF1Kernel(__global float* dstF1, __global float* srcF1, - ConstBuffer cb) -{ - int gIdx = GET_GLOBAL_IDX; - - if( gIdx < cb.m_n ) - { - float a0 = srcF1[gIdx]; - - dstF1[ gIdx ] = a0; - } -} - -__kernel -__attribute__((reqd_work_group_size(64,1,1))) -void CopyF2Kernel(__global float2* dstF2, __global float2* srcF2, - ConstBuffer cb) -{ - int gIdx = GET_GLOBAL_IDX; - - if( gIdx < cb.m_n ) - { - float2 a0 = srcF2[gIdx]; - - dstF2[ gIdx ] = a0; - } -} - diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/CopyKernelsCL.h b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/CopyKernelsCL.h deleted file mode 100644 index 33c9279462..0000000000 --- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/CopyKernelsCL.h +++ /dev/null @@ -1,131 +0,0 @@ -//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"; diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/FillKernels.cl b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/FillKernels.cl deleted file mode 100644 index 71c31075dd..0000000000 --- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/FillKernels.cl +++ /dev/null @@ -1,107 +0,0 @@ -/* -Copyright (c) 2012 Advanced Micro Devices, Inc. - -This software is provided 'as-is', without any express or implied warranty. -In no event will the authors be held liable for any damages arising from the use of this software. -Permission is granted to anyone to use this software for any purpose, -including commercial applications, and to alter it and redistribute it freely, -subject to the following restrictions: - -1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required. -2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software. -3. This notice may not be removed or altered from any source distribution. -*/ -//Originally written by Takahiro Harada - - -#pragma OPENCL EXTENSION cl_amd_printf : enable -#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable - -typedef unsigned int u32; -#define GET_GROUP_IDX get_group_id(0) -#define GET_LOCAL_IDX get_local_id(0) -#define GET_GLOBAL_IDX get_global_id(0) -#define GET_GROUP_SIZE get_local_size(0) -#define GROUP_LDS_BARRIER barrier(CLK_LOCAL_MEM_FENCE) -#define GROUP_MEM_FENCE mem_fence(CLK_LOCAL_MEM_FENCE) -#define AtomInc(x) atom_inc(&(x)) -#define AtomInc1(x, out) out = atom_inc(&(x)) - -#define make_uint4 (uint4) -#define make_uint2 (uint2) -#define make_int2 (int2) - -typedef struct -{ - union - { - int4 m_data; - uint4 m_unsignedData; - float m_floatData; - }; - int m_offset; - int m_n; - int m_padding[2]; -} ConstBuffer; - - -__kernel -__attribute__((reqd_work_group_size(64,1,1))) -void FillIntKernel(__global int* dstInt, int num_elements, int value, const int offset) -{ - int gIdx = GET_GLOBAL_IDX; - - if( gIdx < num_elements ) - { - dstInt[ offset+gIdx ] = value; - } -} - -__kernel -__attribute__((reqd_work_group_size(64,1,1))) -void FillFloatKernel(__global float* dstFloat, int num_elements, float value, const int offset) -{ - int gIdx = GET_GLOBAL_IDX; - - if( gIdx < num_elements ) - { - dstFloat[ offset+gIdx ] = value; - } -} - -__kernel -__attribute__((reqd_work_group_size(64,1,1))) -void FillUnsignedIntKernel(__global unsigned int* dstInt, const int num, const unsigned int value, const int offset) -{ - int gIdx = GET_GLOBAL_IDX; - - if( gIdx < num ) - { - dstInt[ offset+gIdx ] = value; - } -} - -__kernel -__attribute__((reqd_work_group_size(64,1,1))) -void FillInt2Kernel(__global int2* dstInt2, const int num, const int2 value, const int offset) -{ - int gIdx = GET_GLOBAL_IDX; - - if( gIdx < num ) - { - dstInt2[ gIdx + offset] = make_int2( value.x, value.y ); - } -} - -__kernel -__attribute__((reqd_work_group_size(64,1,1))) -void FillInt4Kernel(__global int4* dstInt4, const int num, const int4 value, const int offset) -{ - int gIdx = GET_GLOBAL_IDX; - - if( gIdx < num ) - { - dstInt4[ offset+gIdx ] = value; - } -} - diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/FillKernelsCL.h b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/FillKernelsCL.h deleted file mode 100644 index 983e652270..0000000000 --- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/FillKernelsCL.h +++ /dev/null @@ -1,90 +0,0 @@ -//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"; diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/PrefixScanFloat4Kernels.cl b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/PrefixScanFloat4Kernels.cl deleted file mode 100644 index c9da79854a..0000000000 --- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/PrefixScanFloat4Kernels.cl +++ /dev/null @@ -1,154 +0,0 @@ -/* -Copyright (c) 2012 Advanced Micro Devices, Inc. - -This software is provided 'as-is', without any express or implied warranty. -In no event will the authors be held liable for any damages arising from the use of this software. -Permission is granted to anyone to use this software for any purpose, -including commercial applications, and to alter it and redistribute it freely, -subject to the following restrictions: - -1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required. -2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software. -3. This notice may not be removed or altered from any source distribution. -*/ -//Originally written by Takahiro Harada - - -typedef unsigned int u32; -#define GET_GROUP_IDX get_group_id(0) -#define GET_LOCAL_IDX get_local_id(0) -#define GET_GLOBAL_IDX get_global_id(0) -#define GET_GROUP_SIZE get_local_size(0) -#define GROUP_LDS_BARRIER barrier(CLK_LOCAL_MEM_FENCE) - -// takahiro end -#define WG_SIZE 128 -#define m_numElems x -#define m_numBlocks y -#define m_numScanBlocks z - -/*typedef struct -{ - uint m_numElems; - uint m_numBlocks; - uint m_numScanBlocks; - uint m_padding[1]; -} ConstBuffer; -*/ - -float4 ScanExclusiveFloat4(__local float4* data, u32 n, int lIdx, int lSize) -{ - float4 blocksum; - int offset = 1; - for(int nActive=n>>1; nActive>0; nActive>>=1, offset<<=1) - { - GROUP_LDS_BARRIER; - for(int iIdx=lIdx; iIdx<nActive; iIdx+=lSize) - { - int ai = offset*(2*iIdx+1)-1; - int bi = offset*(2*iIdx+2)-1; - data[bi] += data[ai]; - } - } - - GROUP_LDS_BARRIER; - - if( lIdx == 0 ) - { - blocksum = data[ n-1 ]; - data[ n-1 ] = 0; - } - - GROUP_LDS_BARRIER; - - offset >>= 1; - for(int nActive=1; nActive<n; nActive<<=1, offset>>=1 ) - { - GROUP_LDS_BARRIER; - for( int iIdx = lIdx; iIdx<nActive; iIdx += lSize ) - { - int ai = offset*(2*iIdx+1)-1; - int bi = offset*(2*iIdx+2)-1; - float4 temp = data[ai]; - data[ai] = data[bi]; - data[bi] += temp; - } - } - GROUP_LDS_BARRIER; - - return blocksum; -} - -__attribute__((reqd_work_group_size(WG_SIZE,1,1))) -__kernel -void LocalScanKernel(__global float4* dst, __global float4* src, __global float4* sumBuffer, uint4 cb) -{ - __local float4 ldsData[WG_SIZE*2]; - - int gIdx = GET_GLOBAL_IDX; - int lIdx = GET_LOCAL_IDX; - - ldsData[2*lIdx] = ( 2*gIdx < cb.m_numElems )? src[2*gIdx]: 0; - ldsData[2*lIdx + 1] = ( 2*gIdx+1 < cb.m_numElems )? src[2*gIdx + 1]: 0; - - float4 sum = ScanExclusiveFloat4(ldsData, WG_SIZE*2, GET_LOCAL_IDX, GET_GROUP_SIZE); - - if( lIdx == 0 ) - sumBuffer[GET_GROUP_IDX] = sum; - - if( (2*gIdx) < cb.m_numElems ) - { - dst[2*gIdx] = ldsData[2*lIdx]; - } - if( (2*gIdx + 1) < cb.m_numElems ) - { - dst[2*gIdx + 1] = ldsData[2*lIdx + 1]; - } -} - -__attribute__((reqd_work_group_size(WG_SIZE,1,1))) -__kernel -void AddOffsetKernel(__global float4* dst, __global float4* blockSum, uint4 cb) -{ - const u32 blockSize = WG_SIZE*2; - - int myIdx = GET_GROUP_IDX+1; - int lIdx = GET_LOCAL_IDX; - - float4 iBlockSum = blockSum[myIdx]; - - int endValue = min((myIdx+1)*(blockSize), cb.m_numElems); - for(int i=myIdx*blockSize+lIdx; i<endValue; i+=GET_GROUP_SIZE) - { - dst[i] += iBlockSum; - } -} - -__attribute__((reqd_work_group_size(WG_SIZE,1,1))) -__kernel -void TopLevelScanKernel(__global float4* dst, uint4 cb) -{ - __local float4 ldsData[2048]; - int gIdx = GET_GLOBAL_IDX; - int lIdx = GET_LOCAL_IDX; - int lSize = GET_GROUP_SIZE; - - for(int i=lIdx; i<cb.m_numScanBlocks; i+=lSize ) - { - ldsData[i] = (i<cb.m_numBlocks)? dst[i]:0; - } - - GROUP_LDS_BARRIER; - - float4 sum = ScanExclusiveFloat4(ldsData, cb.m_numScanBlocks, GET_LOCAL_IDX, GET_GROUP_SIZE); - - for(int i=lIdx; i<cb.m_numBlocks; i+=lSize ) - { - dst[i] = ldsData[i]; - } - - if( gIdx == 0 ) - { - dst[cb.m_numBlocks] = sum; - } -} diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/PrefixScanKernels.cl b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/PrefixScanKernels.cl deleted file mode 100644 index 963cc1e48e..0000000000 --- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/PrefixScanKernels.cl +++ /dev/null @@ -1,154 +0,0 @@ -/* -Copyright (c) 2012 Advanced Micro Devices, Inc. - -This software is provided 'as-is', without any express or implied warranty. -In no event will the authors be held liable for any damages arising from the use of this software. -Permission is granted to anyone to use this software for any purpose, -including commercial applications, and to alter it and redistribute it freely, -subject to the following restrictions: - -1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required. -2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software. -3. This notice may not be removed or altered from any source distribution. -*/ -//Originally written by Takahiro Harada - - -typedef unsigned int u32; -#define GET_GROUP_IDX get_group_id(0) -#define GET_LOCAL_IDX get_local_id(0) -#define GET_GLOBAL_IDX get_global_id(0) -#define GET_GROUP_SIZE get_local_size(0) -#define GROUP_LDS_BARRIER barrier(CLK_LOCAL_MEM_FENCE) - -// takahiro end -#define WG_SIZE 128 -#define m_numElems x -#define m_numBlocks y -#define m_numScanBlocks z - -/*typedef struct -{ - uint m_numElems; - uint m_numBlocks; - uint m_numScanBlocks; - uint m_padding[1]; -} ConstBuffer; -*/ - -u32 ScanExclusive(__local u32* data, u32 n, int lIdx, int lSize) -{ - u32 blocksum; - int offset = 1; - for(int nActive=n>>1; nActive>0; nActive>>=1, offset<<=1) - { - GROUP_LDS_BARRIER; - for(int iIdx=lIdx; iIdx<nActive; iIdx+=lSize) - { - int ai = offset*(2*iIdx+1)-1; - int bi = offset*(2*iIdx+2)-1; - data[bi] += data[ai]; - } - } - - GROUP_LDS_BARRIER; - - if( lIdx == 0 ) - { - blocksum = data[ n-1 ]; - data[ n-1 ] = 0; - } - - GROUP_LDS_BARRIER; - - offset >>= 1; - for(int nActive=1; nActive<n; nActive<<=1, offset>>=1 ) - { - GROUP_LDS_BARRIER; - for( int iIdx = lIdx; iIdx<nActive; iIdx += lSize ) - { - int ai = offset*(2*iIdx+1)-1; - int bi = offset*(2*iIdx+2)-1; - u32 temp = data[ai]; - data[ai] = data[bi]; - data[bi] += temp; - } - } - GROUP_LDS_BARRIER; - - return blocksum; -} - -__attribute__((reqd_work_group_size(WG_SIZE,1,1))) -__kernel -void LocalScanKernel(__global u32* dst, __global u32 *src, __global u32 *sumBuffer, - uint4 cb) -{ - __local u32 ldsData[WG_SIZE*2]; - - int gIdx = GET_GLOBAL_IDX; - int lIdx = GET_LOCAL_IDX; - - ldsData[2*lIdx] = ( 2*gIdx < cb.m_numElems )? src[2*gIdx]: 0; - ldsData[2*lIdx + 1] = ( 2*gIdx+1 < cb.m_numElems )? src[2*gIdx + 1]: 0; - - u32 sum = ScanExclusive(ldsData, WG_SIZE*2, GET_LOCAL_IDX, GET_GROUP_SIZE); - - if( lIdx == 0 ) sumBuffer[GET_GROUP_IDX] = sum; - - if( (2*gIdx) < cb.m_numElems ) - { - dst[2*gIdx] = ldsData[2*lIdx]; - } - if( (2*gIdx + 1) < cb.m_numElems ) - { - dst[2*gIdx + 1] = ldsData[2*lIdx + 1]; - } -} - -__attribute__((reqd_work_group_size(WG_SIZE,1,1))) -__kernel -void AddOffsetKernel(__global u32 *dst, __global u32 *blockSum, uint4 cb) -{ - const u32 blockSize = WG_SIZE*2; - - int myIdx = GET_GROUP_IDX+1; - int lIdx = GET_LOCAL_IDX; - - u32 iBlockSum = blockSum[myIdx]; - - int endValue = min((myIdx+1)*(blockSize), cb.m_numElems); - for(int i=myIdx*blockSize+lIdx; i<endValue; i+=GET_GROUP_SIZE) - { - dst[i] += iBlockSum; - } -} - -__attribute__((reqd_work_group_size(WG_SIZE,1,1))) -__kernel -void TopLevelScanKernel(__global u32* dst, uint4 cb) -{ - __local u32 ldsData[2048]; - int gIdx = GET_GLOBAL_IDX; - int lIdx = GET_LOCAL_IDX; - int lSize = GET_GROUP_SIZE; - - for(int i=lIdx; i<cb.m_numScanBlocks; i+=lSize ) - { - ldsData[i] = (i<cb.m_numBlocks)? dst[i]:0; - } - - GROUP_LDS_BARRIER; - - u32 sum = ScanExclusive(ldsData, cb.m_numScanBlocks, GET_LOCAL_IDX, GET_GROUP_SIZE); - - for(int i=lIdx; i<cb.m_numBlocks; i+=lSize ) - { - dst[i] = ldsData[i]; - } - - if( gIdx == 0 ) - { - dst[cb.m_numBlocks] = sum; - } -} diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/PrefixScanKernelsCL.h b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/PrefixScanKernelsCL.h deleted file mode 100644 index fc5e7b865c..0000000000 --- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/PrefixScanKernelsCL.h +++ /dev/null @@ -1,128 +0,0 @@ -//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"; diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/PrefixScanKernelsFloat4CL.h b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/PrefixScanKernelsFloat4CL.h deleted file mode 100644 index 15d1bc5195..0000000000 --- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/PrefixScanKernelsFloat4CL.h +++ /dev/null @@ -1,128 +0,0 @@ -//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"; diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/RadixSort32Kernels.cl b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/RadixSort32Kernels.cl deleted file mode 100644 index 7402e2f3b3..0000000000 --- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/RadixSort32Kernels.cl +++ /dev/null @@ -1,1071 +0,0 @@ -/* -Bullet Continuous Collision Detection and Physics Library -Copyright (c) 2011 Advanced Micro Devices, Inc. http://bulletphysics.org - -This software is provided 'as-is', without any express or implied warranty. -In no event will the authors be held liable for any damages arising from the use of this software. -Permission is granted to anyone to use this software for any purpose, -including commercial applications, and to alter it and redistribute it freely, -subject to the following restrictions: - -1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required. -2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software. -3. This notice may not be removed or altered from any source distribution. -*/ -//Author Takahiro Harada - - -//#pragma OPENCL EXTENSION cl_amd_printf : enable -#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable -#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable - -typedef unsigned int u32; -#define GET_GROUP_IDX get_group_id(0) -#define GET_LOCAL_IDX get_local_id(0) -#define GET_GLOBAL_IDX get_global_id(0) -#define GET_GROUP_SIZE get_local_size(0) -#define GROUP_LDS_BARRIER barrier(CLK_LOCAL_MEM_FENCE) -#define GROUP_MEM_FENCE mem_fence(CLK_LOCAL_MEM_FENCE) -#define AtomInc(x) atom_inc(&(x)) -#define AtomInc1(x, out) out = atom_inc(&(x)) -#define AtomAdd(x, value) atom_add(&(x), value) - -#define SELECT_UINT4( b, a, condition ) select( b,a,condition ) - - -#define make_uint4 (uint4) -#define make_uint2 (uint2) -#define make_int2 (int2) - -#define WG_SIZE 64 -#define ELEMENTS_PER_WORK_ITEM (256/WG_SIZE) -#define BITS_PER_PASS 4 -#define NUM_BUCKET (1<<BITS_PER_PASS) -typedef uchar u8; - -// this isn't optimization for VLIW. But just reducing writes. -#define USE_2LEVEL_REDUCE 1 - -//#define CHECK_BOUNDARY 1 - -//#define NV_GPU 1 - - -// Cypress -#define nPerWI 16 -// Cayman -//#define nPerWI 20 - -#define m_n x -#define m_nWGs y -#define m_startBit z -#define m_nBlocksPerWG w - -/* -typedef struct -{ - int m_n; - int m_nWGs; - int m_startBit; - int m_nBlocksPerWG; -} ConstBuffer; -*/ - -typedef struct -{ - unsigned int m_key; - unsigned int m_value; -} SortDataCL; - - -uint prefixScanVectorEx( uint4* data ) -{ - u32 sum = 0; - u32 tmp = data[0].x; - data[0].x = sum; - sum += tmp; - tmp = data[0].y; - data[0].y = sum; - sum += tmp; - tmp = data[0].z; - data[0].z = sum; - sum += tmp; - tmp = data[0].w; - data[0].w = sum; - sum += tmp; - return sum; -} - -u32 localPrefixSum( u32 pData, uint lIdx, uint* totalSum, __local u32* sorterSharedMemory, int wgSize /*64 or 128*/ ) -{ - { // Set data - sorterSharedMemory[lIdx] = 0; - sorterSharedMemory[lIdx+wgSize] = pData; - } - - GROUP_LDS_BARRIER; - - { // Prefix sum - int idx = 2*lIdx + (wgSize+1); -#if defined(USE_2LEVEL_REDUCE) - if( lIdx < 64 ) - { - u32 u0, u1, u2; - u0 = sorterSharedMemory[idx-3]; - u1 = sorterSharedMemory[idx-2]; - u2 = sorterSharedMemory[idx-1]; - AtomAdd( sorterSharedMemory[idx], u0+u1+u2 ); - GROUP_MEM_FENCE; - - u0 = sorterSharedMemory[idx-12]; - u1 = sorterSharedMemory[idx-8]; - u2 = sorterSharedMemory[idx-4]; - AtomAdd( sorterSharedMemory[idx], u0+u1+u2 ); - GROUP_MEM_FENCE; - - u0 = sorterSharedMemory[idx-48]; - u1 = sorterSharedMemory[idx-32]; - u2 = sorterSharedMemory[idx-16]; - AtomAdd( sorterSharedMemory[idx], u0+u1+u2 ); - GROUP_MEM_FENCE; - if( wgSize > 64 ) - { - sorterSharedMemory[idx] += sorterSharedMemory[idx-64]; - GROUP_MEM_FENCE; - } - - sorterSharedMemory[idx-1] += sorterSharedMemory[idx-2]; - GROUP_MEM_FENCE; - } -#else - if( lIdx < 64 ) - { - sorterSharedMemory[idx] += sorterSharedMemory[idx-1]; - GROUP_MEM_FENCE; - sorterSharedMemory[idx] += sorterSharedMemory[idx-2]; - GROUP_MEM_FENCE; - sorterSharedMemory[idx] += sorterSharedMemory[idx-4]; - GROUP_MEM_FENCE; - sorterSharedMemory[idx] += sorterSharedMemory[idx-8]; - GROUP_MEM_FENCE; - sorterSharedMemory[idx] += sorterSharedMemory[idx-16]; - GROUP_MEM_FENCE; - sorterSharedMemory[idx] += sorterSharedMemory[idx-32]; - GROUP_MEM_FENCE; - if( wgSize > 64 ) - { - sorterSharedMemory[idx] += sorterSharedMemory[idx-64]; - GROUP_MEM_FENCE; - } - - sorterSharedMemory[idx-1] += sorterSharedMemory[idx-2]; - GROUP_MEM_FENCE; - } -#endif - } - - GROUP_LDS_BARRIER; - - *totalSum = sorterSharedMemory[wgSize*2-1]; - u32 addValue = sorterSharedMemory[lIdx+wgSize-1]; - return addValue; -} - -//__attribute__((reqd_work_group_size(128,1,1))) -uint4 localPrefixSum128V( uint4 pData, uint lIdx, uint* totalSum, __local u32* sorterSharedMemory ) -{ - u32 s4 = prefixScanVectorEx( &pData ); - u32 rank = localPrefixSum( s4, lIdx, totalSum, sorterSharedMemory, 128 ); - return pData + make_uint4( rank, rank, rank, rank ); -} - - -//__attribute__((reqd_work_group_size(64,1,1))) -uint4 localPrefixSum64V( uint4 pData, uint lIdx, uint* totalSum, __local u32* sorterSharedMemory ) -{ - u32 s4 = prefixScanVectorEx( &pData ); - u32 rank = localPrefixSum( s4, lIdx, totalSum, sorterSharedMemory, 64 ); - return pData + make_uint4( rank, rank, rank, rank ); -} - -u32 unpack4Key( u32 key, int keyIdx ){ return (key>>(keyIdx*8)) & 0xff;} - -u32 bit8Scan(u32 v) -{ - return (v<<8) + (v<<16) + (v<<24); -} - -//=== - - - - -#define MY_HISTOGRAM(idx) localHistogramMat[(idx)*WG_SIZE+lIdx] - - -__kernel -__attribute__((reqd_work_group_size(WG_SIZE,1,1))) -void StreamCountKernel( __global u32* gSrc, __global u32* histogramOut, int4 cb ) -{ - __local u32 localHistogramMat[NUM_BUCKET*WG_SIZE]; - - u32 gIdx = GET_GLOBAL_IDX; - u32 lIdx = GET_LOCAL_IDX; - u32 wgIdx = GET_GROUP_IDX; - u32 wgSize = GET_GROUP_SIZE; - const int startBit = cb.m_startBit; - const int n = cb.m_n; - const int nWGs = cb.m_nWGs; - const int nBlocksPerWG = cb.m_nBlocksPerWG; - - for(int i=0; i<NUM_BUCKET; i++) - { - MY_HISTOGRAM(i) = 0; - } - - GROUP_LDS_BARRIER; - - const int blockSize = ELEMENTS_PER_WORK_ITEM*WG_SIZE; - u32 localKey; - - int nBlocks = (n)/blockSize - nBlocksPerWG*wgIdx; - - int addr = blockSize*nBlocksPerWG*wgIdx + ELEMENTS_PER_WORK_ITEM*lIdx; - - for(int iblock=0; iblock<min(nBlocksPerWG, nBlocks); iblock++, addr+=blockSize) - { - // 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 i=0; i<ELEMENTS_PER_WORK_ITEM; i++) - { -#if defined(CHECK_BOUNDARY) - if( addr+i < n ) -#endif - { - localKey = (gSrc[addr+i]>>startBit) & 0xf; -#if defined(NV_GPU) - MY_HISTOGRAM( localKey )++; -#else - AtomInc( MY_HISTOGRAM( localKey ) ); -#endif - } - } - } - - GROUP_LDS_BARRIER; - - if( lIdx < NUM_BUCKET ) - { - u32 sum = 0; - for(int i=0; i<GET_GROUP_SIZE; i++) - { - sum += localHistogramMat[lIdx*WG_SIZE+(i+lIdx)%GET_GROUP_SIZE]; - } - histogramOut[lIdx*nWGs+wgIdx] = sum; - } -} - -__kernel -__attribute__((reqd_work_group_size(WG_SIZE,1,1))) -void StreamCountSortDataKernel( __global SortDataCL* gSrc, __global u32* histogramOut, int4 cb ) -{ - __local u32 localHistogramMat[NUM_BUCKET*WG_SIZE]; - - u32 gIdx = GET_GLOBAL_IDX; - u32 lIdx = GET_LOCAL_IDX; - u32 wgIdx = GET_GROUP_IDX; - u32 wgSize = GET_GROUP_SIZE; - const int startBit = cb.m_startBit; - const int n = cb.m_n; - const int nWGs = cb.m_nWGs; - const int nBlocksPerWG = cb.m_nBlocksPerWG; - - for(int i=0; i<NUM_BUCKET; i++) - { - MY_HISTOGRAM(i) = 0; - } - - GROUP_LDS_BARRIER; - - const int blockSize = ELEMENTS_PER_WORK_ITEM*WG_SIZE; - u32 localKey; - - int nBlocks = (n)/blockSize - nBlocksPerWG*wgIdx; - - int addr = blockSize*nBlocksPerWG*wgIdx + ELEMENTS_PER_WORK_ITEM*lIdx; - - for(int iblock=0; iblock<min(nBlocksPerWG, nBlocks); iblock++, addr+=blockSize) - { - // 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 i=0; i<ELEMENTS_PER_WORK_ITEM; i++) - { -#if defined(CHECK_BOUNDARY) - if( addr+i < n ) -#endif - { - localKey = (gSrc[addr+i].m_key>>startBit) & 0xf; -#if defined(NV_GPU) - MY_HISTOGRAM( localKey )++; -#else - AtomInc( MY_HISTOGRAM( localKey ) ); -#endif - } - } - } - - GROUP_LDS_BARRIER; - - if( lIdx < NUM_BUCKET ) - { - u32 sum = 0; - for(int i=0; i<GET_GROUP_SIZE; i++) - { - sum += localHistogramMat[lIdx*WG_SIZE+(i+lIdx)%GET_GROUP_SIZE]; - } - histogramOut[lIdx*nWGs+wgIdx] = sum; - } -} - -#define nPerLane (nPerWI/4) - -// NUM_BUCKET*nWGs < 128*nPerWI -__kernel -__attribute__((reqd_work_group_size(128,1,1))) -void PrefixScanKernel( __global u32* wHistogram1, int4 cb ) -{ - __local u32 ldsTopScanData[128*2]; - - u32 lIdx = GET_LOCAL_IDX; - u32 wgIdx = GET_GROUP_IDX; - const int nWGs = cb.m_nWGs; - - u32 data[nPerWI]; - for(int i=0; i<nPerWI; i++) - { - data[i] = 0; - if( (nPerWI*lIdx+i) < NUM_BUCKET*nWGs ) - data[i] = wHistogram1[nPerWI*lIdx+i]; - } - - uint4 myData = make_uint4(0,0,0,0); - - for(int i=0; i<nPerLane; i++) - { - myData.x += data[nPerLane*0+i]; - myData.y += data[nPerLane*1+i]; - myData.z += data[nPerLane*2+i]; - myData.w += data[nPerLane*3+i]; - } - - uint totalSum; - uint4 scanned = localPrefixSum128V( myData, lIdx, &totalSum, ldsTopScanData ); - -// for(int j=0; j<4; j++) // somehow it introduces a lot of branches - { int j = 0; - u32 sum = 0; - for(int i=0; i<nPerLane; i++) - { - u32 tmp = data[nPerLane*j+i]; - data[nPerLane*j+i] = sum; - sum += tmp; - } - } - { int j = 1; - u32 sum = 0; - for(int i=0; i<nPerLane; i++) - { - u32 tmp = data[nPerLane*j+i]; - data[nPerLane*j+i] = sum; - sum += tmp; - } - } - { int j = 2; - u32 sum = 0; - for(int i=0; i<nPerLane; i++) - { - u32 tmp = data[nPerLane*j+i]; - data[nPerLane*j+i] = sum; - sum += tmp; - } - } - { int j = 3; - u32 sum = 0; - for(int i=0; i<nPerLane; i++) - { - u32 tmp = data[nPerLane*j+i]; - data[nPerLane*j+i] = sum; - sum += tmp; - } - } - - for(int i=0; i<nPerLane; i++) - { - data[nPerLane*0+i] += scanned.x; - data[nPerLane*1+i] += scanned.y; - data[nPerLane*2+i] += scanned.z; - data[nPerLane*3+i] += scanned.w; - } - - for(int i=0; i<nPerWI; i++) - { - int index = nPerWI*lIdx+i; - if (index < NUM_BUCKET*nWGs) - wHistogram1[nPerWI*lIdx+i] = data[i]; - } -} - -// 4 scan, 4 exchange -void sort4Bits(u32 sortData[4], int startBit, int lIdx, __local u32* ldsSortData) -{ - for(int bitIdx=0; bitIdx<BITS_PER_PASS; bitIdx++) - { - u32 mask = (1<<bitIdx); - uint4 cmpResult = make_uint4( (sortData[0]>>startBit) & mask, (sortData[1]>>startBit) & mask, (sortData[2]>>startBit) & mask, (sortData[3]>>startBit) & mask ); - uint4 prefixSum = SELECT_UINT4( make_uint4(1,1,1,1), make_uint4(0,0,0,0), cmpResult != make_uint4(0,0,0,0) ); - u32 total; - prefixSum = localPrefixSum64V( prefixSum, lIdx, &total, ldsSortData ); - { - uint4 localAddr = make_uint4(lIdx*4+0,lIdx*4+1,lIdx*4+2,lIdx*4+3); - uint4 dstAddr = localAddr - prefixSum + make_uint4( total, total, total, total ); - dstAddr = SELECT_UINT4( prefixSum, dstAddr, cmpResult != make_uint4(0, 0, 0, 0) ); - - GROUP_LDS_BARRIER; - - ldsSortData[dstAddr.x] = sortData[0]; - ldsSortData[dstAddr.y] = sortData[1]; - ldsSortData[dstAddr.z] = sortData[2]; - ldsSortData[dstAddr.w] = sortData[3]; - - GROUP_LDS_BARRIER; - - sortData[0] = ldsSortData[localAddr.x]; - sortData[1] = ldsSortData[localAddr.y]; - sortData[2] = ldsSortData[localAddr.z]; - sortData[3] = ldsSortData[localAddr.w]; - - GROUP_LDS_BARRIER; - } - } -} - -// 2 scan, 2 exchange -void sort4Bits1(u32 sortData[4], int startBit, int lIdx, __local u32* ldsSortData) -{ - for(uint ibit=0; ibit<BITS_PER_PASS; ibit+=2) - { - uint4 b = make_uint4((sortData[0]>>(startBit+ibit)) & 0x3, - (sortData[1]>>(startBit+ibit)) & 0x3, - (sortData[2]>>(startBit+ibit)) & 0x3, - (sortData[3]>>(startBit+ibit)) & 0x3); - - u32 key4; - u32 sKeyPacked[4] = { 0, 0, 0, 0 }; - { - sKeyPacked[0] |= 1<<(8*b.x); - sKeyPacked[1] |= 1<<(8*b.y); - sKeyPacked[2] |= 1<<(8*b.z); - sKeyPacked[3] |= 1<<(8*b.w); - - key4 = sKeyPacked[0] + sKeyPacked[1] + sKeyPacked[2] + sKeyPacked[3]; - } - - u32 rankPacked; - u32 sumPacked; - { - rankPacked = localPrefixSum( key4, lIdx, &sumPacked, ldsSortData, WG_SIZE ); - } - - GROUP_LDS_BARRIER; - - u32 newOffset[4] = { 0,0,0,0 }; - { - u32 sumScanned = bit8Scan( sumPacked ); - - u32 scannedKeys[4]; - scannedKeys[0] = 1<<(8*b.x); - scannedKeys[1] = 1<<(8*b.y); - scannedKeys[2] = 1<<(8*b.z); - scannedKeys[3] = 1<<(8*b.w); - { // 4 scans at once - u32 sum4 = 0; - for(int ie=0; ie<4; ie++) - { - u32 tmp = scannedKeys[ie]; - scannedKeys[ie] = sum4; - sum4 += tmp; - } - } - - { - u32 sumPlusRank = sumScanned + rankPacked; - { u32 ie = b.x; - scannedKeys[0] += sumPlusRank; - newOffset[0] = unpack4Key( scannedKeys[0], ie ); - } - { u32 ie = b.y; - scannedKeys[1] += sumPlusRank; - newOffset[1] = unpack4Key( scannedKeys[1], ie ); - } - { u32 ie = b.z; - scannedKeys[2] += sumPlusRank; - newOffset[2] = unpack4Key( scannedKeys[2], ie ); - } - { u32 ie = b.w; - scannedKeys[3] += sumPlusRank; - newOffset[3] = unpack4Key( scannedKeys[3], ie ); - } - } - } - - - GROUP_LDS_BARRIER; - - { - ldsSortData[newOffset[0]] = sortData[0]; - ldsSortData[newOffset[1]] = sortData[1]; - ldsSortData[newOffset[2]] = sortData[2]; - ldsSortData[newOffset[3]] = sortData[3]; - - GROUP_LDS_BARRIER; - - u32 dstAddr = 4*lIdx; - sortData[0] = ldsSortData[dstAddr+0]; - sortData[1] = ldsSortData[dstAddr+1]; - sortData[2] = ldsSortData[dstAddr+2]; - sortData[3] = ldsSortData[dstAddr+3]; - - GROUP_LDS_BARRIER; - } - } -} - -#define SET_HISTOGRAM(setIdx, key) ldsSortData[(setIdx)*NUM_BUCKET+key] - -__kernel -__attribute__((reqd_work_group_size(WG_SIZE,1,1))) -void SortAndScatterKernel( __global const u32* restrict gSrc, __global const u32* rHistogram, __global u32* restrict gDst, int4 cb ) -{ - __local u32 ldsSortData[WG_SIZE*ELEMENTS_PER_WORK_ITEM+16]; - __local u32 localHistogramToCarry[NUM_BUCKET]; - __local u32 localHistogram[NUM_BUCKET*2]; - - u32 gIdx = GET_GLOBAL_IDX; - u32 lIdx = GET_LOCAL_IDX; - u32 wgIdx = GET_GROUP_IDX; - u32 wgSize = GET_GROUP_SIZE; - - const int n = cb.m_n; - const int nWGs = cb.m_nWGs; - const int startBit = cb.m_startBit; - const int nBlocksPerWG = cb.m_nBlocksPerWG; - - if( lIdx < (NUM_BUCKET) ) - { - localHistogramToCarry[lIdx] = rHistogram[lIdx*nWGs + wgIdx]; - } - - GROUP_LDS_BARRIER; - - const int blockSize = ELEMENTS_PER_WORK_ITEM*WG_SIZE; - - int nBlocks = n/blockSize - nBlocksPerWG*wgIdx; - - int addr = blockSize*nBlocksPerWG*wgIdx + ELEMENTS_PER_WORK_ITEM*lIdx; - - for(int iblock=0; iblock<min(nBlocksPerWG, nBlocks); iblock++, addr+=blockSize) - { - u32 myHistogram = 0; - - u32 sortData[ELEMENTS_PER_WORK_ITEM]; - for(int i=0; i<ELEMENTS_PER_WORK_ITEM; i++) -#if defined(CHECK_BOUNDARY) - sortData[i] = ( addr+i < n )? gSrc[ addr+i ] : 0xffffffff; -#else - sortData[i] = gSrc[ addr+i ]; -#endif - - sort4Bits(sortData, startBit, lIdx, ldsSortData); - - u32 keys[ELEMENTS_PER_WORK_ITEM]; - for(int i=0; i<ELEMENTS_PER_WORK_ITEM; i++) - keys[i] = (sortData[i]>>startBit) & 0xf; - - { // create histogram - u32 setIdx = lIdx/16; - if( lIdx < NUM_BUCKET ) - { - localHistogram[lIdx] = 0; - } - ldsSortData[lIdx] = 0; - GROUP_LDS_BARRIER; - - for(int i=0; i<ELEMENTS_PER_WORK_ITEM; i++) -#if defined(CHECK_BOUNDARY) - if( addr+i < n ) -#endif - -#if defined(NV_GPU) - SET_HISTOGRAM( setIdx, keys[i] )++; -#else - AtomInc( SET_HISTOGRAM( setIdx, keys[i] ) ); -#endif - - GROUP_LDS_BARRIER; - - uint hIdx = NUM_BUCKET+lIdx; - if( lIdx < NUM_BUCKET ) - { - u32 sum = 0; - for(int i=0; i<WG_SIZE/16; i++) - { - sum += SET_HISTOGRAM( i, lIdx ); - } - myHistogram = sum; - localHistogram[hIdx] = sum; - } - GROUP_LDS_BARRIER; - -#if defined(USE_2LEVEL_REDUCE) - if( lIdx < NUM_BUCKET ) - { - localHistogram[hIdx] = localHistogram[hIdx-1]; - GROUP_MEM_FENCE; - - u32 u0, u1, u2; - u0 = localHistogram[hIdx-3]; - u1 = localHistogram[hIdx-2]; - u2 = localHistogram[hIdx-1]; - AtomAdd( localHistogram[hIdx], u0 + u1 + u2 ); - GROUP_MEM_FENCE; - u0 = localHistogram[hIdx-12]; - u1 = localHistogram[hIdx-8]; - u2 = localHistogram[hIdx-4]; - AtomAdd( localHistogram[hIdx], u0 + u1 + u2 ); - GROUP_MEM_FENCE; - } -#else - if( lIdx < NUM_BUCKET ) - { - localHistogram[hIdx] = localHistogram[hIdx-1]; - GROUP_MEM_FENCE; - localHistogram[hIdx] += localHistogram[hIdx-1]; - GROUP_MEM_FENCE; - localHistogram[hIdx] += localHistogram[hIdx-2]; - GROUP_MEM_FENCE; - localHistogram[hIdx] += localHistogram[hIdx-4]; - GROUP_MEM_FENCE; - localHistogram[hIdx] += localHistogram[hIdx-8]; - GROUP_MEM_FENCE; - } -#endif - GROUP_LDS_BARRIER; - } - - { - for(int ie=0; ie<ELEMENTS_PER_WORK_ITEM; ie++) - { - int dataIdx = ELEMENTS_PER_WORK_ITEM*lIdx+ie; - int binIdx = keys[ie]; - int groupOffset = localHistogramToCarry[binIdx]; - int myIdx = dataIdx - localHistogram[NUM_BUCKET+binIdx]; -#if defined(CHECK_BOUNDARY) - if( addr+ie < n ) -#endif - gDst[ groupOffset + myIdx ] = sortData[ie]; - } - } - - GROUP_LDS_BARRIER; - - if( lIdx < NUM_BUCKET ) - { - localHistogramToCarry[lIdx] += myHistogram; - } - GROUP_LDS_BARRIER; - } -} - -// 2 scan, 2 exchange -void sort4Bits1KeyValue(u32 sortData[4], int sortVal[4], int startBit, int lIdx, __local u32* ldsSortData, __local int *ldsSortVal) -{ - for(uint ibit=0; ibit<BITS_PER_PASS; ibit+=2) - { - uint4 b = make_uint4((sortData[0]>>(startBit+ibit)) & 0x3, - (sortData[1]>>(startBit+ibit)) & 0x3, - (sortData[2]>>(startBit+ibit)) & 0x3, - (sortData[3]>>(startBit+ibit)) & 0x3); - - u32 key4; - u32 sKeyPacked[4] = { 0, 0, 0, 0 }; - { - sKeyPacked[0] |= 1<<(8*b.x); - sKeyPacked[1] |= 1<<(8*b.y); - sKeyPacked[2] |= 1<<(8*b.z); - sKeyPacked[3] |= 1<<(8*b.w); - - key4 = sKeyPacked[0] + sKeyPacked[1] + sKeyPacked[2] + sKeyPacked[3]; - } - - u32 rankPacked; - u32 sumPacked; - { - rankPacked = localPrefixSum( key4, lIdx, &sumPacked, ldsSortData, WG_SIZE ); - } - - GROUP_LDS_BARRIER; - - u32 newOffset[4] = { 0,0,0,0 }; - { - u32 sumScanned = bit8Scan( sumPacked ); - - u32 scannedKeys[4]; - scannedKeys[0] = 1<<(8*b.x); - scannedKeys[1] = 1<<(8*b.y); - scannedKeys[2] = 1<<(8*b.z); - scannedKeys[3] = 1<<(8*b.w); - { // 4 scans at once - u32 sum4 = 0; - for(int ie=0; ie<4; ie++) - { - u32 tmp = scannedKeys[ie]; - scannedKeys[ie] = sum4; - sum4 += tmp; - } - } - - { - u32 sumPlusRank = sumScanned + rankPacked; - { u32 ie = b.x; - scannedKeys[0] += sumPlusRank; - newOffset[0] = unpack4Key( scannedKeys[0], ie ); - } - { u32 ie = b.y; - scannedKeys[1] += sumPlusRank; - newOffset[1] = unpack4Key( scannedKeys[1], ie ); - } - { u32 ie = b.z; - scannedKeys[2] += sumPlusRank; - newOffset[2] = unpack4Key( scannedKeys[2], ie ); - } - { u32 ie = b.w; - scannedKeys[3] += sumPlusRank; - newOffset[3] = unpack4Key( scannedKeys[3], ie ); - } - } - } - - - GROUP_LDS_BARRIER; - - { - ldsSortData[newOffset[0]] = sortData[0]; - ldsSortData[newOffset[1]] = sortData[1]; - ldsSortData[newOffset[2]] = sortData[2]; - ldsSortData[newOffset[3]] = sortData[3]; - - ldsSortVal[newOffset[0]] = sortVal[0]; - ldsSortVal[newOffset[1]] = sortVal[1]; - ldsSortVal[newOffset[2]] = sortVal[2]; - ldsSortVal[newOffset[3]] = sortVal[3]; - - GROUP_LDS_BARRIER; - - u32 dstAddr = 4*lIdx; - sortData[0] = ldsSortData[dstAddr+0]; - sortData[1] = ldsSortData[dstAddr+1]; - sortData[2] = ldsSortData[dstAddr+2]; - sortData[3] = ldsSortData[dstAddr+3]; - - sortVal[0] = ldsSortVal[dstAddr+0]; - sortVal[1] = ldsSortVal[dstAddr+1]; - sortVal[2] = ldsSortVal[dstAddr+2]; - sortVal[3] = ldsSortVal[dstAddr+3]; - - GROUP_LDS_BARRIER; - } - } -} - - - - -__kernel -__attribute__((reqd_work_group_size(WG_SIZE,1,1))) -void SortAndScatterSortDataKernel( __global const SortDataCL* restrict gSrc, __global const u32* rHistogram, __global SortDataCL* restrict gDst, int4 cb) -{ - __local int ldsSortData[WG_SIZE*ELEMENTS_PER_WORK_ITEM+16]; - __local int ldsSortVal[WG_SIZE*ELEMENTS_PER_WORK_ITEM+16]; - __local u32 localHistogramToCarry[NUM_BUCKET]; - __local u32 localHistogram[NUM_BUCKET*2]; - - u32 gIdx = GET_GLOBAL_IDX; - u32 lIdx = GET_LOCAL_IDX; - u32 wgIdx = GET_GROUP_IDX; - u32 wgSize = GET_GROUP_SIZE; - - const int n = cb.m_n; - const int nWGs = cb.m_nWGs; - const int startBit = cb.m_startBit; - const int nBlocksPerWG = cb.m_nBlocksPerWG; - - if( lIdx < (NUM_BUCKET) ) - { - localHistogramToCarry[lIdx] = rHistogram[lIdx*nWGs + wgIdx]; - } - - GROUP_LDS_BARRIER; - - - const int blockSize = ELEMENTS_PER_WORK_ITEM*WG_SIZE; - - int nBlocks = n/blockSize - nBlocksPerWG*wgIdx; - - int addr = blockSize*nBlocksPerWG*wgIdx + ELEMENTS_PER_WORK_ITEM*lIdx; - - for(int iblock=0; iblock<min(nBlocksPerWG, nBlocks); iblock++, addr+=blockSize) - { - - u32 myHistogram = 0; - - int sortData[ELEMENTS_PER_WORK_ITEM]; - int sortVal[ELEMENTS_PER_WORK_ITEM]; - - for(int i=0; i<ELEMENTS_PER_WORK_ITEM; i++) -#if defined(CHECK_BOUNDARY) - { - sortData[i] = ( addr+i < n )? gSrc[ addr+i ].m_key : 0xffffffff; - sortVal[i] = ( addr+i < n )? gSrc[ addr+i ].m_value : 0xffffffff; - } -#else - { - sortData[i] = gSrc[ addr+i ].m_key; - sortVal[i] = gSrc[ addr+i ].m_value; - } -#endif - - sort4Bits1KeyValue(sortData, sortVal, startBit, lIdx, ldsSortData, ldsSortVal); - - u32 keys[ELEMENTS_PER_WORK_ITEM]; - for(int i=0; i<ELEMENTS_PER_WORK_ITEM; i++) - keys[i] = (sortData[i]>>startBit) & 0xf; - - { // create histogram - u32 setIdx = lIdx/16; - if( lIdx < NUM_BUCKET ) - { - localHistogram[lIdx] = 0; - } - ldsSortData[lIdx] = 0; - GROUP_LDS_BARRIER; - - for(int i=0; i<ELEMENTS_PER_WORK_ITEM; i++) -#if defined(CHECK_BOUNDARY) - if( addr+i < n ) -#endif - -#if defined(NV_GPU) - SET_HISTOGRAM( setIdx, keys[i] )++; -#else - AtomInc( SET_HISTOGRAM( setIdx, keys[i] ) ); -#endif - - GROUP_LDS_BARRIER; - - uint hIdx = NUM_BUCKET+lIdx; - if( lIdx < NUM_BUCKET ) - { - u32 sum = 0; - for(int i=0; i<WG_SIZE/16; i++) - { - sum += SET_HISTOGRAM( i, lIdx ); - } - myHistogram = sum; - localHistogram[hIdx] = sum; - } - GROUP_LDS_BARRIER; - -#if defined(USE_2LEVEL_REDUCE) - if( lIdx < NUM_BUCKET ) - { - localHistogram[hIdx] = localHistogram[hIdx-1]; - GROUP_MEM_FENCE; - - u32 u0, u1, u2; - u0 = localHistogram[hIdx-3]; - u1 = localHistogram[hIdx-2]; - u2 = localHistogram[hIdx-1]; - AtomAdd( localHistogram[hIdx], u0 + u1 + u2 ); - GROUP_MEM_FENCE; - u0 = localHistogram[hIdx-12]; - u1 = localHistogram[hIdx-8]; - u2 = localHistogram[hIdx-4]; - AtomAdd( localHistogram[hIdx], u0 + u1 + u2 ); - GROUP_MEM_FENCE; - } -#else - if( lIdx < NUM_BUCKET ) - { - localHistogram[hIdx] = localHistogram[hIdx-1]; - GROUP_MEM_FENCE; - localHistogram[hIdx] += localHistogram[hIdx-1]; - GROUP_MEM_FENCE; - localHistogram[hIdx] += localHistogram[hIdx-2]; - GROUP_MEM_FENCE; - localHistogram[hIdx] += localHistogram[hIdx-4]; - GROUP_MEM_FENCE; - localHistogram[hIdx] += localHistogram[hIdx-8]; - GROUP_MEM_FENCE; - } -#endif - GROUP_LDS_BARRIER; - } - - { - for(int ie=0; ie<ELEMENTS_PER_WORK_ITEM; ie++) - { - int dataIdx = ELEMENTS_PER_WORK_ITEM*lIdx+ie; - int binIdx = keys[ie]; - int groupOffset = localHistogramToCarry[binIdx]; - int myIdx = dataIdx - localHistogram[NUM_BUCKET+binIdx]; -#if defined(CHECK_BOUNDARY) - if( addr+ie < n ) - { - if ((groupOffset + myIdx)<n) - { - if (sortData[ie]==sortVal[ie]) - { - - SortDataCL tmp; - tmp.m_key = sortData[ie]; - tmp.m_value = sortVal[ie]; - if (tmp.m_key == tmp.m_value) - gDst[groupOffset + myIdx ] = tmp; - } - - } - } -#else - if ((groupOffset + myIdx)<n) - { - gDst[ groupOffset + myIdx ].m_key = sortData[ie]; - gDst[ groupOffset + myIdx ].m_value = sortVal[ie]; - } -#endif - } - } - - GROUP_LDS_BARRIER; - - if( lIdx < NUM_BUCKET ) - { - localHistogramToCarry[lIdx] += myHistogram; - } - GROUP_LDS_BARRIER; - } -} - - - - - - - -__kernel -__attribute__((reqd_work_group_size(WG_SIZE,1,1))) -void SortAndScatterSortDataKernelSerial( __global const SortDataCL* restrict gSrc, __global const u32* rHistogram, __global SortDataCL* restrict gDst, int4 cb) -{ - - u32 gIdx = GET_GLOBAL_IDX; - u32 realLocalIdx = GET_LOCAL_IDX; - u32 wgIdx = GET_GROUP_IDX; - u32 wgSize = GET_GROUP_SIZE; - const int startBit = cb.m_startBit; - const int n = cb.m_n; - const int nWGs = cb.m_nWGs; - const int nBlocksPerWG = cb.m_nBlocksPerWG; - - int counter[NUM_BUCKET]; - - if (realLocalIdx>0) - return; - - for (int c=0;c<NUM_BUCKET;c++) - counter[c]=0; - - const int blockSize = ELEMENTS_PER_WORK_ITEM*WG_SIZE; - - int nBlocks = (n)/blockSize - nBlocksPerWG*wgIdx; - - for(int iblock=0; iblock<min(nBlocksPerWG, nBlocks); iblock++) - { - for (int lIdx=0;lIdx<WG_SIZE;lIdx++) - { - int addr2 = iblock*blockSize + blockSize*nBlocksPerWG*wgIdx + ELEMENTS_PER_WORK_ITEM*lIdx; - - for(int j=0; j<ELEMENTS_PER_WORK_ITEM; j++) - { - int i = addr2+j; - if( i < n ) - { - int tableIdx; - tableIdx = (gSrc[i].m_key>>startBit) & 0xf;//0xf = NUM_TABLES-1 - gDst[rHistogram[tableIdx*nWGs+wgIdx] + counter[tableIdx]] = gSrc[i]; - counter[tableIdx] ++; - } - } - } - } - -} - - -__kernel -__attribute__((reqd_work_group_size(WG_SIZE,1,1))) -void SortAndScatterKernelSerial( __global const u32* restrict gSrc, __global const u32* rHistogram, __global u32* restrict gDst, int4 cb ) -{ - - u32 gIdx = GET_GLOBAL_IDX; - u32 realLocalIdx = GET_LOCAL_IDX; - u32 wgIdx = GET_GROUP_IDX; - u32 wgSize = GET_GROUP_SIZE; - const int startBit = cb.m_startBit; - const int n = cb.m_n; - const int nWGs = cb.m_nWGs; - const int nBlocksPerWG = cb.m_nBlocksPerWG; - - int counter[NUM_BUCKET]; - - if (realLocalIdx>0) - return; - - for (int c=0;c<NUM_BUCKET;c++) - counter[c]=0; - - const int blockSize = ELEMENTS_PER_WORK_ITEM*WG_SIZE; - - int nBlocks = (n)/blockSize - nBlocksPerWG*wgIdx; - - for(int iblock=0; iblock<min(nBlocksPerWG, nBlocks); iblock++) - { - for (int lIdx=0;lIdx<WG_SIZE;lIdx++) - { - int addr2 = iblock*blockSize + blockSize*nBlocksPerWG*wgIdx + ELEMENTS_PER_WORK_ITEM*lIdx; - - for(int j=0; j<ELEMENTS_PER_WORK_ITEM; j++) - { - int i = addr2+j; - if( i < n ) - { - int tableIdx; - tableIdx = (gSrc[i]>>startBit) & 0xf;//0xf = NUM_TABLES-1 - gDst[rHistogram[tableIdx*nWGs+wgIdx] + counter[tableIdx]] = gSrc[i]; - counter[tableIdx] ++; - } - } - } - } - -}
\ No newline at end of file diff --git a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/RadixSort32KernelsCL.h b/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/RadixSort32KernelsCL.h deleted file mode 100644 index fb4bdda303..0000000000 --- a/thirdparty/bullet/Bullet3OpenCL/ParallelPrimitives/kernels/RadixSort32KernelsCL.h +++ /dev/null @@ -1,909 +0,0 @@ -//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"; |