Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

fix bug for knnSearchGpu when knn!=ostride #374

Closed
wants to merge 7 commits into from

Conversation

yuyi1005
Copy link
Contributor

@yuyi1005 yuyi1005 commented May 9, 2018

knnSearchGpu gives wrong results if knn!=ostride.

For example, when we define the output matrix as
flann::Matrix idx(data, rows, 3, 4);
which means we need to find 3 nearest neighbor, but we want to save the results to a matrix of stride=4, then the result is wrong.

This bug can be fixed changing only one place in kdtree_cuda_3d_index.cu.

yuyi1005 added 7 commits May 9, 2018 19:34
fix bug for knnSearch when knn!=ostride
Now most nvidia gpu supports kernel concurrency, but flann always uses default stream, which means gpu can not process two indexs simutaneously when we use multithread. Creating cuda stream for each index can fix this problem.
fix errors in comment
knnSearch should be overrided. If not, Index initialize with KDTreeCuda3dIndexParams will call the default knnSearch, not the gpu one
flann::KDTreeCuda3dIndexParams params;
params["gpu_stream"] = mystream;//cudaStream_t mystream
params["gpu_stream"] = cudaStreamLegacy;//use default stream
params["gpu_stream"] = cudaStreamPerThread;//one stream per thread
if params["gpu_stream"] is not set, stream will be created each time calling buildIndex()
note that there is no internal synchronization, that is, cudaStreamSynchronize() should be called manually if another stream is used to process the matrices (only in case that matrices_in_gpu_ram==true)
-gencode=arch=compute_52,code=sm_52
-gencode=arch=compute_52,code=\"sm_52,compute_52\";
-gencode=arch=compute_61,code=\"sm_61,compute_61\"
@yuyi1005 yuyi1005 closed this Sep 6, 2018
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

1 participant