Merge various commits into a single commit.
Commits after:
2014-03-03 Draft PLBVH construction using binary radix tree.
f19f853685
Are merged into a single commit; this includes:
03-10 Remove single launch build AABB kernel.
03-10 Add kernels for setting PLBVH AABBs using distance from root.
03-10 Use faster morton code, remove convertChildNodeFormat kernel.
03-09 Add duplicate morton code handling to binary radix construct.
03-09 Remove slower PLBVH constructors.
03-08 Add binary radix tree construct using binary search.
03-06 Remove slowest PLBVH constructor, fix implicit construct AABB.
03-04 Test various optimizations for PLBVH binary radix tree construct.
This commit is contained in:
Jackson Lee
@@ -41,34 +41,29 @@ static const char* parallelLinearBvhCL= \
|
||||
" //........ ........ ......12 3456789A //x\n"
|
||||
" //....1..2 ..3..4.. 5..6..7. .8..9..A //x after interleaving bits\n"
|
||||
" \n"
|
||||
" //........ ....1234 56789A12 3456789A //x |= (x << 10)\n"
|
||||
" //........ ....1111 1....... ...11111 //0x 00 0F 80 1F\n"
|
||||
" //........ ....1234 5....... ...6789A //x = ( x | (x << 10) ) & 0x000F801F; \n"
|
||||
" //......12 3456789A ......12 3456789A //x ^ (x << 16)\n"
|
||||
" //11111111 ........ ........ 11111111 //0x FF 00 00 FF\n"
|
||||
" //......12 ........ ........ 3456789A //x = (x ^ (x << 16)) & 0xFF0000FF;\n"
|
||||
" \n"
|
||||
" //.......1 23451234 5.....67 89A6789A //x |= (x << 5)\n"
|
||||
" //.......1 1.....11 1.....11 .....111 //0x 01 83 83 07\n"
|
||||
" //.......1 2.....34 5.....67 .....89A //x = ( x | (x << 5) ) & 0x01838307;\n"
|
||||
" //......12 ........ 3456789A 3456789A //x ^ (x << 8)\n"
|
||||
" //......11 ........ 1111.... ....1111 //0x 03 00 F0 0F\n"
|
||||
" //......12 ........ 3456.... ....789A //x = (x ^ (x << 8)) & 0x0300F00F;\n"
|
||||
" \n"
|
||||
" //....12.1 2..34534 5..67.67 ..89A89A //x |= (x << 3)\n"
|
||||
" //....1... 1..1...1 1..1...1 ..1...11 //0x 08 91 91 23\n"
|
||||
" //....1... 2..3...4 5..6...7 ..8...9A //x = ( x | (x << 3) ) & 0x08919123;\n"
|
||||
" //..12..12 ....3456 3456.... 789A789A //x ^ (x << 4)\n"
|
||||
" //......11 ....11.. ..11.... 11....11 //0x 03 0C 30 C3\n"
|
||||
" //......12 ....34.. ..56.... 78....9A //x = (x ^ (x << 4)) & 0x030C30C3;\n"
|
||||
" \n"
|
||||
" //...11..2 2.33..4N 5.66..77 .88..9NA //x |= (x << 1) ( N indicates overlapping bits, first overlap is bit {4,5} second is {9,A} )\n"
|
||||
" //....1..1 ..1...1. 1..1..1. .1...1.1 //0x 09 22 92 45\n"
|
||||
" //....1..2 ..3...4. 5..6..7. .8...9.A //x = ( x | (x << 1) ) & 0x09229245;\n"
|
||||
" \n"
|
||||
" //...11.22 .33..445 5.66.77. 88..99AA //x |= (x << 1)\n"
|
||||
" //....1..1 ..1..1.. 1..1..1. .1..1..1 //0x 09 34 92 29\n"
|
||||
" //....1..2 ..3..4.. 5..6..7. .8..9..A //x = ( x | (x << 1) ) & 0x09349229;\n"
|
||||
" //....1212 ..3434.. 5656..78 78..9A9A //x ^ (x << 2)\n"
|
||||
" //....1..1 ..1..1.. 1..1..1. .1..1..1 //0x 09 24 92 49\n"
|
||||
" //....1..2 ..3..4.. 5..6..7. .8..9..A //x = (x ^ (x << 2)) & 0x09249249;\n"
|
||||
" \n"
|
||||
" //........ ........ ......11 11111111 //0x000003FF\n"
|
||||
" x &= 0x000003FF; //Clear all bits above bit 10\n"
|
||||
" \n"
|
||||
" x = ( x | (x << 10) ) & 0x000F801F;\n"
|
||||
" x = ( x | (x << 5) ) & 0x01838307;\n"
|
||||
" x = ( x | (x << 3) ) & 0x08919123;\n"
|
||||
" x = ( x | (x << 1) ) & 0x09229245;\n"
|
||||
" x = ( x | (x << 1) ) & 0x09349229;\n"
|
||||
" x = (x ^ (x << 16)) & 0xFF0000FF;\n"
|
||||
" x = (x ^ (x << 8)) & 0x0300F00F;\n"
|
||||
" x = (x ^ (x << 4)) & 0x030C30C3;\n"
|
||||
" x = (x ^ (x << 2)) & 0x09249249;\n"
|
||||
" \n"
|
||||
" return x;\n"
|
||||
"}\n"
|
||||
@@ -150,143 +145,10 @@ static const char* parallelLinearBvhCL= \
|
||||
"#define B3_PLVBH_TRAVERSE_MAX_STACK_SIZE 128\n"
|
||||
"//The most significant bit(0x80000000) of a int32 is used to distinguish between leaf and internal nodes.\n"
|
||||
"//If it is set, then the index is for an internal node; otherwise, it is a leaf node. \n"
|
||||
"//In both cases, the bit should be cleared to access the index.\n"
|
||||
"//In both cases, the bit should be cleared to access the actual node index.\n"
|
||||
"int isLeafNode(int index) { return (index >> 31 == 0); }\n"
|
||||
"int getIndexWithInternalNodeMarkerRemoved(int index) { return index & (~0x80000000); }\n"
|
||||
"int getIndexWithInternalNodeMarkerSet(int isLeaf, int index) { return (isLeaf) ? index : (index | 0x80000000); }\n"
|
||||
"__kernel void constructBinaryTree(__global int* firstIndexOffsetPerLevel,\n"
|
||||
" __global int* numNodesPerLevel,\n"
|
||||
" __global int2* out_internalNodeChildIndices, \n"
|
||||
" __global int* out_internalNodeParentNodes, \n"
|
||||
" __global int* out_leafNodeParentNodes, \n"
|
||||
" int numLevels, int numInternalNodes)\n"
|
||||
"{\n"
|
||||
" int internalNodeIndex = get_global_id(0);\n"
|
||||
" if(internalNodeIndex >= numInternalNodes) return;\n"
|
||||
" \n"
|
||||
" //Find the level that this node is in, using linear search(could replace with binary search)\n"
|
||||
" int level = 0;\n"
|
||||
" int numInternalLevels = numLevels - 1; //All levels except the last are internal nodes\n"
|
||||
" for(; level < numInternalLevels; ++level)\n"
|
||||
" {\n"
|
||||
" if( firstIndexOffsetPerLevel[level] <= internalNodeIndex && internalNodeIndex < firstIndexOffsetPerLevel[level + 1]) break;\n"
|
||||
" }\n"
|
||||
" \n"
|
||||
" //Check lower levels to find child nodes\n"
|
||||
" //Left child is always in the next level, but the same does not apply to the right child\n"
|
||||
" int indexInLevel = internalNodeIndex - firstIndexOffsetPerLevel[level];\n"
|
||||
" int firstIndexInNextLevel = firstIndexOffsetPerLevel[level + 1]; //Should never be out of bounds(see for loop above)\n"
|
||||
" \n"
|
||||
" int leftChildLevel = level + 1;\n"
|
||||
" int leftChildIndex = firstIndexInNextLevel + indexInLevel * 2;\n"
|
||||
" \n"
|
||||
" int rightChildLevel = level + 1;\n"
|
||||
" int rightChildIndex = leftChildIndex + 1;\n"
|
||||
" \n"
|
||||
" //Under certain conditions, the right child index as calculated above is invalid; need to find the correct index\n"
|
||||
" //\n"
|
||||
" //First condition: must be at least 2 levels apart from the leaf node level;\n"
|
||||
" //if the current level is right next to the leaf node level, then the right child\n"
|
||||
" //will never be invalid due to the way the nodes are allocated (also avoid a out-of-bounds memory access)\n"
|
||||
" //\n"
|
||||
" //Second condition: not enough nodes in the next level for each parent to have 2 children, so the right child is invalid\n"
|
||||
" //\n"
|
||||
" //Third condition: must be the last node in its level\n"
|
||||
" if( level < numLevels - 2 \n"
|
||||
" && numNodesPerLevel[level] * 2 > numNodesPerLevel[level + 1] \n"
|
||||
" && indexInLevel == numNodesPerLevel[level] - 1 )\n"
|
||||
" {\n"
|
||||
" //Check lower levels until we find a node without a parent\n"
|
||||
" for(; rightChildLevel < numLevels - 1; ++rightChildLevel)\n"
|
||||
" {\n"
|
||||
" int rightChildNextLevel = rightChildLevel + 1;\n"
|
||||
" \n"
|
||||
" //If this branch is taken, it means that the last node in rightChildNextLevel has no parent\n"
|
||||
" if( numNodesPerLevel[rightChildLevel] * 2 < numNodesPerLevel[rightChildNextLevel] )\n"
|
||||
" {\n"
|
||||
" //Set the node to the last node in rightChildNextLevel\n"
|
||||
" rightChildLevel = rightChildNextLevel;\n"
|
||||
" rightChildIndex = firstIndexOffsetPerLevel[rightChildNextLevel] + numNodesPerLevel[rightChildNextLevel] - 1;\n"
|
||||
" break;\n"
|
||||
" }\n"
|
||||
" }\n"
|
||||
" }\n"
|
||||
" \n"
|
||||
" int isLeftChildLeaf = (leftChildLevel >= numLevels - 1);\n"
|
||||
" int isRightChildLeaf = (rightChildLevel >= numLevels - 1);\n"
|
||||
" \n"
|
||||
" //If left/right child is a leaf node, the index needs to be corrected\n"
|
||||
" //the way the index is calculated assumes that the leaf and internal nodes are in a contiguous array,\n"
|
||||
" //with leaf nodes at the end of the array; in actuality, the leaf and internal nodes are in separate arrays\n"
|
||||
" {\n"
|
||||
" int leafNodeLevel = numLevels - 1;\n"
|
||||
" leftChildIndex = (isLeftChildLeaf) ? leftChildIndex - firstIndexOffsetPerLevel[leafNodeLevel] : leftChildIndex;\n"
|
||||
" rightChildIndex = (isRightChildLeaf) ? rightChildIndex - firstIndexOffsetPerLevel[leafNodeLevel] : rightChildIndex;\n"
|
||||
" }\n"
|
||||
" \n"
|
||||
" //Set the negative sign bit if the node is internal\n"
|
||||
" int2 childIndices;\n"
|
||||
" childIndices.x = getIndexWithInternalNodeMarkerSet(isLeftChildLeaf, leftChildIndex);\n"
|
||||
" childIndices.y = getIndexWithInternalNodeMarkerSet(isRightChildLeaf, rightChildIndex);\n"
|
||||
" out_internalNodeChildIndices[internalNodeIndex] = childIndices;\n"
|
||||
" \n"
|
||||
" //Assign parent node index to children\n"
|
||||
" __global int* out_leftChildParentNodeIndices = (isLeftChildLeaf) ? out_leafNodeParentNodes : out_internalNodeParentNodes;\n"
|
||||
" out_leftChildParentNodeIndices[leftChildIndex] = internalNodeIndex;\n"
|
||||
" \n"
|
||||
" __global int* out_rightChildParentNodeIndices = (isRightChildLeaf) ? out_leafNodeParentNodes : out_internalNodeParentNodes;\n"
|
||||
" out_rightChildParentNodeIndices[rightChildIndex] = internalNodeIndex;\n"
|
||||
"}\n"
|
||||
"__kernel void determineInternalNodeAabbs(__global int* firstIndexOffsetPerLevel,\n"
|
||||
" __global int* numNodesPerLevel, \n"
|
||||
" __global int2* internalNodeChildIndices,\n"
|
||||
" __global SortDataCL* mortonCodesAndAabbIndices,\n"
|
||||
" __global b3AabbCL* leafNodeAabbs, \n"
|
||||
" __global int2* out_internalNodeLeafIndexRanges,\n"
|
||||
" __global b3AabbCL* out_internalNodeAabbs, \n"
|
||||
" int numLevels, int numInternalNodes, int level)\n"
|
||||
"{\n"
|
||||
" int i = get_global_id(0);\n"
|
||||
" if(i >= numInternalNodes) return;\n"
|
||||
" \n"
|
||||
" //For each node in a level, check its child nodes to determine its AABB\n"
|
||||
" {\n"
|
||||
" int indexInLevel = i; //Index relative to firstIndexOffsetPerLevel[level]\n"
|
||||
" \n"
|
||||
" int numNodesInLevel = numNodesPerLevel[level];\n"
|
||||
" if(indexInLevel < numNodesInLevel)\n"
|
||||
" {\n"
|
||||
" int internalNodeIndexGlobal = indexInLevel + firstIndexOffsetPerLevel[level];\n"
|
||||
" int2 childIndicies = internalNodeChildIndices[internalNodeIndexGlobal];\n"
|
||||
" \n"
|
||||
" int leftChildIndex = getIndexWithInternalNodeMarkerRemoved(childIndicies.x);\n"
|
||||
" int rightChildIndex = getIndexWithInternalNodeMarkerRemoved(childIndicies.y);\n"
|
||||
" \n"
|
||||
" int isLeftChildLeaf = isLeafNode(childIndicies.x);\n"
|
||||
" int isRightChildLeaf = isLeafNode(childIndicies.y);\n"
|
||||
" \n"
|
||||
" //left/RightChildLeafIndex == Rigid body indicies\n"
|
||||
" int leftChildLeafIndex = (isLeftChildLeaf) ? mortonCodesAndAabbIndices[leftChildIndex].m_value : -1;\n"
|
||||
" int rightChildLeafIndex = (isRightChildLeaf) ? mortonCodesAndAabbIndices[rightChildIndex].m_value : -1;\n"
|
||||
" \n"
|
||||
" b3AabbCL leftChildAabb = (isLeftChildLeaf) ? leafNodeAabbs[leftChildLeafIndex] : out_internalNodeAabbs[leftChildIndex];\n"
|
||||
" b3AabbCL rightChildAabb = (isRightChildLeaf) ? leafNodeAabbs[rightChildLeafIndex] : out_internalNodeAabbs[rightChildIndex];\n"
|
||||
" \n"
|
||||
" //\n"
|
||||
" b3AabbCL internalNodeAabb;\n"
|
||||
" internalNodeAabb.m_min = b3Min(leftChildAabb.m_min, rightChildAabb.m_min);\n"
|
||||
" internalNodeAabb.m_max = b3Max(leftChildAabb.m_max, rightChildAabb.m_max);\n"
|
||||
" out_internalNodeAabbs[internalNodeIndexGlobal] = internalNodeAabb;\n"
|
||||
" \n"
|
||||
" //For index range, x == min and y == max; left child always has lower index\n"
|
||||
" int2 leafIndexRange;\n"
|
||||
" leafIndexRange.x = (isLeftChildLeaf) ? leftChildIndex : out_internalNodeLeafIndexRanges[leftChildIndex].x;\n"
|
||||
" leafIndexRange.y = (isRightChildLeaf) ? rightChildIndex : out_internalNodeLeafIndexRanges[rightChildIndex].y;\n"
|
||||
" \n"
|
||||
" out_internalNodeLeafIndexRanges[internalNodeIndexGlobal] = leafIndexRange;\n"
|
||||
" }\n"
|
||||
" }\n"
|
||||
"}\n"
|
||||
"//From sap.cl\n"
|
||||
"#define NEW_PAIR_MARKER -1\n"
|
||||
"bool TestAabbAgainstAabb2(const b3AabbCL* aabb1, const b3AabbCL* aabb2)\n"
|
||||
@@ -539,78 +401,57 @@ static const char* parallelLinearBvhCL= \
|
||||
" }\n"
|
||||
" }\n"
|
||||
"}\n"
|
||||
"#define B3_PLBVH_LINKED_LIST_INVALID_NODE -1\n"
|
||||
"int longestCommonPrefix(int i, int j) { return clz(i ^ j); }\n"
|
||||
"__kernel void computePrefixAndInitPointers(__global SortDataCL* mortonCodesAndAabbIndices,\n"
|
||||
" __global int* out_commonPrefixes,\n"
|
||||
" __global int* out_leftInternalNodePointers, \n"
|
||||
" __global int* out_rightInternalNodePointers,\n"
|
||||
"//Set so that it is always greater than the actual common prefixes, and never selected as a parent node.\n"
|
||||
"//If there are no duplicates, then the highest common prefix is 32 or 64, depending on the number of bits used for the z-curve.\n"
|
||||
"//Duplicates common prefixes increase the highest common prefix by N, where 2^N is the number of duplicate nodes.\n"
|
||||
"#define B3_PLBVH_INVALID_COMMON_PREFIX 128\n"
|
||||
"#define B3_PLBVH_ROOT_NODE_MARKER -1\n"
|
||||
"#define b3Int64 long\n"
|
||||
"int computeCommonPrefixLength(b3Int64 i, b3Int64 j) { return (int)clz(i ^ j); }\n"
|
||||
"b3Int64 computeCommonPrefix(b3Int64 i, b3Int64 j) \n"
|
||||
"{\n"
|
||||
" //This function only needs to return (i & j) in order for the algorithm to work,\n"
|
||||
" //but it may help with debugging to mask out the lower bits.\n"
|
||||
" b3Int64 commonPrefixLength = (b3Int64)computeCommonPrefixLength(i, j);\n"
|
||||
" b3Int64 sharedBits = i & j;\n"
|
||||
" b3Int64 bitmask = ((b3Int64)(~0)) << (64 - commonPrefixLength); //Set all bits after the common prefix to 0\n"
|
||||
" \n"
|
||||
" return sharedBits & bitmask;\n"
|
||||
"}\n"
|
||||
"int getSharedPrefixLength(b3Int64 prefixA, int prefixLengthA, b3Int64 prefixB, int prefixLengthB)\n"
|
||||
"{\n"
|
||||
" return b3Min( computeCommonPrefixLength(prefixA, prefixB), b3Min(prefixLengthA, prefixLengthB) );\n"
|
||||
"}\n"
|
||||
"__kernel void computeAdjacentPairCommonPrefix(__global SortDataCL* mortonCodesAndAabbIndices,\n"
|
||||
" __global b3Int64* out_commonPrefixes,\n"
|
||||
" __global int* out_commonPrefixLengths,\n"
|
||||
" int numInternalNodes)\n"
|
||||
"{\n"
|
||||
" int internalNodeIndex = get_global_id(0);\n"
|
||||
" if (internalNodeIndex >= numInternalNodes) return;\n"
|
||||
" \n"
|
||||
" //Compute common prefix\n"
|
||||
" {\n"
|
||||
" //Here, (internalNodeIndex + 1) is never out of bounds since it is a leaf node index,\n"
|
||||
" //and the number of internal nodes is always numLeafNodes - 1\n"
|
||||
" int leftLeafMortonCode = mortonCodesAndAabbIndices[internalNodeIndex].m_key;\n"
|
||||
" int rightLeafMortonCode = mortonCodesAndAabbIndices[internalNodeIndex + 1].m_key;\n"
|
||||
" //Here, (internalNodeIndex + 1) is never out of bounds since it is a leaf node index,\n"
|
||||
" //and the number of internal nodes is always numLeafNodes - 1\n"
|
||||
" int leftLeafIndex = internalNodeIndex;\n"
|
||||
" int rightLeafIndex = internalNodeIndex + 1;\n"
|
||||
" \n"
|
||||
" out_commonPrefixes[internalNodeIndex] = longestCommonPrefix(leftLeafMortonCode, rightLeafMortonCode);\n"
|
||||
" }\n"
|
||||
" int leftLeafMortonCode = mortonCodesAndAabbIndices[leftLeafIndex].m_key;\n"
|
||||
" int rightLeafMortonCode = mortonCodesAndAabbIndices[rightLeafIndex].m_key;\n"
|
||||
" \n"
|
||||
" //Assign neighbor pointers of this node\n"
|
||||
" {\n"
|
||||
" int leftInternalIndex = internalNodeIndex - 1;\n"
|
||||
" int rightInternalIndex = internalNodeIndex + 1;\n"
|
||||
" \n"
|
||||
" out_leftInternalNodePointers[internalNodeIndex] = (leftInternalIndex >= 0) ? leftInternalIndex : B3_PLBVH_LINKED_LIST_INVALID_NODE;\n"
|
||||
" out_rightInternalNodePointers[internalNodeIndex] = (rightInternalIndex < numInternalNodes) ? rightInternalIndex : B3_PLBVH_LINKED_LIST_INVALID_NODE;\n"
|
||||
" }\n"
|
||||
" //Binary radix tree construction algorithm does not work if there are duplicate morton codes.\n"
|
||||
" //Append the index of each leaf node to each morton code so that there are no duplicates.\n"
|
||||
" //The algorithm also requires that the morton codes are sorted in ascending order; this requirement\n"
|
||||
" //is also satisfied with this method, as (leftLeafIndex < rightLeafIndex) is always true.\n"
|
||||
" //\n"
|
||||
" //upsample(a, b) == ( ((b3Int64)a) << 32) | b\n"
|
||||
" b3Int64 nonduplicateLeftMortonCode = upsample(leftLeafMortonCode, leftLeafIndex);\n"
|
||||
" b3Int64 nonduplicateRightMortonCode = upsample(rightLeafMortonCode, rightLeafIndex);\n"
|
||||
" \n"
|
||||
" out_commonPrefixes[internalNodeIndex] = computeCommonPrefix(nonduplicateLeftMortonCode, nonduplicateRightMortonCode);\n"
|
||||
" out_commonPrefixLengths[internalNodeIndex] = computeCommonPrefixLength(nonduplicateLeftMortonCode, nonduplicateRightMortonCode);\n"
|
||||
"}\n"
|
||||
"__kernel void correctDuplicatePrefixes(__global int* commonPrefixes, __global int* out_maxCommonPrefix, int numInternalNodes)\n"
|
||||
"{\n"
|
||||
" int internalNodeIndex = get_global_id(0);\n"
|
||||
" if (internalNodeIndex >= numInternalNodes) return;\n"
|
||||
" \n"
|
||||
" int commonPrefix = commonPrefixes[internalNodeIndex];\n"
|
||||
" \n"
|
||||
" //Linear search to find the size of the subtree\n"
|
||||
" int firstSubTreeIndex = internalNodeIndex;\n"
|
||||
" int lastSubTreeIndex = internalNodeIndex;\n"
|
||||
" \n"
|
||||
" while(firstSubTreeIndex - 1 >= 0 && commonPrefix == commonPrefixes[firstSubTreeIndex - 1]) --firstSubTreeIndex;\n"
|
||||
" while(lastSubTreeIndex + 1 < numInternalNodes && commonPrefix == commonPrefixes[lastSubTreeIndex + 1]) ++lastSubTreeIndex;\n"
|
||||
" \n"
|
||||
" //Fix duplicate common prefixes by incrementing them so that a subtree is formed.\n"
|
||||
" //Recursively divide the tree until the position of the split is this node's index.\n"
|
||||
" //Every time this node is not the split node, increment the common prefix.\n"
|
||||
" int isCurrentSplitNode = false;\n"
|
||||
" int correctedCommonPrefix = commonPrefix;\n"
|
||||
" \n"
|
||||
" while(!isCurrentSplitNode)\n"
|
||||
" {\n"
|
||||
" int numInternalNodesInSubTree = lastSubTreeIndex - firstSubTreeIndex + 1;\n"
|
||||
" int splitNodeIndex = firstSubTreeIndex + numInternalNodesInSubTree / 2;\n"
|
||||
" \n"
|
||||
" if(internalNodeIndex > splitNodeIndex) firstSubTreeIndex = splitNodeIndex + 1;\n"
|
||||
" else if(internalNodeIndex < splitNodeIndex) lastSubTreeIndex = splitNodeIndex - 1;\n"
|
||||
" //else if(internalNodeIndex == splitNodeIndex) break;\n"
|
||||
" \n"
|
||||
" isCurrentSplitNode = (internalNodeIndex == splitNodeIndex);\n"
|
||||
" if(!isCurrentSplitNode) correctedCommonPrefix++;\n"
|
||||
" }\n"
|
||||
" \n"
|
||||
" commonPrefixes[internalNodeIndex] = correctedCommonPrefix;\n"
|
||||
" atomic_max(out_maxCommonPrefix, correctedCommonPrefix);\n"
|
||||
"}\n"
|
||||
"//Set so that it is always greater than the actual common prefixes, and never selected as a parent node.\n"
|
||||
"//If there are no duplicates, then the highest common prefix is 32 or 64, depending on the number of bits used for the z-curve.\n"
|
||||
"//Duplicates common prefixes increase the highest common prefix by N, where 2^N is the number of duplicate nodes.\n"
|
||||
"#define B3_PLBVH_INVALID_COMMON_PREFIX 128\n"
|
||||
"__kernel void buildBinaryRadixTreeLeafNodes(__global int* commonPrefixes, __global int* out_leftChildNodes, \n"
|
||||
" __global int* out_rightChildNodes, int numLeafNodes)\n"
|
||||
"__kernel void buildBinaryRadixTreeLeafNodes(__global int* commonPrefixLengths, __global int* out_leafNodeParentNodes,\n"
|
||||
" __global int2* out_childNodes, int numLeafNodes)\n"
|
||||
"{\n"
|
||||
" int leafNodeIndex = get_global_id(0);\n"
|
||||
" if (leafNodeIndex >= numLeafNodes) return;\n"
|
||||
@@ -620,8 +461,8 @@ static const char* parallelLinearBvhCL= \
|
||||
" int leftSplitIndex = leafNodeIndex - 1;\n"
|
||||
" int rightSplitIndex = leafNodeIndex;\n"
|
||||
" \n"
|
||||
" int leftCommonPrefix = (leftSplitIndex >= 0) ? commonPrefixes[leftSplitIndex] : B3_PLBVH_INVALID_COMMON_PREFIX;\n"
|
||||
" int rightCommonPrefix = (rightSplitIndex < numInternalNodes) ? commonPrefixes[rightSplitIndex] : B3_PLBVH_INVALID_COMMON_PREFIX;\n"
|
||||
" int leftCommonPrefix = (leftSplitIndex >= 0) ? commonPrefixLengths[leftSplitIndex] : B3_PLBVH_INVALID_COMMON_PREFIX;\n"
|
||||
" int rightCommonPrefix = (rightSplitIndex < numInternalNodes) ? commonPrefixLengths[rightSplitIndex] : B3_PLBVH_INVALID_COMMON_PREFIX;\n"
|
||||
" \n"
|
||||
" //Parent node is the highest adjacent common prefix that is lower than the node's common prefix\n"
|
||||
" //Leaf nodes are considered as having the highest common prefix\n"
|
||||
@@ -633,71 +474,218 @@ static const char* parallelLinearBvhCL= \
|
||||
" if(rightCommonPrefix == B3_PLBVH_INVALID_COMMON_PREFIX) isLeftHigherCommonPrefix = true;\n"
|
||||
" \n"
|
||||
" int parentNodeIndex = (isLeftHigherCommonPrefix) ? leftSplitIndex : rightSplitIndex;\n"
|
||||
" out_leafNodeParentNodes[leafNodeIndex] = parentNodeIndex;\n"
|
||||
" \n"
|
||||
" //If the left node is the parent, then this node is its right child and vice versa\n"
|
||||
" __global int* out_childNode = (isLeftHigherCommonPrefix) ? out_rightChildNodes : out_leftChildNodes;\n"
|
||||
" int isRightChild = (isLeftHigherCommonPrefix); //If the left node is the parent, then this node is its right child and vice versa\n"
|
||||
" \n"
|
||||
" //out_childNodesAsInt[0] == int2.x == left child\n"
|
||||
" //out_childNodesAsInt[1] == int2.y == right child\n"
|
||||
" int isLeaf = 1;\n"
|
||||
" out_childNode[parentNodeIndex] = getIndexWithInternalNodeMarkerSet(isLeaf, leafNodeIndex);\n"
|
||||
" __global int* out_childNodesAsInt = (__global int*)(&out_childNodes[parentNodeIndex]);\n"
|
||||
" out_childNodesAsInt[isRightChild] = getIndexWithInternalNodeMarkerSet(isLeaf, leafNodeIndex);\n"
|
||||
"}\n"
|
||||
"__kernel void buildBinaryRadixTreeInternalNodes(__global int* commonPrefixes, __global SortDataCL* mortonCodesAndAabbIndices,\n"
|
||||
" __global int* leftChildNodes, __global int* rightChildNodes,\n"
|
||||
" __global int* leftNeighborPointers, __global int* rightNeighborPointers,\n"
|
||||
" __global b3AabbCL* leafNodeAabbs, __global b3AabbCL* internalNodeAabbs,\n"
|
||||
" __global int* out_rootNodeIndex,\n"
|
||||
" int processedCommonPrefix, int numInternalNodes)\n"
|
||||
"__kernel void buildBinaryRadixTreeInternalNodes(__global b3Int64* commonPrefixes, __global int* commonPrefixLengths,\n"
|
||||
" __global int2* out_childNodes,\n"
|
||||
" __global int* out_internalNodeParentNodes, __global int* out_rootNodeIndex,\n"
|
||||
" __global int* TEMP_out_leftLowerPrefix, __global int* TEMP_out_rightLowerPrefix,\n"
|
||||
" __global int* TEMP_spl_left, __global int* TEMP_spl_right,\n"
|
||||
" int numInternalNodes)\n"
|
||||
"{\n"
|
||||
" int internalNodeIndex = get_global_id(0);\n"
|
||||
" if (internalNodeIndex >= numInternalNodes) return;\n"
|
||||
" int internalNodeIndex = get_group_id(0) * get_local_size(0) + get_local_id(0);\n"
|
||||
" if(internalNodeIndex >= numInternalNodes) return;\n"
|
||||
" \n"
|
||||
" int commonPrefix = commonPrefixes[internalNodeIndex];\n"
|
||||
" if (commonPrefix == processedCommonPrefix)\n"
|
||||
" b3Int64 nodePrefix = commonPrefixes[internalNodeIndex];\n"
|
||||
" int nodePrefixLength = commonPrefixLengths[internalNodeIndex];\n"
|
||||
" \n"
|
||||
"//#define USE_LINEAR_SEARCH\n"
|
||||
"#ifdef USE_LINEAR_SEARCH\n"
|
||||
" int leftIndex = -1;\n"
|
||||
" int rightIndex = -1;\n"
|
||||
" \n"
|
||||
" for(int i = internalNodeIndex - 1; i >= 0; --i)\n"
|
||||
" {\n"
|
||||
" //Check neighbors and compare the common prefix to select the parent node\n"
|
||||
" int leftNodeIndex = leftNeighborPointers[internalNodeIndex];\n"
|
||||
" int rightNodeIndex = rightNeighborPointers[internalNodeIndex];\n"
|
||||
" \n"
|
||||
" int leftCommonPrefix = (leftNodeIndex != B3_PLBVH_LINKED_LIST_INVALID_NODE) ? commonPrefixes[leftNodeIndex] : B3_PLBVH_INVALID_COMMON_PREFIX;\n"
|
||||
" int rightCommonPrefix = (rightNodeIndex != B3_PLBVH_LINKED_LIST_INVALID_NODE) ? commonPrefixes[rightNodeIndex] : B3_PLBVH_INVALID_COMMON_PREFIX;\n"
|
||||
" \n"
|
||||
" //Parent node is the highest common prefix that is lower than the node's common prefix\n"
|
||||
" //Since the nodes with lower common prefixes are removed, that condition does not have to be tested for,\n"
|
||||
" //and we only need to pick the node with the higher prefix.\n"
|
||||
" int isLeftHigherCommonPrefix = (leftCommonPrefix > rightCommonPrefix);\n"
|
||||
" \n"
|
||||
" //\n"
|
||||
" if(leftCommonPrefix == B3_PLBVH_INVALID_COMMON_PREFIX) isLeftHigherCommonPrefix = false;\n"
|
||||
" else if(rightCommonPrefix == B3_PLBVH_INVALID_COMMON_PREFIX) isLeftHigherCommonPrefix = true;\n"
|
||||
" \n"
|
||||
" int isRootNode = false;\n"
|
||||
" if(leftCommonPrefix == B3_PLBVH_INVALID_COMMON_PREFIX && rightCommonPrefix == B3_PLBVH_INVALID_COMMON_PREFIX) isRootNode = true;\n"
|
||||
" \n"
|
||||
" int parentNodeIndex = (isLeftHigherCommonPrefix) ? leftNodeIndex : rightNodeIndex;\n"
|
||||
" \n"
|
||||
" //If the left node is the parent, then this node is its right child and vice versa\n"
|
||||
" __global int* out_childNode = (isLeftHigherCommonPrefix) ? rightChildNodes : leftChildNodes;\n"
|
||||
" \n"
|
||||
" int isLeaf = 0;\n"
|
||||
" if(!isRootNode) out_childNode[parentNodeIndex] = getIndexWithInternalNodeMarkerSet(isLeaf, internalNodeIndex);\n"
|
||||
" \n"
|
||||
" if(isRootNode) *out_rootNodeIndex = getIndexWithInternalNodeMarkerSet(isLeaf, internalNodeIndex);\n"
|
||||
" \n"
|
||||
" //Remove this node from the linked list, \n"
|
||||
" //so that the left and right nodes point at each other instead of this node\n"
|
||||
" if(leftNodeIndex != B3_PLBVH_LINKED_LIST_INVALID_NODE) rightNeighborPointers[leftNodeIndex] = rightNodeIndex;\n"
|
||||
" if(rightNodeIndex != B3_PLBVH_LINKED_LIST_INVALID_NODE) leftNeighborPointers[rightNodeIndex] = leftNodeIndex;\n"
|
||||
" \n"
|
||||
" //For debug\n"
|
||||
" leftNeighborPointers[internalNodeIndex] = -2;\n"
|
||||
" rightNeighborPointers[internalNodeIndex] = -2;\n"
|
||||
" int nodeLeftSharedPrefixLength = getSharedPrefixLength(nodePrefix, nodePrefixLength, commonPrefixes[i], commonPrefixLengths[i]);\n"
|
||||
" if(nodeLeftSharedPrefixLength < nodePrefixLength)\n"
|
||||
" {\n"
|
||||
" leftIndex = i;\n"
|
||||
" break;\n"
|
||||
" }\n"
|
||||
" }\n"
|
||||
" \n"
|
||||
" //Processing occurs from highest common prefix to lowest common prefix\n"
|
||||
" //Nodes in the previously processed level have their children set, so we merge their child AABBs here\n"
|
||||
" if (commonPrefix == processedCommonPrefix + 1)\n"
|
||||
" for(int i = internalNodeIndex + 1; i < numInternalNodes; ++i)\n"
|
||||
" {\n"
|
||||
" int leftChildIndex = leftChildNodes[internalNodeIndex];\n"
|
||||
" int rightChildIndex = rightChildNodes[internalNodeIndex];\n"
|
||||
" int nodeRightSharedPrefixLength = getSharedPrefixLength(nodePrefix, nodePrefixLength, commonPrefixes[i], commonPrefixLengths[i]);\n"
|
||||
" if(nodeRightSharedPrefixLength < nodePrefixLength)\n"
|
||||
" {\n"
|
||||
" rightIndex = i;\n"
|
||||
" break;\n"
|
||||
" }\n"
|
||||
" }\n"
|
||||
" \n"
|
||||
"#else //Use binary search\n"
|
||||
" //Find nearest element to left with a lower common prefix\n"
|
||||
" int leftIndex = -1;\n"
|
||||
" {\n"
|
||||
" int lower = 0;\n"
|
||||
" int upper = internalNodeIndex - 1;\n"
|
||||
" \n"
|
||||
" while(lower <= upper)\n"
|
||||
" {\n"
|
||||
" int mid = (lower + upper) / 2;\n"
|
||||
" b3Int64 midPrefix = commonPrefixes[mid];\n"
|
||||
" int midPrefixLength = commonPrefixLengths[mid];\n"
|
||||
" \n"
|
||||
" int nodeMidSharedPrefixLength = getSharedPrefixLength(nodePrefix, nodePrefixLength, midPrefix, midPrefixLength);\n"
|
||||
" if(nodeMidSharedPrefixLength < nodePrefixLength) \n"
|
||||
" {\n"
|
||||
" int right = mid + 1;\n"
|
||||
" if(right < internalNodeIndex)\n"
|
||||
" {\n"
|
||||
" b3Int64 rightPrefix = commonPrefixes[right];\n"
|
||||
" int rightPrefixLength = commonPrefixLengths[right];\n"
|
||||
" \n"
|
||||
" int nodeRightSharedPrefixLength = getSharedPrefixLength(nodePrefix, nodePrefixLength, rightPrefix, rightPrefixLength);\n"
|
||||
" if(nodeRightSharedPrefixLength < nodePrefixLength) \n"
|
||||
" {\n"
|
||||
" lower = right;\n"
|
||||
" leftIndex = right;\n"
|
||||
" }\n"
|
||||
" else \n"
|
||||
" {\n"
|
||||
" leftIndex = mid;\n"
|
||||
" break;\n"
|
||||
" }\n"
|
||||
" }\n"
|
||||
" else \n"
|
||||
" {\n"
|
||||
" leftIndex = mid;\n"
|
||||
" break;\n"
|
||||
" }\n"
|
||||
" }\n"
|
||||
" else upper = mid - 1;\n"
|
||||
" }\n"
|
||||
" }\n"
|
||||
" \n"
|
||||
" //Find nearest element to right with a lower common prefix\n"
|
||||
" int rightIndex = -1;\n"
|
||||
" {\n"
|
||||
" int lower = internalNodeIndex + 1;\n"
|
||||
" int upper = numInternalNodes - 1;\n"
|
||||
" \n"
|
||||
" while(lower <= upper)\n"
|
||||
" {\n"
|
||||
" int mid = (lower + upper) / 2;\n"
|
||||
" b3Int64 midPrefix = commonPrefixes[mid];\n"
|
||||
" int midPrefixLength = commonPrefixLengths[mid];\n"
|
||||
" \n"
|
||||
" int nodeMidSharedPrefixLength = getSharedPrefixLength(nodePrefix, nodePrefixLength, midPrefix, midPrefixLength);\n"
|
||||
" if(nodeMidSharedPrefixLength < nodePrefixLength) \n"
|
||||
" {\n"
|
||||
" int left = mid - 1;\n"
|
||||
" if(left > internalNodeIndex)\n"
|
||||
" {\n"
|
||||
" b3Int64 leftPrefix = commonPrefixes[left];\n"
|
||||
" int leftPrefixLength = commonPrefixLengths[left];\n"
|
||||
" \n"
|
||||
" int nodeLeftSharedPrefixLength = getSharedPrefixLength(nodePrefix, nodePrefixLength, leftPrefix, leftPrefixLength);\n"
|
||||
" if(nodeLeftSharedPrefixLength < nodePrefixLength) \n"
|
||||
" {\n"
|
||||
" upper = left;\n"
|
||||
" rightIndex = left;\n"
|
||||
" }\n"
|
||||
" else \n"
|
||||
" {\n"
|
||||
" rightIndex = mid;\n"
|
||||
" break;\n"
|
||||
" }\n"
|
||||
" }\n"
|
||||
" else \n"
|
||||
" {\n"
|
||||
" rightIndex = mid;\n"
|
||||
" break;\n"
|
||||
" }\n"
|
||||
" }\n"
|
||||
" else lower = mid + 1;\n"
|
||||
" }\n"
|
||||
" }\n"
|
||||
"#endif\n"
|
||||
" \n"
|
||||
" TEMP_out_leftLowerPrefix[internalNodeIndex] = leftIndex;\n"
|
||||
" TEMP_out_rightLowerPrefix[internalNodeIndex] = rightIndex;\n"
|
||||
" TEMP_spl_left[internalNodeIndex] = (leftIndex != -1) ? getSharedPrefixLength(nodePrefix, nodePrefixLength, commonPrefixes[leftIndex], commonPrefixLengths[leftIndex]) : -1;\n"
|
||||
" TEMP_spl_right[internalNodeIndex] = (rightIndex != -1) ? getSharedPrefixLength(nodePrefix, nodePrefixLength, commonPrefixes[rightIndex], commonPrefixLengths[rightIndex]) : -1;\n"
|
||||
" \n"
|
||||
" //Select parent\n"
|
||||
" {\n"
|
||||
" int leftPrefixLength = (leftIndex != -1) ? commonPrefixLengths[leftIndex] : B3_PLBVH_INVALID_COMMON_PREFIX;\n"
|
||||
" int rightPrefixLength = (rightIndex != -1) ? commonPrefixLengths[rightIndex] : B3_PLBVH_INVALID_COMMON_PREFIX;\n"
|
||||
" \n"
|
||||
" int isLeftHigherPrefixLength = (leftPrefixLength > rightPrefixLength);\n"
|
||||
" \n"
|
||||
" if(leftPrefixLength == B3_PLBVH_INVALID_COMMON_PREFIX) isLeftHigherPrefixLength = false;\n"
|
||||
" else if(rightPrefixLength == B3_PLBVH_INVALID_COMMON_PREFIX) isLeftHigherPrefixLength = true;\n"
|
||||
" \n"
|
||||
" int parentNodeIndex = (isLeftHigherPrefixLength) ? leftIndex : rightIndex;\n"
|
||||
" \n"
|
||||
" int isRootNode = (leftIndex == -1 && rightIndex == -1);\n"
|
||||
" out_internalNodeParentNodes[internalNodeIndex] = (!isRootNode) ? parentNodeIndex : B3_PLBVH_ROOT_NODE_MARKER;\n"
|
||||
" \n"
|
||||
" int isLeaf = 0;\n"
|
||||
" if(!isRootNode)\n"
|
||||
" {\n"
|
||||
" int isRightChild = (isLeftHigherPrefixLength); //If the left node is the parent, then this node is its right child and vice versa\n"
|
||||
" \n"
|
||||
" //out_childNodesAsInt[0] == int2.x == left child\n"
|
||||
" //out_childNodesAsInt[1] == int2.y == right child\n"
|
||||
" __global int* out_childNodesAsInt = (__global int*)(&out_childNodes[parentNodeIndex]);\n"
|
||||
" out_childNodesAsInt[isRightChild] = getIndexWithInternalNodeMarkerSet(isLeaf, internalNodeIndex);\n"
|
||||
" }\n"
|
||||
" else *out_rootNodeIndex = getIndexWithInternalNodeMarkerSet(isLeaf, internalNodeIndex);\n"
|
||||
" }\n"
|
||||
"}\n"
|
||||
"__kernel void findDistanceFromRoot(__global int* rootNodeIndex, __global int* internalNodeParentNodes,\n"
|
||||
" __global int* out_maxDistanceFromRoot, __global int* out_distanceFromRoot, int numInternalNodes)\n"
|
||||
"{\n"
|
||||
" if( get_global_id(0) == 0 ) atomic_xchg(out_maxDistanceFromRoot, 0);\n"
|
||||
" int internalNodeIndex = get_global_id(0);\n"
|
||||
" if(internalNodeIndex >= numInternalNodes) return;\n"
|
||||
" \n"
|
||||
" //\n"
|
||||
" int distanceFromRoot = 0;\n"
|
||||
" {\n"
|
||||
" int parentIndex = internalNodeParentNodes[internalNodeIndex];\n"
|
||||
" while(parentIndex != B3_PLBVH_ROOT_NODE_MARKER)\n"
|
||||
" {\n"
|
||||
" parentIndex = internalNodeParentNodes[parentIndex];\n"
|
||||
" ++distanceFromRoot;\n"
|
||||
" }\n"
|
||||
" }\n"
|
||||
" out_distanceFromRoot[internalNodeIndex] = distanceFromRoot;\n"
|
||||
" \n"
|
||||
" //\n"
|
||||
" __local int localMaxDistanceFromRoot;\n"
|
||||
" if( get_local_id(0) == 0 ) localMaxDistanceFromRoot = 0;\n"
|
||||
" barrier(CLK_LOCAL_MEM_FENCE);\n"
|
||||
" \n"
|
||||
" atomic_max(&localMaxDistanceFromRoot, distanceFromRoot);\n"
|
||||
" barrier(CLK_LOCAL_MEM_FENCE);\n"
|
||||
" \n"
|
||||
" if( get_local_id(0) == 0 ) atomic_max(out_maxDistanceFromRoot, localMaxDistanceFromRoot);\n"
|
||||
"}\n"
|
||||
"__kernel void buildBinaryRadixTreeAabbsRecursive(__global int* distanceFromRoot, __global SortDataCL* mortonCodesAndAabbIndices,\n"
|
||||
" __global int2* childNodes,\n"
|
||||
" __global b3AabbCL* leafNodeAabbs, __global b3AabbCL* internalNodeAabbs,\n"
|
||||
" int maxDistanceFromRoot, int processedDistance, int numInternalNodes)\n"
|
||||
"{\n"
|
||||
" int internalNodeIndex = get_global_id(0);\n"
|
||||
" if(internalNodeIndex >= numInternalNodes) return;\n"
|
||||
" \n"
|
||||
" int distance = distanceFromRoot[internalNodeIndex];\n"
|
||||
" \n"
|
||||
" if(distance == processedDistance)\n"
|
||||
" {\n"
|
||||
" int leftChildIndex = childNodes[internalNodeIndex].x;\n"
|
||||
" int rightChildIndex = childNodes[internalNodeIndex].y;\n"
|
||||
" \n"
|
||||
" int isLeftChildLeaf = isLeafNode(leftChildIndex);\n"
|
||||
" int isRightChildLeaf = isLeafNode(rightChildIndex);\n"
|
||||
@@ -718,16 +706,4 @@ static const char* parallelLinearBvhCL= \
|
||||
" internalNodeAabbs[internalNodeIndex] = mergedAabb;\n"
|
||||
" }\n"
|
||||
"}\n"
|
||||
"__kernel void convertChildNodeFormat(__global int* leftChildNodes, __global int* rightChildNodes, \n"
|
||||
" __global int2* out_childNodes, int numInternalNodes)\n"
|
||||
"{\n"
|
||||
" int internalNodeIndex = get_global_id(0);\n"
|
||||
" if (internalNodeIndex >= numInternalNodes) return;\n"
|
||||
" \n"
|
||||
" int2 childNodesIndices;\n"
|
||||
" childNodesIndices.x = leftChildNodes[internalNodeIndex];\n"
|
||||
" childNodesIndices.y = rightChildNodes[internalNodeIndex];\n"
|
||||
" \n"
|
||||
" out_childNodes[internalNodeIndex] = childNodesIndices;\n"
|
||||
"}\n"
|
||||
;
|
||||
|
||||
Reference in New Issue
Block a user