diff --git a/CMakeLists.txt b/CMakeLists.txt index e63a444..6e1477d 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -166,6 +166,11 @@ add_subdirectory(cuBQL) if (NOT CUBQL_IS_SUBPROJECT) add_subdirectory(samples) + option(CUBQL_BUILD_TESTS "Build cuBQL correctness tests" OFF) + if (CUBQL_BUILD_TESTS) + enable_testing() + add_subdirectory(tests) + endif() endif() #add_subdirectory(testing) diff --git a/cuBQL/builder/cuda.h b/cuBQL/builder/cuda.h index a8f9e7f..b30a106 100644 --- a/cuBQL/builder/cuda.h +++ b/cuBQL/builder/cuda.h @@ -10,6 +10,7 @@ #endif # include "cuBQL/math/box.h" # include +# include namespace cuBQL { // ------------------------------------------------------------------ @@ -179,6 +180,18 @@ namespace cuBQL { BuildConfig buildConfig, cudaStream_t s=0, GpuMemoryResource &memResource=defaultGpuMemResource()); + + /*! Select uint32_t for compact Morton keys, or uint64_t for higher + precision. Requires finite, nonempty primitive bounds. */ + template + typename std::enable_if::value + || std::is_same::value,void>::type + radixBuilder(BinaryBVH &bvh, + const box_t *boxes, + uint32_t numPrims, + BuildConfig buildConfig, + cudaStream_t s=0, + GpuMemoryResource &memResource=defaultGpuMemResource()); // ------------------------------------------------------------------ /*! fast radix/morton builder with automatic rebinning where diff --git a/cuBQL/builder/cuda/instantiate_builders.cu b/cuBQL/builder/cuda/instantiate_builders.cu index 7d9c667..6b3f9a4 100644 --- a/cuBQL/builder/cuda/instantiate_builders.cu +++ b/cuBQL/builder/cuda/instantiate_builders.cu @@ -36,6 +36,20 @@ cudaStream_t s, \ GpuMemoryResource &mem_resource); \ template \ + void radixBuilder(BinaryBVH &bvh, \ + const box_t *boxes, \ + uint32_t numBoxes, \ + BuildConfig buildConfig, \ + cudaStream_t s, \ + GpuMemoryResource &mem_resource); \ + template \ + void radixBuilder(BinaryBVH &bvh, \ + const box_t *boxes, \ + uint32_t numBoxes, \ + BuildConfig buildConfig, \ + cudaStream_t s, \ + GpuMemoryResource &mem_resource); \ + template \ void rebinRadixBuilder(BinaryBVH &bvh, \ const box_t *boxes, \ uint32_t numBoxes, \ diff --git a/cuBQL/builder/cuda/radix.h b/cuBQL/builder/cuda/radix.h index 1a4c74a..471adf4 100644 --- a/cuBQL/builder/cuda/radix.h +++ b/cuBQL/builder/cuda/radix.h @@ -4,6 +4,7 @@ #pragma once #include "cuBQL/builder/cuda/sm_builder.h" +#include namespace cuBQL { namespace radixBuilder_impl { @@ -11,8 +12,6 @@ namespace cuBQL { using gpuBuilder_impl::_ALLOC; using gpuBuilder_impl::_FREE; - template struct Quantizer; - template struct numMortonBits; #if 0 template<> struct numMortonBits<2> { enum { value = 31 }; }; @@ -24,9 +23,27 @@ namespace cuBQL { template<> struct numMortonBits<3> { enum { value = 10 }; }; template<> struct numMortonBits<4> { enum { value = 7 }; }; #endif + + template + struct MortonKeyTraits { + static_assert(std::is_same::value + || std::is_same::value, + "Morton keys must be uint32_t or uint64_t"); + static_assert(D >= 2 && D <= 4, + "Morton keys require two, three, or four dimensions"); + enum { + numBits = (8*sizeof(MortonKey)-1)/D, + meaningfulBits = D*numBits, + sortEndBit = sizeof(MortonKey) == sizeof(uint32_t) + ? meaningfulBits : 64 + }; + }; + + template::value> + struct Quantizer; - template - struct Quantizer { + template + struct Quantizer { using vec_t = cuBQL::vec_t; using box_t = cuBQL::box_t; @@ -35,7 +52,7 @@ namespace cuBQL { quantizeBias = centBounds.lower; quantizeScale - = vec_t(1u<::value) + = vec_t(1u<; vec_ui cell = vec_ui((P-quantizeBias)*quantizeScale); - cell = min(cell,vec_ui(uint32_t(((1u<::value)-1)))); + cell = min(cell,vec_ui(uint32_t(((1u< - struct Quantizer { + template + struct Quantizer { using vec_t = cuBQL::vec_t; using box_t = cuBQL::box_t; @@ -66,7 +83,7 @@ namespace cuBQL { using vec_ui = cuBQL::vec_t; vec_ui cell = vec_ui((P-quantizeBias)*quantizeScale); - cell = min(cell,vec_ui(uint32_t(((1u<::value)-1)))); + cell = min(cell,vec_ui(uint32_t(((1u<::value) + = vec_t(1u< - struct Quantizer { + template + struct Quantizer { using vec_t = cuBQL::vec_t; using box_t = cuBQL::box_t; @@ -104,7 +121,7 @@ namespace cuBQL { cuBQL::vec_t cell = cuBQL::vec_t(P-quantizeBias); // move all relevant bits to top cell = cell << shlBits; - return cell >> (32-numMortonBits::value); + return cell >> (32-NumBits); } /*! coefficients of `scale*(x-bias)` in the 21-bit fixed-point @@ -115,8 +132,8 @@ namespace cuBQL { int shlBits; }; - template - struct Quantizer { + template + struct Quantizer { using vec_t = cuBQL::vec_t; using box_t = cuBQL::box_t; @@ -132,7 +149,7 @@ namespace cuBQL { cuBQL::vec_t cell = cuBQL::vec_t(P-quantizeBias); // move all relevant bits to top cell = cell << shlBits; - return cuBQL::vec_t(cell >> (64-numMortonBits::value)); + return cuBQL::vec_t(cell >> (64-NumBits)); } /*! coefficients of `scale*(x-bias)` in the 21-bit fixed-point @@ -147,7 +164,8 @@ namespace cuBQL { /*! maintains high-level summary of the build process */ - template + template::value> struct CUBQL_ALIGN(16) BuildState { using vec_t = cuBQL::vec_t;//float,3>; using box_t = cuBQL::box_t;//float,3>; @@ -166,12 +184,13 @@ namespace cuBQL { morton codes */ atomic_box_t a_centBounds; box_t centBounds; - Quantizer quantizer; + Quantizer quantizer; }; - template + template::value> __global__ - void clearBuildState(BuildState *buildState, + void clearBuildState(BuildState *buildState, int numPrims) { if (threadIdx.x != 0) return; @@ -183,14 +202,15 @@ namespace cuBQL { buildState->numNodesAlloced = 0; } - template + template::value> __global__ - void fillBuildState(BuildState *buildState, - const typename BuildState::box_t *prims, + void fillBuildState(BuildState *buildState, + const typename BuildState::box_t *prims, int numPrims) { - using atomic_box_t = typename BuildState::atomic_box_t; - using box_t = typename BuildState::box_t; + using atomic_box_t = typename BuildState::atomic_box_t; + using box_t = typename BuildState::box_t; __shared__ atomic_box_t l_centBounds; if (threadIdx.x == 0) @@ -213,11 +233,12 @@ namespace cuBQL { atomic_grow(buildState->a_centBounds,l_centBounds); } - template + template::value> __global__ - void finishBuildState(BuildState *buildState) + void finishBuildState(BuildState *buildState) { - using ctx_t = BuildState; + using ctx_t = BuildState; using atomic_box_t = typename ctx_t::atomic_box_t; using box_t = typename ctx_t::box_t; @@ -372,24 +393,26 @@ namespace cuBQL { (bitInterleave11(yw) << 1); } - template + template::value> inline __device__ - uint64_t computeMortonCode(typename BuildState::vec_t P, - const Quantizer quantizer) + MortonKey computeMortonCode(typename BuildState::vec_t P, + const Quantizer quantizer) { - return interleaveBits64(quantizer.quantize(P)); + return MortonKey(interleaveBits64(quantizer.quantize(P))); } - template + template::value> __global__ - void computeUnsortedKeysAndPrimIDs(uint64_t *mortonCodes, + void computeUnsortedKeysAndPrimIDs(MortonKey *mortonCodes, uint32_t *primIDs, - BuildState *buildState, - const typename BuildState::box_t *prims, + BuildState *buildState, + const typename BuildState::box_t *prims, int numPrims) { - using atomic_box_t = typename BuildState::atomic_box_t; - using box_t = typename BuildState::box_t; + using atomic_box_t = typename BuildState::atomic_box_t; + using box_t = typename BuildState::box_t; int tid = threadIdx.x + blockIdx.x*blockDim.x; if (tid >= numPrims) return; @@ -404,7 +427,7 @@ namespace cuBQL { primIDs[tid] = primID; mortonCodes[tid] - = computeMortonCode(prim.center(),buildState->quantizer); + = computeMortonCode(prim.center(),buildState->quantizer); } struct TempNode { @@ -430,9 +453,10 @@ namespace cuBQL { }; + template inline __device__ bool findSplit(int &split, - const uint64_t *__restrict__ keys, + const MortonKey *__restrict__ keys, int begin, int end, int maxAllowedLeafSize) { @@ -468,9 +492,10 @@ namespace cuBQL { return true; } - template + template::value> __global__ - void initNodes(BuildState *buildState, + void initNodes(BuildState *buildState, TempNode *nodes, int numValidPrims) { @@ -485,17 +510,18 @@ namespace cuBQL { nodes[1] = n1; } - template + template::value> __global__ - void createNodes(BuildState *buildState, + void createNodes(BuildState *buildState, int leafThreshold, int maxAllowedLeafSize, TempNode *nodes, int begin, int end, - const uint64_t *keys) + const MortonKey *keys) { - using atomic_box_t = typename BuildState::atomic_box_t; - using box_t = typename BuildState::box_t; + using atomic_box_t = typename BuildState::atomic_box_t; + using box_t = typename BuildState::box_t; __shared__ int l_allocOffset; @@ -575,14 +601,25 @@ namespace cuBQL { finalNodes[tid].admin.offsetAndCountBits = node.admin.offsetAndCountBits; } - template + template::value, + int SortEndBit=8*sizeof(MortonKey)> void build(BinaryBVH &bvh, - const typename BuildState::box_t *boxes, + const box_t *boxes, uint32_t numPrims, BuildConfig buildConfig, cudaStream_t s, GpuMemoryResource &memResource) { + static_assert(std::is_same::value + || std::is_same::value, + "Morton keys must be uint32_t or uint64_t"); + static_assert(NumBits > 0 && NumBits <= 31 + && D*NumBits <= 8*sizeof(MortonKey), + "Morton key precision must fit in the key type"); + static_assert(SortEndBit >= D*NumBits + && SortEndBit <= 8*sizeof(MortonKey), + "Morton sort range must include every populated key bit"); const int makeLeafThreshold = (buildConfig.makeLeafThreshold > 0) ? std::min(buildConfig.makeLeafThreshold,buildConfig.maxAllowedLeafSize) @@ -595,7 +632,7 @@ namespace cuBQL { /* step 1.1, init build state; in particular, clear the shared centbounds we need to atomically grow centroid bounds in next step */ - BuildState *d_buildState = 0; + BuildState *d_buildState = 0; _ALLOC(d_buildState,1,s,memResource); clearBuildState<<<32,1,0,s>>> (d_buildState,numPrims); @@ -609,7 +646,7 @@ namespace cuBQL { finishBuildState<<<32,1,0,s>>> (d_buildState); - static BuildState *h_buildState = 0; + static BuildState *h_buildState = 0; if (!h_buildState) CUBQL_CUDA_CALL(MallocHost((void**)&h_buildState, sizeof(*h_buildState))); @@ -634,7 +671,7 @@ namespace cuBQL { /* 2.1, allocate mem for _unsorted_ prim IDs and morton codes, then compute initial primID array (will already exclude prims that are invalid) and (unsorted) morton code array */ - uint64_t *d_primKeys_unsorted; + MortonKey *d_primKeys_unsorted; uint32_t *d_primIDs_unsorted; _ALLOC(d_primKeys_unsorted,numPrims,s,memResource); _ALLOC(d_primIDs_unsorted,numPrims,s,memResource); @@ -646,7 +683,7 @@ namespace cuBQL { /* 2.2: ask cub radix sorter for how much temp mem it needs, and allocate */ size_t cub_tempMemSize; - uint64_t *d_primKeys_sorted = 0; + MortonKey *d_primKeys_sorted = 0; uint32_t *d_primIDs_inMortonOrder = 0; // with tempMem ptr null this won't do anything but return reqd // temp size*/ @@ -657,7 +694,7 @@ namespace cuBQL { /*keys out:*/ d_primKeys_sorted, /*values in:*/ d_primIDs_unsorted, /*values out:*/d_primIDs_inMortonOrder, - numValidPrims,0,64,s); + numValidPrims,0,SortEndBit,s); // 2.3: allocate temp mem and output arrays void *d_tempMem = 0; @@ -673,7 +710,7 @@ namespace cuBQL { /*keys out:*/ d_primKeys_sorted, /*values in:*/ d_primIDs_unsorted, /*values out:*/d_primIDs_inMortonOrder, - numValidPrims,0,64,s); + numValidPrims,0,SortEndBit,s); rc = rc; // 2.5 - cleanup after sort: no longer need tempmem, or unsorted inputs _FREE(d_primKeys_unsorted,s,memResource); @@ -757,6 +794,21 @@ namespace cuBQL { { radixBuilder_impl::build(bvh,boxes,numPrims,buildConfig,s,memResource); } + + template + typename std::enable_if::value + || std::is_same::value,void>::type + radixBuilder(cuBQL::BinaryBVH &bvh, + const cuBQL::box_t *boxes, + uint32_t numPrims, + cuBQL::BuildConfig buildConfig, + cudaStream_t s, + cuBQL::GpuMemoryResource &memResource) + { + using traits = radixBuilder_impl::MortonKeyTraits; + radixBuilder_impl::build + (bvh,boxes,numPrims,buildConfig,s,memResource); + } } } diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt new file mode 100644 index 0000000..74fe1ed --- /dev/null +++ b/tests/CMakeLists.txt @@ -0,0 +1,11 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +if (CUBQL_HAVE_CUDA) + add_executable(cubql_morton_sort_range morton_sort_range.cu) + target_link_libraries(cubql_morton_sort_range PRIVATE cuBQL) + target_compile_definitions(cubql_morton_sort_range PRIVATE CUBQL_HAVE_CUDA=1) + set_target_properties(cubql_morton_sort_range PROPERTIES + CUDA_STANDARD 17 CUDA_STANDARD_REQUIRED ON) + add_test(NAME cubql_morton_sort_range COMMAND cubql_morton_sort_range) +endif() diff --git a/tests/data/eda_layout.h b/tests/data/eda_layout.h new file mode 100644 index 0000000..5ce42e7 --- /dev/null +++ b/tests/data/eda_layout.h @@ -0,0 +1,231 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "cuBQL/math/box.h" +#include +#include +#include +#include +#include +#include +#include +#include + +namespace eda_layout { + +enum class Pattern { Layout, Motif }; + +struct Rectangle { int x0, y0, x1, y1; }; + +// Original integer geometry: fingers, routing rails, vertical routes, contacts, +// and a concentric pair that deliberately produces equal centroid keys. +inline constexpr Rectangle motif[] = { + { 8, 12, 12, 36}, {20, 12, 24, 36}, {32, 12, 36, 36}, {44, 12, 48, 36}, + {56, 12, 60, 36}, {68, 12, 72, 36}, {80, 12, 84, 36}, {92, 12, 96, 36}, + { 8, 72, 12, 96}, {20, 72, 24, 96}, {32, 72, 36, 96}, {44, 72, 48, 96}, + {56, 72, 60, 96}, {68, 72, 72, 96}, {80, 72, 84, 96}, {92, 72, 96, 96}, + { 4, 4,100, 8}, { 4, 40,100, 44}, { 4, 64,100, 68}, { 4,100,100,104}, + { 2, 4, 4,104}, {104, 4,106,104}, {16, 44, 18, 64}, {88, 44, 90, 64}, + { 8, 46, 14, 52}, {26, 46, 32, 52}, {44, 46, 50, 52}, {62, 46, 68, 52}, + {80, 46, 86, 52}, {98, 46,104, 52}, {48, 54, 56, 62}, {50, 56, 54, 60} +}; +inline constexpr uint32_t rectanglesPerTile = sizeof(motif)/sizeof(motif[0]); +inline constexpr int tilePitch = 128; + +inline Rectangle layoutRectangle(uint32_t x, uint32_t y, uint32_t r) +{ + const uint32_t groupX = x/3; + const uint32_t groupY = y/2; + const uint32_t selection = (5*groupX+7*groupY+3*(groupX^groupY))%19; + uint32_t kind = selection < 8 ? 0 : selection < 13 ? 1 : selection < 18 ? 2 : 3; + if ((x%11 == 5 && y%5 != 1) || (y%9 == 4 && x%7 != 2)) kind = 3; + Rectangle rect; + if (kind == 0) { + if (r < 16) { + const int x0 = 8+14*int(r%8); + const int y0 = 8+72*int(r/8); + rect = {x0,y0,x0+10,y0+34}; + } else if (r < 20) { + const int i = int(r-16); + const int y0 = 4+72*(i/2)+38*(i%2); + rect = {2,y0,122,y0+4}; + } else if (r < 28) { + const int x0 = 10+14*int(r-20); + rect = {x0,54,x0+8,62}; + } else if (r < 30) { + const int x0 = 2+118*int(r-28); + rect = {x0,4,x0+4,118}; + } else { + rect = r == 30 ? Rectangle{56,64,72,74} : Rectangle{58,66,70,72}; + } + } else if (kind == 1) { + if (r < 16) { + const int group = int(r/4); + const int x0 = 8+60*(group%2)+12*int(r%4); + const int y0 = 8+60*(group/2); + rect = {x0,y0,x0+8,y0+28}; + } else if (r < 24) { + const int i = int(r-16); + const int x0 = 4+60*((i/2)%2); + const int y0 = 4+60*(i/4)+32*(i%2); + rect = {x0,y0,x0+52,y0+4}; + } else if (r < 28) { + const int i = int(r-24); + const int x0 = 4+60*(i%2); + const int y0 = 4+60*(i/2); + rect = {x0,y0,x0+4,y0+36}; + } else if (r < 30) { + const int y0 = 48+72*int(r-28); + rect = {2,y0,124,y0+4}; + } else { + rect = r == 30 ? Rectangle{56,48,72,60} : Rectangle{58,50,70,58}; + } + } else if (kind == 2) { + if (r < 16) { + const int i = int(r/2); + const int x0 = 2+30*(i%4); + const int y0 = 8+60*(i/4); + rect = r%2 == 0 ? Rectangle{x0,y0,x0+24,y0+10} + : Rectangle{x0+16,y0+10,x0+24,y0+38}; + } else if (r < 24) { + const int i = int(r-16); + const int x0 = 8+30*(i%4); + const int y0 = 28+60*(i/4); + rect = {x0,y0,x0+8,y0+8}; + } else if (r < 28) { + const int i = int(r-24); + const int y0 = 4+56*(i/2)+4*(i%2); + rect = {4,y0,124,y0+3}; + } else if (r < 30) { + const int x0 = 4+116*int(r-28); + rect = {x0,4,x0+4,120}; + } else { + rect = r == 30 ? Rectangle{56,50,72,58} : Rectangle{58,52,70,56}; + } + } else { + if (r < 16) { + const int i = int(r/2); + const int x0 = 8+72*(i%2); + const int y0 = 8+28*(i/2); + rect = r%2 == 0 ? Rectangle{x0,y0,x0+36,y0+6} + : Rectangle{x0+28,y0+6,x0+36,y0+20}; + } else if (r < 28) { + const int i = int(r-16); + const int x0 = 12+12*(i%3)+72*(i/6); + const int y0 = 30+32*((i%6)/3); + rect = {x0,y0,x0+6,y0+6}; + } else if (r < 30) { + const int x0 = 2+120*int(r-28); + rect = {x0,4,x0+4,124}; + } else { + rect = r == 30 ? Rectangle{58,58,70,70} : Rectangle{60,60,68,68}; + } + } + if ((groupX+groupY)%2) { + const int x0 = tilePitch-rect.x1; + rect.x1 = tilePitch-rect.x0; + rect.x0 = x0; + } + if ((x/5+y/3)%2) { + const int y0 = tilePitch-rect.y1; + rect.y1 = tilePitch-rect.y0; + rect.y0 = y0; + } + return rect; +} + +template +std::vector> makeBoxes(uint32_t tilesX, uint32_t tilesY, + Pattern pattern=Pattern::Layout) +{ + static_assert(D >= 2 && D <= 4); + if (tilesX == 0 || tilesY == 0 || tilesX > 131072 || tilesY > 131072) + throw std::invalid_argument("Synthetic layout requires 1 to 131072 tiles per axis"); + if (uint64_t(tilesX)*tilesY*rectanglesPerTile > uint64_t(std::numeric_limits::max())/2) + throw std::invalid_argument("Synthetic layout exceeds the radix builder node limit"); + std::vector> boxes(std::size_t(tilesX)*tilesY*rectanglesPerTile); + std::size_t index = 0; + for (uint32_t y=0;y +uint64_t fingerprint(const std::vector> &boxes) +{ + uint64_t hash = UINT64_C(14695981039346656037); + for (const auto &box : boxes) + for (int axis=0;axis> (8*byte)) & 255u; + hash *= UINT64_C(1099511628211); + } + } + return hash; +} + +inline void writeSvg(const char *path, uint32_t tilesX=16, uint32_t tilesY=8, + Pattern pattern=Pattern::Layout) +{ + if (!path || !*path) + throw std::invalid_argument("SVG output path must not be empty"); + if (tilesX == 0 || tilesY == 0 || tilesX > 32 || tilesY > 32) + throw std::invalid_argument("SVG preview requires 1 to 32 tiles per axis"); + const auto boxes = makeBoxes(tilesX,tilesY,pattern); + auto bounds = boxes.front(); + for (const auto &box : boxes) + bounds.grow(box); + const int left = int(bounds.lower[0])-18; + const int top = int(bounds.lower[1])-46; + const int width = std::max(340,int(bounds.upper[0]-bounds.lower[0])+36); + const int height = int(bounds.upper[1]-bounds.lower[1])+64; + std::ofstream out; + out.exceptions(std::ios::failbit | std::ios::badbit); + try { + out.open(path); + out << "\n" + << "Original synthetic layout\n" + << "\n" + << "" + << "Original synthetic layout: " << boxes.size() << " rectangles\n"; + for (std::size_t i=0;i= 16 && r < 24) || r == 31 ? "#52c7bd" : "#e7be62"; + out << "\n"; + } + out << "\n"; + out.close(); + } catch (const std::ios_base::failure &) { + throw std::runtime_error(std::string("Failed to write SVG preview: ")+path); + } +} + +} // namespace eda_layout diff --git a/tests/data/sample.jpg b/tests/data/sample.jpg new file mode 100644 index 0000000..16623a5 Binary files /dev/null and b/tests/data/sample.jpg differ diff --git a/tests/morton_sort_range.cu b/tests/morton_sort_range.cu new file mode 100644 index 0000000..b03e5b3 --- /dev/null +++ b/tests/morton_sort_range.cu @@ -0,0 +1,636 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#include "cuBQL/bvh.h" +#include "cuBQL/builder/cuda.h" +#include "cuBQL/builder/cuda/radix.h" +#include "data/eda_layout.h" + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace { + +void check(cudaError_t result) +{ + if (result != cudaSuccess) + throw std::runtime_error(cudaGetErrorString(result)); +} + +template +struct DeviceBuffer { + T *data = nullptr; + explicit DeviceBuffer(size_t count) { + check(cudaMalloc(reinterpret_cast(&data),count*sizeof(T))); + } + ~DeviceBuffer() { cudaFree(data); } + DeviceBuffer(const DeviceBuffer &) = delete; + DeviceBuffer &operator=(const DeviceBuffer &) = delete; +}; + +cudaError_t freeTreeBuffer(void *ptr) +{ +#if CUDART_VERSION >= 11020 + return cudaFreeAsync(ptr,0); +#else + return cudaFree(ptr); +#endif +} + +template +struct BuiltTree { + cuBQL::BinaryBVH bvh; + BuiltTree() = default; + BuiltTree(const BuiltTree &) = delete; + BuiltTree &operator=(const BuiltTree &) = delete; + ~BuiltTree() noexcept { + if (bvh.nodes) freeTreeBuffer(bvh.nodes); + if (bvh.primIDs) freeTreeBuffer(bvh.primIDs); + } + void release() { + if (bvh.nodes) { check(freeTreeBuffer(bvh.nodes)); bvh.nodes = nullptr; } + if (bvh.primIDs) { check(freeTreeBuffer(bvh.primIDs)); bvh.primIDs = nullptr; } + } +}; + +struct Event { + cudaEvent_t handle; + Event() { check(cudaEventCreate(&handle)); } + ~Event() noexcept { cudaEventDestroy(handle); } + Event(const Event &) = delete; + Event &operator=(const Event &) = delete; +}; + +template +struct HostTree { + std::vector::node_t> nodes; + std::vector primIDs; +}; + +template +HostTree snapshot(const cuBQL::BinaryBVH &bvh) +{ + HostTree result; + result.nodes.resize(bvh.numNodes); + result.primIDs.resize(bvh.numPrims); + check(cudaMemcpy(result.nodes.data(),bvh.nodes, + result.nodes.size()*sizeof(result.nodes[0]),cudaMemcpyDeviceToHost)); + check(cudaMemcpy(result.primIDs.data(),bvh.primIDs, + result.primIDs.size()*sizeof(uint32_t),cudaMemcpyDeviceToHost)); + return result; +} + +template +__global__ void auditKeyBits(const MortonKey *keys, uint32_t count, + int bits, unsigned *flags) +{ + const uint32_t i = blockIdx.x*blockDim.x+threadIdx.x; + if (i >= count) return; + if ((keys[i] >> bits) != 0) atomicOr(flags,1u); + if constexpr (sizeof(MortonKey) == 8) + if ((keys[i] >> 32) != 0) atomicOr(flags,2u); +} + +template +std::vector auditGeneratedKeys(const cuBQL::box_t *boxes, + uint32_t count, bool requireHighBits=false) +{ + namespace radix = cuBQL::radixBuilder_impl; + constexpr int bits = D*NumBits; + static_assert(bits > 0 && bits < 8*sizeof(MortonKey)); + DeviceBuffer> state(1); + DeviceBuffer keys(count); + DeviceBuffer ids(count); + DeviceBuffer flags(1); + check(cudaMemset(flags.data,0,sizeof(unsigned))); + radix::clearBuildState<<<32,1>>>(state.data,count); + radix::fillBuildState<<<(count+1023)/1024,1024>>>(state.data,boxes,count); + radix::finishBuildState<<<32,1>>>(state.data); + radix::computeUnsortedKeysAndPrimIDs<<<(count+1023)/1024,1024>>> + (keys.data,ids.data,state.data,boxes,count); + auditKeyBits<<<(count+255)/256,256>>>(keys.data,count,bits,flags.data); + check(cudaGetLastError()); + unsigned flagValue = 0; + check(cudaMemcpy(&flagValue,flags.data,sizeof(flagValue),cudaMemcpyDeviceToHost)); + if (flagValue & 1u) throw std::runtime_error("Morton key exceeds populated bit range"); + if (requireHighBits && !(flagValue & 2u)) + throw std::runtime_error("64-bit Morton keys do not populate their upper 32 bits"); + std::vector result(count); + check(cudaMemcpy(result.data(),keys.data,count*sizeof(MortonKey),cudaMemcpyDeviceToHost)); + return result; +} + +template +void validateTree(const HostTree &tree, const std::vector> &boxes) +{ + if (tree.nodes.empty() || tree.primIDs.size() != boxes.size()) + throw std::runtime_error("BVH size differs from the input"); + std::vector seen(boxes.size(),0), covered(boxes.size(),0); + for (uint32_t id : tree.primIDs) + if (id >= boxes.size() || seen[id]++) + throw std::runtime_error("primitive IDs are not a permutation of the input"); + std::vector visited(tree.nodes.size(),0); + std::vector> bounds(tree.nodes.size()); + std::vector> pending{{0,false}}; + size_t numVisited = 0; + while (!pending.empty()) { + const auto entry = pending.back(); pending.pop_back(); + const uint32_t id = entry.first; + if (id >= tree.nodes.size()) throw std::runtime_error("BVH child exceeds node array"); + const auto &node = tree.nodes[id]; + if (!entry.second) { + if (visited[id]++) throw std::runtime_error("BVH contains a cycle or shared child"); + ++numVisited; + pending.emplace_back(id,true); + if (node.admin.count) { + const uint64_t end = uint64_t(node.admin.offset)+node.admin.count; + if (end > tree.primIDs.size()) throw std::runtime_error("BVH leaf exceeds primitive array"); + for (uint64_t i=node.admin.offset;i +void compareTrees(const HostTree &a, const HostTree &b, size_t count) +{ + if (a.nodes.size() != b.nodes.size() || a.primIDs != b.primIDs || + a.primIDs.size() != count) + throw std::runtime_error("BVH size or sorted primitive IDs differ"); + std::vector seen(count,0), covered(count,0); + for (uint32_t id : a.primIDs) { + if (id >= count || seen[id]++) + throw std::runtime_error("primitive IDs are not a permutation of the input"); + } + std::vector> pending{{0,0}}; + size_t visited = 0; + while (!pending.empty()) { + const auto pair = pending.back(); pending.pop_back(); + if (pair.first >= a.nodes.size() || pair.second >= b.nodes.size() || + ++visited > a.nodes.size()) + throw std::runtime_error("invalid BVH topology"); + const auto &left = a.nodes[pair.first]; + const auto &right = b.nodes[pair.second]; + if (left.admin.count != right.admin.count) + throw std::runtime_error("leaf counts differ"); + for (int axis=0;axis count) + throw std::runtime_error("leaf ranges differ or exceed primitive array"); + for (uint64_t i=left.admin.offset;i &a, + const cuBQL::box_t &b) +{ + return a.lower[0] < b.upper[0] && b.lower[0] < a.upper[0] && + a.lower[1] < b.upper[1] && b.lower[1] < a.upper[1]; +} + +__global__ void queryAreas(cuBQL::BinaryBVH bvh, + const cuBQL::box_t *boxes, + const cuBQL::box_t *queries, + QueryResult *results, uint32_t count) +{ + const uint32_t q = blockIdx.x*blockDim.x+threadIdx.x; + if (q >= count) return; + uint32_t stack[64], depth = 1; + stack[0] = 0; + QueryResult result{0,0,0}; + while (depth) { + const auto node = bvh.nodes[stack[--depth]]; + if (!positiveOverlap(node.bounds,queries[q])) continue; + if (!node.admin.count) { + if (depth+2 > 64) { result.overflow = 1; break; } + stack[depth++] = uint32_t(node.admin.offset)+1; + stack[depth++] = uint32_t(node.admin.offset); + } else { + for (uint32_t i=0;i> &boxes, + const cuBQL::box_t *deviceBoxes, + const cuBQL::BinaryBVH &a, + const cuBQL::BinaryBVH &b, + const cuBQL::BinaryBVH &c) +{ + using Box = cuBQL::box_t; + const eda_layout::Rectangle windows[] = { + {-140,-140,-100,-90}, {-128,-128,0,0}, {-120,-116,-116,-92}, + {-16,-16,16,16}, {0,0,64,64}, {48,54,56,62}, {112,112,124,124}, + {-256,-256,-200,-200}, {-128,-128,896,384}, {100,4,104,40} + }; + std::vector queries; + for (const auto &r : windows) { + Box box; + box.lower[0] = float(r.x0); box.upper[0] = float(r.x1); + box.lower[1] = float(r.y0); box.upper[1] = float(r.y1); + queries.push_back(box); + } + DeviceBuffer deviceQueries(queries.size()); + DeviceBuffer deviceResults(queries.size()); + check(cudaMemcpy(deviceQueries.data,queries.data(),queries.size()*sizeof(Box),cudaMemcpyHostToDevice)); + for (const auto &tree : {a,b,c}) { + queryAreas<<<1,32>>>(tree,deviceBoxes,deviceQueries.data,deviceResults.data,uint32_t(queries.size())); + check(cudaGetLastError()); + std::vector results(queries.size()); + check(cudaMemcpy(results.data(),deviceResults.data,results.size()*sizeof(QueryResult),cudaMemcpyDeviceToHost)); + for (size_t q=0;q 0 && dy > 0) { area += dx*dy; ++hits; } + } + if (results[q].overflow || results[q].area != double(area) || results[q].hits != hits) + throw std::runtime_error("2D query differs from independent CPU area/overlap oracle"); + } + } + std::cout << "query_oracle=PASS queries=" << queries.size() << " key_modes=3\n"; +} + +template +void auditPrecisionQueries(const std::vector> &boxes, + const cuBQL::box_t *deviceBoxes, + const cuBQL::BinaryBVH &a, + const cuBQL::BinaryBVH &b, + const cuBQL::BinaryBVH &c); + +template +void auditBuilds(const std::vector> &boxes, + const cuBQL::box_t *deviceBoxes, bool pointQueries=false) +{ + namespace radix = cuBQL::radixBuilder_impl; + constexpr int lowBits = radix::MortonKeyTraits::numBits; + constexpr int highBits = radix::MortonKeyTraits::numBits; + static_assert(lowBits == radix::numMortonBits::value); + const uint32_t count = uint32_t(boxes.size()); + const auto legacyKeys = auditGeneratedKeys(deviceBoxes,count); + const auto keys32 = auditGeneratedKeys(deviceBoxes,count); + auditGeneratedKeys(deviceBoxes,count,true); + if (!std::equal(legacyKeys.begin(),legacyKeys.end(),keys32.begin())) + throw std::runtime_error("32-bit keys differ from legacy keys at equal quantization"); + BuiltTree legacy, key32, key64; + cuBQL::BuildConfig config(16); + cuBQL::cuda::radixBuilder(legacy.bvh,deviceBoxes,count,config); + cuBQL::cuda::radixBuilder(key32.bvh,deviceBoxes,count,config); + cuBQL::cuda::radixBuilder(key64.bvh,deviceBoxes,count,config); + check(cudaDeviceSynchronize()); + const auto legacyHost = snapshot(legacy.bvh); + compareTrees(legacyHost,snapshot(key32.bvh),boxes.size()); + validateTree(legacyHost,boxes); + validateTree(snapshot(key64.bvh),boxes); + if constexpr (std::is_same::value && D == 2) + if (!pointQueries) auditQueries(boxes,deviceBoxes,legacy.bvh,key32.bvh,key64.bvh); + if constexpr (D == 2) + if (pointQueries) auditPrecisionQueries(boxes,deviceBoxes,legacy.bvh,key32.bvh,key64.bvh); + legacy.release(); + key32.release(); + key64.release(); +} + +const char *patternName(eda_layout::Pattern pattern) +{ + return pattern == eda_layout::Pattern::Layout ? "layout" : "motif"; +} + +template +void runCase(const char *name, eda_layout::Pattern pattern) +{ + const auto boxes = eda_layout::makeBoxes(8,4,pattern); + DeviceBuffer> deviceBoxes(boxes.size()); + check(cudaMemcpy(deviceBoxes.data,boxes.data(),boxes.size()*sizeof(boxes[0]),cudaMemcpyHostToDevice)); + auditBuilds(boxes,deviceBoxes.data); + std::cout << "geometry=" << patternName(pattern) << " case=" << name + << " boxes=" << boxes.size() << " geometry_fnv64=" + << std::hex << eda_layout::fingerprint(boxes) << std::dec + << " key_modes=3 status=PASS\n"; +} + +void runCases(eda_layout::Pattern pattern) +{ + runCase("float2",pattern); runCase("float3",pattern); runCase("float4",pattern); + runCase("double2",pattern); runCase("double3",pattern); runCase("double4",pattern); + runCase("int2",pattern); runCase("int3",pattern); runCase("int4",pattern); + runCase("int64_2",pattern); runCase("int64_3",pattern); runCase("int64_4",pattern); +} + +template +void auditQuantizerBias(const cuBQL::box_t *boxes, uint32_t count, + const cuBQL::vec_t &expected) +{ + namespace radix = cuBQL::radixBuilder_impl; + using State = radix::BuildState; + DeviceBuffer state(1); + radix::clearBuildState<<<1,1>>>(state.data,count); + radix::fillBuildState<<<1,1024>>>(state.data,boxes,count); + radix::finishBuildState<<<1,1>>>(state.data); + check(cudaGetLastError()); + State hostState; + check(cudaMemcpy(&hostState,state.data,sizeof(State),cudaMemcpyDeviceToHost)); + for (int axis=0;axis<2;++axis) + if (hostState.quantizer.quantizeBias[axis] != expected[axis]) + throw std::runtime_error("quantizer centroid bounds lost input scalar precision"); +} + +template +__global__ void queryPointHits(cuBQL::BinaryBVH bvh, + const cuBQL::box_t *boxes, + uint32_t *hits, unsigned *overflow, uint32_t count) +{ + const uint32_t q = blockIdx.x*blockDim.x+threadIdx.x; + if (q >= count) return; + const auto point = boxes[q].center(); + uint32_t stack[64], depth = 1, numHits = 0; + stack[0] = 0; + while (depth) { + const auto node = bvh.nodes[stack[--depth]]; + if (point[0] < node.bounds.lower[0] || point[0] > node.bounds.upper[0] || + point[1] < node.bounds.lower[1] || point[1] > node.bounds.upper[1]) continue; + if (node.admin.count) { + for (uint32_t i=0;i= box.lower[0] && point[0] <= box.upper[0] && + point[1] >= box.lower[1] && point[1] <= box.upper[1]) ++numHits; + } + } else { + if (depth+2 > 64) { atomicExch(overflow,1u); break; } + stack[depth++] = uint32_t(node.admin.offset)+1; + stack[depth++] = uint32_t(node.admin.offset); + } + } + hits[q] = numHits; +} + +template +void auditPrecisionQueries(const std::vector> &boxes, + const cuBQL::box_t *deviceBoxes, + const cuBQL::BinaryBVH &a, + const cuBQL::BinaryBVH &b, + const cuBQL::BinaryBVH &c) +{ + const uint32_t count = uint32_t(boxes.size()); + DeviceBuffer hits(count); + DeviceBuffer overflow(1); + for (const auto &tree : {a,b,c}) { + check(cudaMemset(overflow.data,0,sizeof(unsigned))); + queryPointHits<<<1,256>>>(tree,deviceBoxes,hits.data,overflow.data,count); + check(cudaGetLastError()); + unsigned overflowValue = 0; + std::vector actual(count); + check(cudaMemcpy(&overflowValue,overflow.data,sizeof(unsigned),cudaMemcpyDeviceToHost)); + check(cudaMemcpy(actual.data(),hits.data,count*sizeof(uint32_t),cudaMemcpyDeviceToHost)); + if (overflowValue) throw std::runtime_error("precision query exceeded traversal stack"); + for (uint32_t q=0;q= box.lower[0] && point[0] <= box.upper[0] && + point[1] >= box.lower[1] && point[1] <= box.upper[1]) ++expected; + if (actual[q] != expected) + throw std::runtime_error("precision query differs from independent CPU point oracle"); + } + } +} + +template +void runPrecisionCase(const char *name, T origin, T span, T step, + bool requireFewerCollisions) +{ + namespace radix = cuBQL::radixBuilder_impl; + constexpr int lowBits = radix::MortonKeyTraits::numBits; + constexpr int highBits = radix::MortonKeyTraits::numBits; + std::vector> boxes(128), translated(128); + for (size_t i=0;i> deviceBoxes(boxes.size()), deviceTranslated(boxes.size()); + check(cudaMemcpy(deviceBoxes.data,boxes.data(),boxes.size()*sizeof(boxes[0]),cudaMemcpyHostToDevice)); + check(cudaMemcpy(deviceTranslated.data,translated.data(),translated.size()*sizeof(translated[0]),cudaMemcpyHostToDevice)); + auditQuantizerBias(deviceBoxes.data,uint32_t(boxes.size()),boxes[0].center()); + auditQuantizerBias(deviceBoxes.data,uint32_t(boxes.size()),boxes[0].center()); + auto keys32 = auditGeneratedKeys(deviceBoxes.data,uint32_t(boxes.size())); + auto keys64 = auditGeneratedKeys(deviceBoxes.data,uint32_t(boxes.size()),true); + if (keys32 != auditGeneratedKeys(deviceTranslated.data,uint32_t(boxes.size())) || + keys64 != auditGeneratedKeys(deviceTranslated.data,uint32_t(boxes.size()),true)) + throw std::runtime_error("Morton quantization changed after an exact coordinate translation"); + if (requireFewerCollisions) { + std::sort(keys32.begin(),keys32.end()); std::sort(keys64.begin(),keys64.end()); + const auto unique32 = std::unique(keys32.begin(),keys32.end())-keys32.begin(); + const auto unique64 = std::unique(keys64.begin(),keys64.end())-keys64.begin(); + if (unique64 <= unique32) + throw std::runtime_error("higher Morton precision did not resolve known low-precision collisions"); + } + auditBuilds(boxes,deviceBoxes.data,true); + std::cout << "precision_case=" << name << " boxes=" << boxes.size() + << " key_modes=3 status=PASS\n"; +} + +float timedBuild(const cuBQL::box_t *boxes, uint32_t count, bool key32) +{ + BuiltTree tree; + cuBQL::BuildConfig config(16); + Event start, stop; + check(cudaEventRecord(start.handle)); + if (key32) + cuBQL::cuda::radixBuilder(tree.bvh,boxes,count,config); + else + cuBQL::cuda::radixBuilder(tree.bvh,boxes,count,config); + check(cudaGetLastError()); + check(cudaEventRecord(stop.handle)); check(cudaEventSynchronize(stop.handle)); + float ms = 0; + check(cudaEventElapsedTime(&ms,start.handle,stop.handle)); + if (!std::isfinite(ms) || ms <= 0) + throw std::runtime_error("CUDA event interval is not finite and positive"); + tree.release(); + return ms; +} + +void benchmark(uint32_t tilesX, uint32_t tilesY, eda_layout::Pattern pattern) +{ + constexpr int meaningfulBits = cuBQL::radixBuilder_impl::MortonKeyTraits::meaningfulBits; + const auto boxes = eda_layout::makeBoxes(tilesX,tilesY,pattern); + cudaDeviceProp device{}; + int currentDevice = 0; + check(cudaGetDevice(¤tDevice)); check(cudaGetDeviceProperties(&device,currentDevice)); + int runtimeVersion = 0, driverVersion = 0; + check(cudaRuntimeGetVersion(&runtimeVersion)); check(cudaDriverGetVersion(&driverVersion)); + std::cout << "geometry=" << patternName(pattern) << " benchmark_device=" << device.name << "\n" + << "compute_capability=" << device.major << '.' << device.minor + << " device_bytes=" << device.totalGlobalMem + << " cuda_runtime=" << runtimeVersion << " cuda_driver=" << driverVersion << "\n" + << "tiles=" << tilesX << 'x' << tilesY << " boxes=" << boxes.size() + << " input_bytes=" << boxes.size()*sizeof(boxes[0]) << " geometry_fnv64=" + << std::hex << eda_layout::fingerprint(boxes) << std::dec + << " leaf=16 warmups_per_arm=1 balanced_pairs=6\n"; + DeviceBuffer> deviceBoxes(boxes.size()); + check(cudaMemcpy(deviceBoxes.data,boxes.data(),boxes.size()*sizeof(boxes[0]),cudaMemcpyHostToDevice)); + auditBuilds(boxes,deviceBoxes.data); + timedBuild(deviceBoxes.data,uint32_t(boxes.size()),false); + timedBuild(deviceBoxes.data,uint32_t(boxes.size()),true); + std::vector times[2]; + for (int round=0;round<6;++round) + for (int position=0;position<2;++position) { + const int arm = (round+position)%2; + const float ms = timedBuild(deviceBoxes.data,uint32_t(boxes.size()),arm != 0); + times[arm].push_back(ms); + std::cout << "round=" << round+1 << " position=" << position+1 + << " key_mode=" << (arm ? "uint32" : "legacy") + << " key_bits=" << (arm ? 32 : 64) + << " populated_bits=" << meaningfulBits + << " end_bit=" << (arm ? meaningfulBits : 64) << " build_ms=" << ms << "\n"; + } + for (auto &arm : times) std::sort(arm.begin(),arm.end()); + const float full = (times[0][2]+times[0][3])*0.5f; + const float narrow = (times[1][2]+times[1][3])*0.5f; + std::cout << "geometry=" << patternName(pattern) + << " median_legacy64_ms=" << full << " median_key32_ms=" << narrow + << " speedup=" << full/narrow << " validation=PASS\n"; +} + +uint32_t parseTileCount(const char *text, const char *name) +{ + const std::string error = std::string(name)+" must be a positive integer"; + uint32_t count = 0; + if (!*text) throw std::invalid_argument(error); + for (const char *p=text;*p;++p) { + if (*p < '0' || *p > '9') throw std::invalid_argument(error); + const uint32_t digit = uint32_t(*p-'0'); + if (count > (std::numeric_limits::max()-digit)/10) + throw std::invalid_argument(error); + count = count*10+digit; + } + if (!count) throw std::invalid_argument(error); + return count; +} + +void validateBenchmarkSize(uint32_t tilesX, uint32_t tilesY) +{ + constexpr uint32_t maxTiles = (1u<<24)/eda_layout::tilePitch; + if (tilesX > maxTiles || tilesY > maxTiles) + throw std::invalid_argument("tile dimensions exceed exact float coordinate range"); + const uint64_t count = uint64_t(tilesX)*tilesY*eda_layout::rectanglesPerTile; + if (count > uint64_t(std::numeric_limits::max()/2)) + throw std::invalid_argument("box count exceeds radix builder integer range"); + if (count > std::numeric_limits::max()/sizeof(cuBQL::box_t)) + throw std::invalid_argument("input exceeds host address range"); +} + +void printHelp(const char *program) +{ + std::cout << "Usage: " << program << " [--help | --benchmark [tilesX tilesY]\n" + << " | --benchmark-motif [tilesX tilesY] | --write-layout output.svg]\n" + << " No arguments: check legacy, uint32_t, and higher-precision uint64_t keys for 12\n" + << " scalar/dimension cases on each pattern (24 cases, 1024 boxes each), plus three\n" + << " focused precision cases covering collisions and large double/int64_t origins.\n" + << " --benchmark: generate an original layout-inspired pattern; default 256x256 tiles\n" + << " produce 2097152 float2 boxes (32 MiB input). Use --benchmark 8 4 for a small run.\n" + << " --benchmark-motif: regenerate the historical A100 motif; default 1024x512 tiles\n" + << " produce 16777216 boxes (256 MiB input). Use --benchmark-motif 8 4 for a small run.\n" + << " Both benchmarks validate all three key modes, warm the legacy and uint32_t modes,\n" + << " then time six balanced pairs of complete radix-builder calls with CUDA events\n" + << " (leaf threshold 16).\n" + << " Timed modes use the same quantization: legacy uint64_t storage/full 64-bit sort\n" + << " versus actual uint32_t storage/populated-bit sort (30 bits in 2D). Explicit\n" + << " uint64_t keys use higher precision and are validated separately. No speedup\n" + << " threshold is required.\n" + << " --write-layout: write a 16x8 preview of the new layout pattern without initializing\n" + << " CUDA. Both geometries use deterministic integer coordinates generated in C++;\n" + << " no image input or external layout data is needed.\n"; +} + +} // namespace + +int main(int argc, char **argv) +{ + try { + std::cout << std::fixed << std::setprecision(6); + if (argc == 2 && std::strcmp(argv[1],"--help") == 0) { + printHelp(argv[0]); + } else if (argc == 3 && std::strcmp(argv[1],"--write-layout") == 0) { + eda_layout::writeSvg(argv[2],16,8); + std::cout << "geometry=layout layout_svg=" << argv[2] << " tiles=16x8 boxes=" + << 16*8*eda_layout::rectanglesPerTile << "\n"; + } else if ((argc == 2 || argc == 4) && + (std::strcmp(argv[1],"--benchmark") == 0 || std::strcmp(argv[1],"--benchmark-motif") == 0)) { + const bool motif = std::strcmp(argv[1],"--benchmark-motif") == 0; + const uint32_t tilesX = argc == 4 ? parseTileCount(argv[2],"tilesX") : (motif ? 1024 : 256); + const uint32_t tilesY = argc == 4 ? parseTileCount(argv[3],"tilesY") : (motif ? 512 : 256); + validateBenchmarkSize(tilesX,tilesY); + benchmark(tilesX,tilesY,motif ? eda_layout::Pattern::Motif : eda_layout::Pattern::Layout); + } else if (argc == 1) { + runCases(eda_layout::Pattern::Layout); + runCases(eda_layout::Pattern::Motif); + runPrecisionCase("float2_collisions",0.f,float(1u<<20),1.f,true); + runPrecisionCase("double2_large_origin",double(uint64_t(1)<<40)+12345.,1024.,2.,false); + runPrecisionCase("int64_2_large_origin",(int64_t(1)<<40)+12345,1024,2,false); + } else { + throw std::invalid_argument("invalid arguments; use --help for usage"); + } + return 0; + } catch (const std::exception &error) { + std::cerr << "Morton sort-range test failed: " << error.what() << '\n'; + return 1; + } +}