### GPU Closest Point Query Example using C++ CUDA

Source: https://context7.com/nvidia/cubql/llms.txt

Demonstrates a complete GPU example for finding the closest point on a point cloud using cuBQL. It includes BVH construction and closest point queries. The example requires defining CUBQL_GPU_BUILDER_IMPLEMENTATION and includes necessary headers for BVH, CUDA builder, and closest point queries.

```cpp
#define CUBQL_GPU_BUILDER_IMPLEMENTATION 1
#include "cuBQL/bvh.h"
#include "cuBQL/builder/cuda.h"
#include "cuBQL/queries/pointData/findClosest.h"
#include <vector>
#include <iostream>

using namespace cuBQL;

__global__
void computeBoxes(box3f *boxes, const vec3f *points, int numPoints)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid >= numPoints) return;
    boxes[tid] = box3f().including(points[tid]);
}

__global__
void runQueries(
    bvh3f bvh,
    const vec3f *dataPoints,
    const vec3f *queryPoints,
    int *results,
    int numQueries)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid >= numQueries) return;

    results[tid] = cuBQL::points::findClosest(
        bvh, dataPoints, queryPoints[tid]
    );
}

int main()
{
    // Generate sample data
    int numData = 100000, numQueries = 1000;
    std::vector<vec3f> dataPoints(numData), queryPoints(numQueries);

    for (int i = 0; i < numData; i++)
        dataPoints[i] = vec3f{float(rand())/RAND_MAX,
                              float(rand())/RAND_MAX,
                              float(rand())/RAND_MAX};
    for (int i = 0; i < numQueries; i++)
        queryPoints[i] = vec3f{float(rand())/RAND_MAX,
                               float(rand())/RAND_MAX,
                               float(rand())/RAND_MAX};

    // Allocate device memory
    vec3f *d_data, *d_queries;
    box3f *d_boxes;
    int *d_results;

    cudaMalloc(&d_data, numData * sizeof(vec3f));
    cudaMalloc(&d_queries, numQueries * sizeof(vec3f));
    cudaMalloc(&d_boxes, numData * sizeof(box3f));
    cudaMalloc(&d_results, numQueries * sizeof(int));

    cudaMemcpy(d_data, dataPoints.data(), numData * sizeof(vec3f), cudaMemcpyHostToDevice);
    cudaMemcpy(d_queries, queryPoints.data(), numQueries * sizeof(vec3f), cudaMemcpyHostToDevice);

    // Compute bounding boxes
    computeBoxes<<<(numData+127)/128, 128>>>(d_boxes, d_data, numData);

    // Build BVH
    bvh3f bvh;
    gpuBuilder(bvh, d_boxes, numData, BuildConfig());

    // Run queries
    runQueries<<<(numQueries+127)/128, 128>>>(bvh, d_data, d_queries, d_results, numQueries);

    // Retrieve results
    std::vector<int> results(numQueries);
    cudaMemcpy(results.data(), d_results, numQueries * sizeof(int), cudaMemcpyDeviceToHost);

    std::cout << "First 5 results: ";
    for (int i = 0; i < 5; i++) std::cout << results[i] << " ";
    std::cout << std::endl;

    // Cleanup
    cuda::free(bvh);
    cudaFree(d_data);
    cudaFree(d_queries);
    cudaFree(d_boxes);
    cudaFree(d_results);

    return 0;
}

```

--------------------------------

### Configure BVH Construction in C++

Source: https://context7.com/nvidia/cubql/llms.txt

Demonstrates the use of the `BuildConfig` class to control BVH construction parameters. Options include selecting build methods (spatial median, SAH, ELH) and setting leaf size thresholds for quality/performance tradeoffs. Examples show how to enable SAH and ELH, and customize leaf creation.

```cpp
#include "cuBQL/bvh.h"

// Default configuration uses spatial median splits
cuBQL::BuildConfig defaultConfig;

// Enable Surface Area Heuristic for better ray tracing performance
cuBQL::BuildConfig sahConfig;
sahConfig.enableSAH();
// Or equivalently:
// sahConfig.buildMethod = cuBQL::BuildConfig::SAH;

// Enable Edge Length Heuristic (experimental)
cuBQL::BuildConfig elhConfig;
elhConfig.enableELH();

// Control leaf creation threshold
cuBQL::BuildConfig customConfig;
customConfig.makeLeafThreshold = 8;    // Create leaves with <= 8 primitives
customConfig.maxAllowedLeafSize = 64;  // Maximum allowed leaf size

// Build with SAH for potentially faster queries (slower build)
cuBQL::BinaryBVH<float, 3> bvh;
cuBQL::box3f *d_boxes; // ... device memory with boxes ...
int numBoxes = 100000;
cuBQL::gpuBuilder(bvh, d_boxes, numBoxes, sahConfig);
```

--------------------------------

### CMake Integration for cuBQL

Source: https://context7.com/nvidia/cubql/llms.txt

Provides an example CMakeLists.txt file for integrating cuBQL into a project. It demonstrates how to set CUDA architectures and include cuBQL either as a subdirectory for header-only usage or by linking against a pre-built library.

```cmake
cmake_minimum_required(VERSION 3.18)
project(MyProject LANGUAGES CXX CUDA)

# Set CUDA architectures before including cuBQL
set(CMAKE_CUDA_ARCHITECTURES "70;80;86")
# Or for auto-detection (CMake 3.24+):
# set(CMAKE_CUDA_ARCHITECTURES "native")

# Include cuBQL as subdirectory
add_subdirectory(external/cuBQL)

# Option 1: Header-only mode (compile builders in your code)
# In your .cu file, add: #define CUBQL_GPU_BUILDER_IMPLEMENTATION 1
add_executable(my_app_headeronly
    main.cu
)
target_link_libraries(my_app_headeronly PRIVATE cuBQL)

```

--------------------------------

### GPU BVH Builder Implementation in C++

Source: https://context7.com/nvidia/cubql/llms.txt

Implements the primary GPU builder for constructing a `BinaryBVH` using adaptive spatial median splits. It requires an array of bounding boxes in device memory and automatically handles invalid primitives by excluding them. The example includes a kernel for computing bounding boxes from point data.

```cpp
#define CUBQL_GPU_BUILDER_IMPLEMENTATION 1
#include "cuBQL/bvh.h"
#include "cuBQL/builder/cuda.h"

// Kernel to compute bounding boxes for point data
__global__
void computeBoxes(cuBQL::box3f *d_boxes, const cuBQL::vec3f *d_points, int numPoints)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid >= numPoints) return;
    d_boxes[tid] = cuBQL::box3f().including(d_points[tid]);
}

int main()
{
    int numPoints = 1000000;

    // Allocate and populate point data (device memory)
    cuBQL::vec3f *d_points;
    cudaMalloc(&d_points, numPoints * sizeof(cuBQL::vec3f));
    // ... populate d_points ...

    // Compute bounding boxes for each primitive
    cuBQL::box3f *d_boxes;
    cudaMalloc(&d_boxes, numPoints * sizeof(cuBQL::box3f));
    computeBoxes<<<(numPoints + 127) / 128, 128>>>(d_boxes, d_points, numPoints);

    // Build BVH with default configuration
    cuBQL::BinaryBVH<float, 3> bvh;
    cuBQL::BuildConfig config;
    cuBQL::gpuBuilder(bvh, d_boxes, numPoints, config);

    // Optional: Use custom stream and memory resource
    cudaStream_t stream;
    cudaStreamCreate(&stream);
    cuBQL::gpuBuilder(bvh, d_boxes, numPoints, config, stream);

    // ... use BVH for queries ...

    // Free BVH resources
    cuBQL::cuda::free(bvh);
    cudaFree(d_boxes);
    cudaFree(d_points);

    return 0;
}
```

--------------------------------

### CuBQL Lambda for Closest Point Search (CUDA C++)

Source: https://github.com/nvidia/cubql/blob/main/README.md

Illustrates a specific application of the traversal template: finding the closest point. The lambda callback tracks the minimum distance and closest primitive ID encountered during the traversal. This example highlights how the callback can update state and return the current closest distance as the shrinking radius.

```cuda_cpp
void findClosestPoint(...)
{
   float closestDist = INFINITY;
   int   closestID   = -1;
   auto myQueryLambda = [&closestDist,&closestID,...](int primID) -> float {
      // Calculate distance to the current primitive
      float dist = distance(queryPoint,myPrims[primID]);
      // Update closest if a nearer primitive is found
      if (dist < closestDist) 
         { closestDist = dist; closestID = primID; }
      // Return the current closest distance to shrink the search radius
      return closestDist;
   };
   // Assume cuBQL::shrinkingRadiusQuery::forEachPrim is called elsewhere with this lambda
}
```

--------------------------------

### Header-only Build: CMake Configuration

Source: https://github.com/nvidia/cubql/blob/main/README.md

This CMakeLists.txt configuration demonstrates how to integrate cuBQL using the header-only approach. It involves adding cuBQL as a subdirectory and linking to the 'cuBQL' target, which acts as an INTERFACE target providing include paths without building a separate library.

```cmake
add_subdirectory(<pathTo>/cuBQL)
	
add_executable(userExec ... 
    userMain.cu ...)
	
target_link_libraries(userExec ...
    cuBQL)
```

--------------------------------

### Build CUDA Sample: cuBQL_sample01_points_closestPoint_cuda

Source: https://github.com/nvidia/cubql/blob/main/samples/s01_closestPoint_points_gpu/CMakeLists.txt

Configures the build for the 'cuBQL_sample01_points_closestPoint_cuda' executable using CUDA. It requires the 'closestPoint.cu' source file and links against 'cuBQL_cuda_float3' and 'cuBQL_samples_common' libraries. This snippet is executed only if CUBQL_HAVE_CUDA is enabled.

```cmake
if (CUBQL_HAVE_CUDA)
  add_executable(cuBQL_sample01_points_closestPoint_cuda
    closestPoint.cu
    )
  target_link_libraries(cuBQL_sample01_points_closestPoint_cuda
    # the cuda-side builders for float3 data
    cuBQL_cuda_float3
    # common samples stuff
    cuBQL_samples_common
    )
endif()
```

--------------------------------

### Build CUDA Sample: cuBQL_sample01_points_closestPoint_wideBVH_cuda

Source: https://github.com/nvidia/cubql/blob/main/samples/s01_closestPoint_points_gpu/CMakeLists.txt

Configures the build for the 'cuBQL_sample01_points_closestPoint_wideBVH_cuda' executable using CUDA. It requires the 'closestPoint_WideBVH.cu' source file and links against 'cuBQL_cuda_float3' and 'cuBQL_samples_common' libraries. This snippet is executed only if CUBQL_HAVE_CUDA is enabled.

```cmake
if (CUBQL_HAVE_CUDA)
  add_executable(cuBQL_sample01_points_closestPoint_wideBVH_cuda
    closestPoint_WideBVH.cu
    )
  target_link_libraries(cuBQL_sample01_points_closestPoint_wideBVH_cuda
    # the cuda-side builders for float3 data
    cuBQL_cuda_float3
    # common samples stuff
    cuBQL_samples_common
    )
endif()
```

--------------------------------

### CPU BVH Builder (cuBQL)

Source: https://context7.com/nvidia/cubql/llms.txt

Illustrates how to construct a BVH on the CPU using cuBQL's `cpuBuilder`. This method is suitable for environments without GPU access or for preprocessing. It utilizes spatial median splits and requires host memory arrays for points and boxes. The BVH nodes and primitive IDs are stored in host memory.

```cpp
#include "cuBQL/bvh.h"
#include "cuBQL/builder/cpu.h"

int main()
{
    int numPoints = 10000;

    // Host memory arrays
    std::vector<cuBQL::vec3f> points(numPoints);
    std::vector<cuBQL::box3f> boxes(numPoints);

    // Generate sample points
    for (int i = 0; i < numPoints; i++) {
        points[i] = cuBQL::vec3f{
            float(rand()) / RAND_MAX,
            float(rand()) / RAND_MAX,
            float(rand()) / RAND_MAX
        };
        boxes[i] = cuBQL::box3f().including(points[i]);
    }

    // Build BVH on CPU
    cuBQL::BinaryBVH<float, 3> bvh;
    cuBQL::BuildConfig config;
    cuBQL::cpuBuilder(bvh, boxes.data(), numPoints, config);

    // Query on CPU (BVH nodes/primIDs are in host memory)
    // ... perform queries ...

    // Free host-allocated BVH
    cuBQL::cpu::freeBVH(bvh);

    return 0;
}

```

--------------------------------

### CUDA GPU BVH Builders (cuBQL)

Source: https://context7.com/nvidia/cubql/llms.txt

Demonstrates the usage of different GPU-accelerated BVH builders in cuBQL. These include fast radix/Morton code builders, a rebinning radix builder for challenging data, and an SAH-based builder for better query performance. Requires the CUBQL_GPU_BUILDER_IMPLEMENTATION macro and relevant headers.

```cpp
#define CUBQL_GPU_BUILDER_IMPLEMENTATION 1
#include "cuBQL/bvh.h"
#include "cuBQL/builder/cuda.h"

cuBQL::BinaryBVH<float, 3> bvh;
cuBQL::box3f *d_boxes; // ... device boxes ...
int numBoxes = 1000000;
cuBQL::BuildConfig config;

// Fast radix/Morton code builder
cuBQL::cuda::radixBuilder(bvh, d_boxes, numBoxes, config);

// Radix builder with rebinning for challenging data distributions
// (handles cases where standard Morton codes produce poor results)
cuBQL::cuda::rebinRadixBuilder(bvh, d_boxes, numBoxes, config);

// SAH-based builder (slower build, better query performance)
cuBQL::cuda::sahBuilder(bvh, d_boxes, numBoxes, config);

// Clean up
cuBQL::cuda::free(bvh);

```

--------------------------------

### GPU BVH Builder Usage in C++

Source: https://github.com/nvidia/cubql/blob/main/README.md

Demonstrates how to use the cuBQL::gpuBuilder function to construct a BinaryBVH on the GPU. It requires including the 'cuBQL/bvh.h' header, preparing an array of bounding boxes in device memory, and then calling the gpuBuilder function with the BVH object, bounding boxes, number of primitives, and build configuration.

```cpp
#include "cuBQL/bvh.h"
...
box3f *d_boxes  = 0;
int    numBoxes = 0;
userCodeForGeneratingPrims(&d_boxes,&numBoxes, ...);
...
cuBQL::BinaryBVH<float,3> bvh;
cuBQL::BuildConfig buildParams;
cuBQL::gpuBuilder(bvh,d_boxes,numBoxes,buildParams);
...
```

--------------------------------

### Build cuBQL for CPU-Only Execution

Source: https://context7.com/nvidia/cubql/llms.txt

This snippet shows how to configure a build for cuBQL that runs entirely on the CPU, without requiring CUDA. It links against the base cuBQL library, suitable for applications where GPU acceleration is not needed or available. The input is 'main.cpp', producing the 'my_app_cpu' executable.

```cmake
add_executable(my_app_cpu
    main.cpp
)
target_link_libraries(my_app_cpu PRIVATE cuBQL)
```

--------------------------------

### Link Pre-built cuBQL Library for Float3 Data (CUDA)

Source: https://context7.com/nvidia/cubql/llms.txt

This snippet demonstrates how to link a pre-built cuBQL library specifically compiled for float3 data types. It assumes the cuBQL library is available and configured in the build system. The primary input is the source file 'main.cu', and the output is an executable 'my_app_float3'.

```cmake
add_executable(my_app_float3
    main.cu
)
target_link_libraries(my_app_float3 PRIVATE cuBQL_cuda_float3)
```

--------------------------------

### Header-only Build: CUDA Source

Source: https://github.com/nvidia/cubql/blob/main/README.md

This snippet shows how to use cuBQL in a header-only mode within your CUDA source files. It requires defining 'CUBQL_GPU_BUILDER_IMPLEMENTATION' before including the cuBQL header to enable explicit instantiation. This method allows cuBQL to be used directly from any compiler and build system by providing include paths.

```cuda
#define CUBQL_GPU_BUILDER_IMPLEMENTATION 1
#include <cuBQL/bvh.h>
...
void foo(...) {
	cuBQL::gpuBuilder(...)
}
```

--------------------------------

### Create cuBQL Samples Common Library (CMake)

Source: https://github.com/nvidia/cubql/blob/main/samples/common/CMakeLists.txt

Defines the CMake target for the cuBQL samples common library. It lists the source files and links the necessary cuBQL library.

```cmake
add_library(cuBQL_samples_common
  Generator.h
  Generator.cpp
  tiny_obj_loader.h
  loadOBJ.h
  loadOBJ.cpp
  loadBinMesh.h
  loadBinMesh.cpp
  )
target_link_libraries(cuBQL_samples_common
  PUBLIC
  cuBQL
  )
```

--------------------------------

### Sign Commits with Developer Certificate of Origin (Bash)

Source: https://github.com/nvidia/cubql/blob/main/CONTRIBUTING.md

This command shows how to sign your Git commits using the `--signoff` or `-s` flag. This action appends a 'Signed-off-by' line to your commit message, certifying your contribution. The appended text confirms your agreement with the Developer Certificate of Origin.

```bash
$ git commit -s -m "Add cool feature."

This will append the following to your commit message:
Signed-off-by: Your Name <your@email.com>
```

--------------------------------

### CuBQL Shrinking Radius Query Traversal Template (CUDA C++)

Source: https://github.com/nvidia/cubql/blob/main/README.md

Demonstrates the core structure of a shrinking radius query using CuBQL's traversal templates. It shows how to define a lambda callback for per-primitive processing and pass it to the `forEachPrim` function. This pattern is fundamental for building custom GPU queries.

```cuda_cpp
inline __device__ 
void myQuery(bvh3f myBVH, 
             MyPrim *myPrims, 
             vec3f myQueryPoint,
             ...) 
{
  auto myQueryLambda = [...](int primID) -> float {
   // Custom logic to process primitive `primID`
   // Return the new maximum search radius
   ...
  };
  // Execute the traversal with the custom lambda
  cuBQL::shrinkingRadiusQuery::forEachPrim
     (myQueryLambda, myQueryPoint, myBVH, ...);
}
```

--------------------------------

### Clone and Push Changes to Forked Repository (Bash)

Source: https://github.com/nvidia/cubql/blob/main/CONTRIBUTING.md

This snippet demonstrates the basic Git commands to clone a forked repository and push local changes to a remote branch on your fork. It assumes you have already forked the upstream repository. Ensure you replace placeholders with your actual username and branch names.

```bash
git clone https://github.com/YOUR_USERNAME/YOUR_FORK.git nvCUBQL
# Checkout the targeted branch and commit changes
# Push the commits to a branch on the fork (remote).
git push -u origin <local-branch>:<remote-branch>
```

--------------------------------

### Predefined Target Build: CUDA Source

Source: https://github.com/nvidia/cubql/blob/main/README.md

This snippet illustrates using cuBQL with a predefined target, such as for float3 data, in your CUDA source files. Unlike the header-only mode, 'CUBQL_GPU_BUILDER_IMPLEMENTATION' should not be defined. This approach simplifies integration by leveraging pre-compiled targets.

```cuda
// do NOT define CUBQL_GPU_BUILDER_IMPLEMENTATION 
#include <cuBQL/bvh.h>
...
void foo(...) {
   cuBQL::gpuBuilder(...)
}
```

--------------------------------

### Define cuBQL Queries Interface Library (CMake)

Source: https://github.com/nvidia/cubql/blob/main/cuBQL/CMakeLists.txt

Defines `cuBQL_queries` as an interface library containing header files for various query types, such as nearest neighbor search (`knn.h`) and finding the closest point (`findClosest.h`). This library is header-only and automatically works on both CPU and CUDA. It links against the main `cuBQL` library.

```cmake
# the collection of all different type-specific queries currently
# supplied by cubql. all these should be header-only and should all
# automatically on both cpu and cuda, so this is a INTERFACE library
add_library(cuBQL_queries INTERFACE
  #
  queries/pointData/knn.h
  #
  queries/pointData/findClosest.h
  queries/pointData/knn.h
  )

set_target_properties(cuBQL_queries PROPERTIES FOLDER "CuBQL")

target_link_libraries(cuBQL_queries
  INTERFACE
  cuBQL
  )

```

--------------------------------

### Wide BVH Builders (cuBQL)

Source: https://context7.com/nvidia/cubql/llms.txt

Shows how to construct Wide BVHs with a fixed branching factor (e.g., 4 or 8) using cuBQL's `gpuBuilder`. Wide BVHs are optimized for SIMD traversal. The process involves building a BinaryBVH first and then collapsing it, or directly constructing a WideBVH. Requires the CUBQL_GPU_BUILDER_IMPLEMENTATION macro.

```cpp
#define CUBQL_GPU_BUILDER_IMPLEMENTATION 1
#include "cuBQL/bvh.h"
#include "cuBQL/builder/cuda.h"

// Build a wide BVH with branching factor 4
cuBQL::WideBVH<float, 3, 4> wideBvh4;
cuBQL::box3f *d_boxes; // ... device boxes ...
int numBoxes = 100000;
cuBQL::BuildConfig config;

// Direct wide BVH construction
cuBQL::gpuBuilder(wideBvh4, d_boxes, numBoxes, config);

// Or build 8-wide BVH
cuBQL::WideBVH<float, 3, 8> wideBvh8;
cuBQL::gpuBuilder(wideBvh8, d_boxes, numBoxes, config);

// Free wide BVH
cuBQL::cuda::free(wideBvh4);
cuBQL::cuda::free(wideBvh8);

```

--------------------------------

### Define stb_image Library and Include Directories (CMake)

Source: https://github.com/nvidia/cubql/blob/main/samples/3rdParty/stb_image/CMakeLists.txt

This snippet defines the 'stb_image' library as an INTERFACE library in CMake and specifies its include directories. It ensures that the 'stb_image' headers are accessible to targets that link against it. This is a common pattern for managing external libraries within a CMake build system.

```cmake
add_library(stb_image INTERFACE)
target_include_directories(stb_image INTERFACE ${CMAKE_CURRENT_LIST_DIR})
```

--------------------------------

### Generate Specific Instantiations (CMake Function)

Source: https://github.com/nvidia/cubql/blob/main/cuBQL/CMakeLists.txt

Defines a CMake function `add_specific_instantiation` to create shared and static libraries for specific type and device combinations (e.g., `cuBQL_cuda_float4`). This function handles setting compile definitions, visibility, and linking against the main `cuBQL` library. It's used to provide pre-compiled builders for users who don't want to use the header-only mechanism.

```cmake
# helper for creating type- and device-specific implementations of
# cubql; i.e., one for host_float4, one for cuda_int3, etc. Since
# everything other than the builders are entirely header only these
# type-specific targets will, in fact, only contain instantiations of
# the specific builders for the given type and device
function(add_specific_instantiation device suffix T D)
  add_library(cuBQL_${device}_${T}${D} SHARED EXCLUDE_FROM_ALL
    builder/${device}/instantiate_builders.${suffix}
    )
  if (${device} STREQUAL "cuda")
    target_compile_definitions(cuBQL_${device}_${T}${D}
      PUBLIC
      -DCUBQL_HAVE_CUDA=1
      )
  endif()
  set_target_properties(cuBQL_${device}_${T}${D}
    PROPERTIES
    POSITION_INDEPENDENT_CODE ON
    WINDOWS_EXPORT_ALL_SYMBOLS TRUE
    FOLDER "CuBQL"
    )
  target_link_libraries(cuBQL_${device}_${T}${D}
    PUBLIC
    cuBQL
    )
  target_compile_definitions(cuBQL_${device}_${T}${D}
    PRIVATE
    -DCUBQL_INSTANTIATE_T=${T}
    -DCUBQL_INSTANTIATE_D=${D}
    )

  add_library(cuBQL_${device}_${T}${D}_static STATIC EXCLUDE_FROM_ALL
    builder/${device}/instantiate_builders.${suffix}
    )
  if (${device} STREQUAL "cuda")
    target_compile_definitions(cuBQL_${device}_${T}${D}_static
      PUBLIC
      -DCUBQL_HAVE_CUDA=1
      )
  endif()
  set_target_properties(cuBQL_${device}_${T}${D}_static
    PROPERTIES
    POSITION_INDEPENDENT_CODE ON
    FOLDER "CuBQL"
    )
  target_link_libraries(cuBQL_${device}_${T}${D}_static
    PUBLIC
    cuBQL
    )
  set_target_properties(cuBQL_${device}_${T}${D}_static
    PROPERTIES
    CXX_VISIBILITY_PRESET  default
    CUDA_VISIBILITY_PRESET default
    POSITION_INDEPENDENT_CODE ON
    CUDA_SEPARABLE_COMPILATION ON
    CUDA_RESOLVE_DEVICE_SYMBOLS ON
    CUDA_USE_STATIC_CUDA_RUNTIME ON
    )
  target_compile_definitions(cuBQL_${device}_${T}${D}_static
    PRIVATE
    -DCUBQL_INSTANTIATE_T=${T}
    -DCUBQL_INSTANTIATE_D=${D}
    )
endfunction()

```

--------------------------------

### Shrinking Radius Query for Closest Point Search (C++)

Source: https://context7.com/nvidia/cubql/llms.txt

Implements a device kernel using the shrinking radius query template to find the closest point to a query point. It traverses a BVH and uses a lambda function to update the closest point found. Dependencies include cuBQL BVH and traversal headers.

```cpp
#include "cuBQL/bvh.h"
#include "cuBQL/traversal/shrinkingRadiusQuery.h"

// Device kernel for custom shrinking radius query
__device__
int customFindClosest(cuBQL::bvh3f bvh,
                      const cuBQL::vec3f *points,
                      cuBQL::vec3f queryPoint)
{
    int closestID = -1;
    float closestSqrDist = CUBQL_INF;

    // Lambda called for each candidate primitive
    auto candidateLambda = [&closestID, &closestSqrDist, points, queryPoint]
        (uint32_t primID) -> float
    {
        float sqrDist = cuBQL::fSqrDistance_rd(points[primID], queryPoint);
        if (sqrDist < closestSqrDist) {
            closestSqrDist = sqrDist;
            closestID = primID;
        }
        // Return new maximum search radius (squared)
        return closestSqrDist;
    };

    // Execute shrinking radius query
    cuBQL::shrinkingRadiusQuery::forEachPrim(
        candidateLambda,    // Per-primitive callback
        bvh,                // BVH to traverse
        queryPoint,         // Query center
        CUBQL_INF           // Initial max radius squared
    );

    return closestID;
}

// Alternative: forEachLeaf processes leaves instead of individual primitives
__device__
int findClosestViaLeaf(cuBQL::bvh3f bvh,
                       const cuBQL::vec3f *points,
                       cuBQL::vec3f queryPoint)
{
    int closestID = -1;
    float closestSqrDist = CUBQL_INF;

    auto leafLambda = [&closestID, &closestSqrDist, points, queryPoint]
        (const uint32_t *primIDs, size_t numPrims) -> float
    {
        for (size_t i = 0; i < numPrims; i++) {
            uint32_t primID = primIDs[i];
            float sqrDist = cuBQL::fSqrDistance_rd(points[primID], queryPoint);
            if (sqrDist < closestSqrDist) {
                closestSqrDist = sqrDist;
                closestID = primID;
            }
        }
        return closestSqrDist;
    };

    cuBQL::shrinkingRadiusQuery::forEachLeaf(
        leafLambda, bvh, queryPoint, CUBQL_INF
    );

    return closestID;
}

```

--------------------------------

### Custom GPU Memory Allocation (cuBQL)

Source: https://context7.com/nvidia/cubql/llms.txt

Demonstrates using `GpuMemoryResource` interfaces in cuBQL for custom GPU memory allocation. Options include `ManagedMemMemoryResource`, `DeviceMemoryResource`, and `AsyncGpuMemoryResource` (default for CUDA 11.2+). This allows for strategies like async allocation or custom memory pools, impacting build performance.

```cpp
#include "cuBQL/bvh.h"
#include "cuBQL/builder/cuda.h"

// Use managed memory (accessible from both CPU and GPU)
cuBQL::ManagedMemMemoryResource managedMem;

// Use regular device memory (explicit control for multi-GPU)
cuBQL::DeviceMemoryResource deviceMem;

// Use async memory allocation (CUDA 11.2+, faster builds)
// AsyncGpuMemoryResource is the default when available
int deviceId = 0;
cuBQL::AsyncGpuMemoryResource asyncMem(deviceId);

// Build BVH with custom memory resource
cuBQL::BinaryBVH<float, 3> bvh;
cuBQL::box3f *d_boxes; // ... device boxes ...
int numBoxes = 100000;
cudaStream_t stream = 0;

cuBQL::gpuBuilder(bvh, d_boxes, numBoxes,
                  cuBQL::BuildConfig(),
                  stream,
                  managedMem);

// Free with same memory resource
cuBQL::cuda::free(bvh, stream, managedMem);

```

--------------------------------

### Predefined Target Build: CMake Configuration

Source: https://github.com/nvidia/cubql/blob/main/README.md

This CMakeLists.txt configuration shows how to build cuBQL using a predefined target, specifically 'cuBQL_cuda_float3'. After adding cuBQL as a subdirectory, you link your executable to this specific target, which is pre-configured for float3 data types.

```cmake
add_subdirectory(<pathTo>/cuBQL)
	
add_executable(userExec ... 
   userMain.cu ...)
	
target_link_libraries(userExec ...
   cuBQL_cuda_float3)
```

--------------------------------

### Generate All Specific Instantiations (CMake Loop)

Source: https://github.com/nvidia/cubql/blob/main/cuBQL/CMakeLists.txt

This CMake code iterates through common data types (`float`, `int`, `double`, `longlong`) and dimensions (`2`, `3`, `4`) to generate specific CPU and CUDA builder instantiations using the `add_specific_instantiation` function. These generated targets are marked `EXCLUDE_FROM_ALL` to avoid building unnecessary targets.

```cmake
# ------------------------------------------------------------------
# generate all type/dim specific targets, ie, one target for each
# {int,float,double,long}x{2,3,4}. Each such target contains the
# pre-compiled builder(s) for that specific type/dimension, for the
# case where the user does NOT want to use the header-only
# mechanism. To avoid "polluting" the user's project with lots of
# different targets that he or she may or may not need we add all
# those as 'EXCLUDE_FROM_ALL', so only those that get actually used
# will actually get built
# ------------------------------------------------------------------
foreach(T IN ITEMS float int double longlong)
  foreach(D IN ITEMS 2 3 4)
    add_specific_instantiation(cpu cpp ${T} ${D})
    if (CUBQL_HAVE_CUDA)
      add_specific_instantiation(cuda cu ${T} ${D})
    endif()
  endforeach()
endforeach()

```

--------------------------------

### Fixed Radius Query for Range Search (C++)

Source: https://context7.com/nvidia/cubql/llms.txt

Demonstrates using the fixed radius query template to count primitives within a specified radius of a query point. It utilizes a lambda for primitive processing and returns the count. Dependencies include cuBQL BVH and traversal headers.

```cpp
#include "cuBQL/bvh.h"
#include "cuBQL/traversal/fixedRadiusQuery.h"

// Device kernel to count points within radius
__device__
int countPointsInRadius(cuBQL::bvh3f bvh,
                        const cuBQL::vec3f *points,
                        cuBQL::vec3f queryPoint,
                        float radius)
{
    int count = 0;
    float radiusSquared = radius * radius;

    // Lambda called for each candidate primitive
    auto countLambda = [&count, points, queryPoint, radiusSquared]
        (uint32_t primID) -> int
    {
        float sqrDist = cuBQL::fSqrDistance_rd(points[primID], queryPoint);
        if (sqrDist <= radiusSquared) {
            count++;
        }
        return CUBQL_CONTINUE_TRAVERSAL;  // Continue searching
    };

    cuBQL::fixedRadiusQuery::forEachPrim(
        countLambda,
        bvh,
        queryPoint,
        radiusSquared  // Note: pass squared radius
    );

    return count;
}

// Find any point within radius (early termination)
__device__
bool hasPointInRadius(cuBQL::bvh3f bvh,
                      const cuBQL::vec3f *points,
                      cuBQL::vec3f queryPoint,
                      float radius)
{
    bool found = false;
    float radiusSquared = radius * radius;

    auto findLambda = [&found, points, queryPoint, radiusSquared]
        (uint32_t primID) -> int
    {
        float sqrDist = cuBQL::fSqrDistance_rd(points[primID], queryPoint);
        if (sqrDist <= radiusSquared) {
            found = true;
            return CUBQL_TERMINATE_TRAVERSAL;  // Stop searching
        }
        return CUBQL_CONTINUE_TRAVERSAL;
    };

    cuBQL::fixedRadiusQuery::forEachPrim(
        findLambda, bvh, queryPoint, radiusSquared
    );

    return found;
}

```

--------------------------------

### Define cuBQL Interface Library (CMake)

Source: https://github.com/nvidia/cubql/blob/main/cuBQL/CMakeLists.txt

Defines the main cuBQL library as an interface-only library. This means it primarily sets include paths and source files without compiling actual object code, acting as a header-only library for users. It requires users to manually set CUBQL_BUILDER_INSTANTIATION in their source files.

```cmake
# a interface-only library that only sets include paths etc; when
# using this the user has to manulaly set CUBQL_BUILDER_INSTANTIATION
# in one of his/her source files
add_library(cuBQL INTERFACE)

target_sources(cuBQL INTERFACE
  # main public "interface" to this library
  bvh.h
  # general math struff to make public stuff work
  math/common.h
  math/math.h
  math/vec.h
  math/box.h
  # general (lambda-templated) traversal routines
  traversal/shrinkingRadiusQuery.h
  traversal/fixedBoxQuery.h
  # host side builder(s)
  builder/cpu.h
  builder/cpu/spatialMedian.h
  # cuda side builder(s)
  builder/cuda.h
  builder/cuda/builder_common.h
  builder/cuda/sm_builder.h
  builder/cuda/sah_builder.h
  builder/cuda/gpu_builder.h
  builder/cuda/radix.h
  builder/cuda/rebinMortonBuilder.h
  builder/cuda/wide_gpu_builder.h
  )
target_include_directories(cuBQL INTERFACE
  ${PROJECT_SOURCE_DIR}
  )
set_target_properties(cuBQL
  PROPERTIES
  CXX_VISIBILITY_PRESET  default
  CUDA_VISIBILITY_PRESET default
  POSITION_INDEPENDENT_CODE ON
  CUDA_SEPARABLE_COMPILATION ON
  CUDA_RESOLVE_DEVICE_SYMBOLS ON
  CUDA_USE_STATIC_CUDA_RUNTIME ON
  )

```

--------------------------------

### K-Nearest Neighbors Query using BVH Traversal in C++

Source: https://context7.com/nvidia/cubql/llms.txt

This snippet illustrates the `cuBQL::points::findKNN` function, which identifies the K nearest neighbors to a query point. It returns results in a provided candidate array, storing the point index (`primID`) and squared distance (`sqrDist`) for each neighbor. The function requires the output candidate array, the number of neighbors (K), the BVH structure, the point data, and the query point. The `result` struct returned indicates the number of neighbors found and the squared distance to the farthest neighbor.

```cpp
#include "cuBQL/bvh.h"
#include "cuBQL/queries/pointData/knn.h"

__global__
void runKNNQueries(cuBQL::bvh3f bvh,
                   const cuBQL::vec3f *dataPoints,
                   const cuBQL::vec3f *queryPoints,
                   cuBQL::knn::Candidate *allResults,
                   int K,
                   int numQueries)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid >= numQueries) return;

    // Each query gets K candidate slots
    cuBQL::knn::Candidate *myResults = allResults + tid * K;

    cuBQL::vec3f query = queryPoints[tid];

    // Find K nearest neighbors
    cuBQL::knn::Result result = cuBQL::points::findKNN(
        myResults,    // Output array for candidates
        K,            // Number of neighbors to find
        bvh,          // BVH over data points
        dataPoints,   // Original point data
        query         // Query point
    );

    // result.numFound = number of points actually found
    // result.sqrDistMax = squared distance to farthest found point

    // Access results
    for (int i = 0; i < result.numFound; i++) {
        int pointID = myResults[i].primID;
        float sqrDist = myResults[i].sqrDist;
        // Use pointID and sqrDist...
    }
}

// KNN candidate structure for reference:
// struct cuBQL::knn::Candidate {
//     int primID;      // Point index
//     float sqrDist;   // Squared distance to query
// };
// struct cuBQL::knn::Result {
//     int numFound;       // Actual number found
//     float sqrDistMax;   // Max squared distance in results
// };

```

--------------------------------

### Find Closest Point Query using BVH Traversal in C++

Source: https://context7.com/nvidia/cubql/llms.txt

This snippet demonstrates the `cuBQL::points::findClosest` function, which finds the nearest data point to a query point using an optimized BVH traversal. It shows how to find the absolute closest point and how to find the closest point within a specified maximum distance. The function requires a BVH structure, the original point data, and the query point(s). It returns the index of the closest point or -1 if no point is found within the distance limit.

```cpp
#include "cuBQL/bvh.h"
#include "cuBQL/queries/pointData/findClosest.h"

__global__
void runClosestPointQueries(cuBQL::bvh3f bvh,
                            const cuBQL::vec3f *dataPoints,
                            const cuBQL::vec3f *queryPoints,
                            int *results,
                            int numQueries)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid >= numQueries) return;

    cuBQL::vec3f query = queryPoints[tid];

    // Find closest point (no distance limit)
    int closestID = cuBQL::points::findClosest(
        bvh,          // BVH built over data points
        dataPoints,   // Original point data
        query         // Query point
    );

    // Find closest with maximum search distance
    float maxDistance = 10.0f;
    float maxDistSquared = maxDistance * maxDistance;
    int closestInRange = cuBQL::points::findClosest(
        bvh, dataPoints, query, maxDistSquared
    );
    // Returns -1 if no point within maxDistance

    results[tid] = closestID;
}

// Variant: exclude a specific point ID from results
__device__
int findClosestExcluding(cuBQL::bvh3f bvh,
                         const cuBQL::vec3f *points,
                         cuBQL::vec3f query,
                         int excludeID)
{
    return cuBQL::points::findClosest_exludeID(
        excludeID,  // Point ID to exclude
        bvh, points, query
    );
}

// Variant: exclude points at same position as query
__device__
int findClosestOther(cuBQL::bvh3f bvh,
                     const cuBQL::vec3f *points,
                     cuBQL::vec3f query)
{
    return cuBQL::points::findClosest_exludeSelf(bvh, points, query);
}

```

--------------------------------

### Ray Query Traversal for Intersections in C++

Source: https://context7.com/nvidia/cubql/llms.txt

Handles ray intersection queries within a BVH. It supports both fixed-length rays for any-hit tests (like shadow rays) and shrinking rays to find the closest intersection. The closest-hit query modifies the ray's tMax to the distance of the nearest hit.

```cpp
#include "cuBQL/bvh.h"
#include "cuBQL/traversal/rayQueries.h"

// Fixed ray query - find any intersection (shadow ray)
__device__
bool anyHit(cuBQL::bvh3f bvh,
            const cuBQL::Triangle *triangles,
            cuBQL::vec3f rayOrigin,
            cuBQL::vec3f rayDirection,
            float tMin, float tMax)
{
    bool hitFound = false;
    cuBQL::ray3f ray{rayOrigin, rayDirection, tMin, tMax};

    auto testLambda = [&hitFound, triangles, ray](uint32_t primID) -> int
    {
        const cuBQL::Triangle &tri = triangles[primID];
        if (cuBQL::rayIntersectsTriangle(ray, tri)) {
            hitFound = true;
            return CUBQL_TERMINATE_TRAVERSAL;
        }
        return CUBQL_CONTINUE_TRAVERSAL;
    };

    cuBQL::fixedRayQuery::forEachPrim(testLambda, bvh, ray);
    return hitFound;
}

// Shrinking ray query - find closest intersection
__device__
int closestHit(cuBQL::bvh3f bvh,
               const cuBQL::Triangle *triangles,
               cuBQL::ray3f &ray)  // Note: ray.tMax will be modified
{
    int closestTriID = -1;

    auto testLambda = [&closestTriID, triangles, &ray](uint32_t primID) -> float
    {
        const cuBQL::Triangle &tri = triangles[primID];
        float t;
        if (cuBQL::rayIntersectsTriangle(ray, tri, t) && t < ray.tMax) {
            ray.tMax = t;
            closestTriID = primID;
        }
        return ray.tMax;
    };

    cuBQL::shrinkingRayQuery::forEachPrim(testLambda, bvh, ray);
    return closestTriID;
}

```

--------------------------------

### Set C++ Standard and CMake Policies in CMake

Source: https://github.com/nvidia/cubql/blob/main/CMakeLists.txt

This snippet sets the C++ standard to 17 and ensures it's required. It also applies a CMake policy (CMP0104) for newer CMake versions (3.18 and above) to ensure consistent behavior.

```cmake
set(CMAKE_CXX_STANDARD 17)
set(CMAKE_CXX_STANDARD_REQUIRED ON)

if(${CMAKE_VERSION} VERSION_GREATER_EQUAL 3.18)
  cmake_policy(SET CMP0104 NEW)
endif()
```

--------------------------------

### Configure CUDA and OpenMP Build Options in CMake

Source: https://github.com/nvidia/cubql/blob/main/CMakeLists.txt

This snippet configures build options for CUDA and OpenMP. It checks for user-defined flags like CUBQL_OMP and CUBQL_DISABLE_CUDA to determine whether to build CUDA-enabled libraries. It also includes logic to find and enable the CUDA compiler if available, falling back to host-only libraries if CUDA is not found or explicitly disabled.

```cmake
cmake_minimum_required(VERSION 3.16)

cmake_policy(SET CMP0048 NEW)
set(CMAKE_BUILD_TYPE_INIT "Release")
project(cuBQL VERSION 1.2.0 LANGUAGES C CXX)

if (CUBQL_OMP)
  set(CUBQL_DISABLE_CUDA ON)
endif()
if (CUBQL_DISABLE_CUDA)
  message("#cuBQL: CUDA _DISABLED_ by user request")
  set(CUBQL_HAVE_CUDA OFF)
else()
if (NOT CMAKE_CUDA_COMPILER)
  include(CheckLanguage)
  check_language(CUDA)
endif()

if (CMAKE_CUDA_COMPILER)
  message("#cuBQL: CUDA _FOUND_! building both cuda and host libs")
  enable_language(CUDA)
  set(CUBQL_HAVE_CUDA ON)
else()
  message(AUTHOR_WARNING
    " ===========================================================\n"
    " #cuBQL: could not find CUDA - going to build only host libs\n"
    " ===========================================================\n"
    )
  set(CUBQL_HAVE_CUDA OFF)
endif()
endif()
```