ingowald / cudakdtree Goto Github PK

Hello,

This is a nice project. Do you have a specific licence on this code?
We are thinking on using it in Advanced Computer Graphics class (UdelaR), where students will code Photon Mapping using OptiX. I'm guessing cudaKDTree should be ideal in this case.

Please let me know if we have permissions for this.

Thanks,
José Aguerre

Hello, I'm having the following problem:
when I run ./cukd_sample and other testing code，it will have some problem:

terminate called after throwing an instance of thrust: system:: system_error'
what(): parallel_for failed: cudErrorNoKernelImageForDevice: no kernel image is available for execution on the device.

The full output is below:
•/cukd_sample

generating 1000000 uniform random points
allocated memory for the world space bounding box ... calling builder...
terminate called after throwing an instance of thrust: system:: system_error'
what(): parallel_for failed: cudErrorNoKernelImageForDevice: no kernel image is available for execution on the device

My code environment is as follows:
wls2,ubuntu20.04，
cuda11.4,cudnn8.6.0,
4090

I compiled the code as with cmake as follows commod:
mkdir build
cd build
cmake ..
make

plz advise me on how ti fix this;
thankx!

hello，when i follow the example like:

/* code */
struct PointPlusPayload {
float3 position;
int payload;
};

struct PointPlusPayload_traits
: public cukd::default_data_traits
{
using point_t = float3;
static inline device host
float3 &get_point(PointPlusPayload &data) { return data.position; }

static inline device host
float get_coord(const PointPlusPayload &data, int dim)
{ return cukd::get_coord(get_point(data),dim); }
};

void test_kdtree()
{

PointPlusPayload data;
cudaMallocManaged((char **)&data,countsizeof(*data));

/... fill the data/

cukd::buildTree
</* type of the data: /PointPlusPayload ,
/ traits for this data: */PointPlusPayload_traits>
(data,count);
std::cout<<"build tree done"<<std::endl;

}
/* code */

I got 3 errors:

.cu：
error: qualifiers dropped in binding reference of type "PointPlusPayload &" to initializer of type "const PointPlusPayload"

builder_common.h(84): error: qualifiers dropped in binding reference of type "PointPlusPayload &" to initializer of type "const PointPlusPayload"
detected during:
instantiation of "void cukd::computeBounds_copyFirst<data_t,data_traits>(cukd::box_t<data_traits::point_t> *, const data_t *) [with data_t=PointPlusPayload, data_traits=PointPlusPayload_traits]"

builder_common.h(137): error: qualifiers dropped in binding reference of type "PointPlusPayload &" to initializer of type "const PointPlusPayload"
detected during:
instantiation of "void cukd::computeBounds_atomicGrow<data_t,data_traits>(cukd::box_t<data_traits::point_t> *, const data_t *, int) [with data_t=PointPlusPayload, data_traits=PointPlusPayload_traits]"

plz tell me how can i fix that? thanx!

Hi!
I am starting to play with your very interesting work. However, I sometimes get the wrong result. Here are some cases:

1. Using builder_thrust() on 90166 random points: the right closest point is found --> OK!
2. Using buildTree_inPlace() on the same 90166 points: the right closest point is found, but the returned ID is different. Normal?
3. Using builder_thrust() on 90167 points (same as above + 1 new): NO closest point is found :(
4. Using buildTree_inPlace() on the same 90167 points: the right closest point is found --> OK!
5. Using buildTree_bitonic() on those 90166 and 90167 points : NO closest point is found :(

I would have expected the tree to be exactly the same with the 3 builders. Should it be the case?
Do you know of an existing problem that could cause this behavior? Where to look for?

Thank you!
Claude

Notes:

• I can compare your result with my own "Closest Point" algorithm working on the CPU. Much slower but proven to be robust.
• I tried all tree query methods. It shows that the problem is with the builder.
• My points are generated at random between -100 and +100 in all 3 dimensions. I can provide them if needed.

Thanks for such a great job, but when I compile the project, an error occurs“spatial-kdtree.h:20:10: fatal error: cub/cub.cuh: No such file or directory
#include <cub/cub.cuh>
”
my cuda version is 10.1

Thanks for your great work! It's been incredibly useful for my projects.

I'm currently integrating it into a new application and need to keep track of the original order of d_points after calling buildTree. Could you advise on how to achieve this or if there's a way to map the original indices to their new positions after buildTree runs?

Appreciate any help you can provide!

What does this function do? i found enable this function but some case verify failed.

I played with your code and I found that the query speed (not the build speed) is affected by the distribution of points. Simply modifying generatePoints in the test code to achieve this (where p_ij is the j-th dimension of the i-th point, r is a random number between 0 and 1 obtained from drand48). Testing with 3M points and 25k query points.

• p_ij = r (baseline): 0.8ms
• p_ij = r * 100: 0.8ms
• p_ij = r + 100 * j: 0.8ms
• p_ij = r * 100 * j: 40ms! <-- here
• p_ij = r * 100 * j + j * j * 200: 40ms

For the time being in my application I will try to scale the points first, but I was wondering if this is a limitation of the algorithm, or whether this can be fixed?

Hi,

I have seen the example for building the kd-tree on host using cukd::buildTree_host.

Is it also possible to perform scan for the nearest neighbor on host, via cukd::stackFree::fcp or cukd::stackBased::fcp ?

Hi, apologies for the big text that follows!

I am trying to implement a find-closest-point in 4 dimensions using this library in the following way.

Initially, I create randomly a number of points in 4D that I use as my search space and then define a query-point that I use to find the closest neighbor using a brute-force O(N) implementation that looks like this:

float sqrtdiff( float4 p, float x, float y, float z, float w ){
    return std::sqrt(std::pow(p.x - x,2) + std::pow(p.y - y,2) + std::pow(p.z - z,2) + std::pow(p.w - w,2));
}

void find_closest( int &closest, float4 p, int numData, const float * x, const float * y, const float * z, const float * w ){
    float Ebest = 1e6;
    for( int i(0); i < numData; ++i ){
        float E = sqrtdiff(p,x[i],y[i],z[i],w[i]);
        if( E < Ebest ){
            Ebest = E;
            closest = i;
        }
    }
}

Then, using the same search data and query point I perform the equivalent implementation using this library as follows:

Initially, I call the routine the builds the k-d tree structure:
cukd::buildTree(data,numData);

and then I launch a very simple kernel with just one cuda thread to retrieve the closest neighbor of the same query-point:
find_closest_kernel<<<1,1>>>(p, data, numData);

the cuda kernal is the following:

__global__ void find_closest_kernel(float4 queryPoint, float4 *data, int numData) {
    int IDofClosestDataPoint = cukd::stackFree::fcp(queryPoint,data,numData);
    printf("cuda-closest: %d\n",IDofClosestDataPoint);
}

For example, with the following 20 search-points and the query point -> x=10, y=20, z=10, w=20
ID DISTANCE X Y Z W

0: 25.7876 26, 8, -2, 9
1: 73.3689 21, 6, 39, -45
2: 51.5946 -30, 18, 33, 43
3: 37.2425 29, 45, 11, 40
4: 71.225 -28, -36, -8, 33
5: 40.3113 32, -3, -14, 26
6: 52.2398 -38, 15, -6, 8
7: 89.4763 25, -36, 27, -46
8: 92.2009 49, -38, 14, -40
9: 66.1967 -43, -11, -14, 14
10: 79.0759 36, 18, -37, -38
11: 43.4396 -7, 5, 47, 18
12: 88.0114 -43, -6, 4, -45
13: 31.289 -9, 15, 2, 43
14: 63.3167 50, 22, -21, -18
15: 75.1332 -6, -48, 16, -7
16: 69.0652 -19, -41, 2, 32
17: 104.8 29, -49, 41, -50
18: 20.445 30, 24, 9, 19
19: 79.1833 -43, 37, -44, 36

The result I get is:
cuda-closest: 6

which is not what I expected since the closest point in search space is point with ID = 18

Did I use the library as I should, or do I miss something?