Commit fcd2eb7c193c6c8f41ab8d5c78b38765168f46fb
1 parent
ac48e0d0
added support for distance fields in stim::kdtree
Showing
5 changed files
with
135 additions
and
36 deletions
Show diff stats
stim/biomodels/network.h
... | ... | @@ -388,6 +388,22 @@ public: |
388 | 388 | return n; |
389 | 389 | } |
390 | 390 | |
391 | + //Copy the point cloud representing the centerline for the network into an array | |
392 | + void centerline_cloud(T* dst) { | |
393 | + size_t p; //stores the current edge point | |
394 | + size_t P; //stores the number of points in an edge | |
395 | + size_t i = 0; //index into the output array of points | |
396 | + for (size_t e = 0; e < E.size(); e++) { //for each edge in the network | |
397 | + P = E[e].size(); //get the number of points in this edge | |
398 | + for (p = 0; p < P; p++) { | |
399 | + dst[i * 3 + 0] = E[e][p][0]; | |
400 | + dst[i * 3 + 1] = E[e][p][1]; | |
401 | + dst[i * 3 + 2] = E[e][p][2]; | |
402 | + i++; | |
403 | + } | |
404 | + } | |
405 | + } | |
406 | + | |
391 | 407 | // gaussian function |
392 | 408 | float gaussianFunction(float x, float std=25){ return exp(-x/(2*std*std));} // by default std = 25 |
393 | 409 | |
... | ... | @@ -418,27 +434,27 @@ public: |
418 | 434 | /// @param sigma is the user-defined tolerance value - smaller values provide a stricter comparison |
419 | 435 | stim::network<T> compare(stim::network<T> A, float sigma, int device){ |
420 | 436 | |
421 | - stim::network<T> R; //generate a network storing the result of the comparison | |
422 | - R = (*this); //initialize the result with the current network | |
437 | + stim::network<T> R; //generate a network storing the result of the comparison | |
438 | + R = (*this); //initialize the result with the current network | |
423 | 439 | |
424 | - T *c; // centerline (array of double pointers) - points on kdtree must be double | |
440 | + T *c; // centerline (array of double pointers) - points on kdtree must be double | |
425 | 441 | size_t n_data = A.total_points(); // set the number of points |
426 | - c = (T*) malloc(sizeof(T) * n_data * 3); | |
442 | + c = (T*) malloc(sizeof(T) * n_data * 3); //allocate an array to store all points in the data set | |
427 | 443 | |
428 | 444 | unsigned t = 0; |
429 | - for(unsigned e = 0; e < A.E.size(); e++){ //for each edge in the network | |
430 | - for(unsigned p = 0; p < A.E[e].size(); p++){ //for each point in the edge | |
445 | + for(unsigned e = 0; e < A.E.size(); e++){ //for each edge in the network | |
446 | + for(unsigned p = 0; p < A.E[e].size(); p++){ //for each point in the edge | |
431 | 447 | for(unsigned d = 0; d < 3; d++){ //for each coordinate |
432 | 448 | |
433 | - c[t * 3 + d] = A.E[e][p][d]; | |
449 | + c[t * 3 + d] = A.E[e][p][d]; //copy the point into the array c | |
434 | 450 | } |
435 | 451 | t++; |
436 | 452 | } |
437 | 453 | } |
438 | 454 | |
439 | 455 | //generate a KD-tree for network A |
440 | - //float metric = 0.0; // initialize metric to be returned after comparing the network | |
441 | - size_t MaxTreeLevels = 3; // max tree level | |
456 | + //float metric = 0.0; // initialize metric to be returned after comparing the network | |
457 | + size_t MaxTreeLevels = 3; // max tree level | |
442 | 458 | |
443 | 459 | #ifdef __CUDACC__ |
444 | 460 | cudaSetDevice(device); |
... | ... | @@ -472,13 +488,13 @@ public: |
472 | 488 | stim::cpu_kdtree<T, 3> kdt; |
473 | 489 | kdt.create(c, n_data, MaxTreeLevels); |
474 | 490 | T *dists = new T[1]; // near neighbor distances |
475 | - size_t *nnIdx = new size_t[1]; // near neighbor indices // allocate near neigh indices | |
491 | + size_t *nnIdx = new size_t[1]; // near neighbor indices // allocate near neigh indices | |
476 | 492 | |
477 | 493 | stim::vec3<T> p0, p1; |
478 | 494 | T m1; |
479 | 495 | T* queryPt = new T[3]; |
480 | - for(unsigned e = 0; e < R.E.size(); e++){ //for each edge in A | |
481 | - R.E[e].add_mag(0); //add a new magnitude for the metric | |
496 | + for(unsigned e = 0; e < R.E.size(); e++){ //for each edge in A | |
497 | + R.E[e].add_mag(0); //add a new magnitude for the metric | |
482 | 498 | |
483 | 499 | for(unsigned p = 0; p < R.E[e].size(); p++){ //for each point in the edge |
484 | 500 | ... | ... |
stim/grids/image_stack.h
stim/math/vector.h
stim/structures/kdtree.cuh
... | ... | @@ -428,12 +428,6 @@ namespace stim { |
428 | 428 | |
429 | 429 | search<T, D>(nodes, indices, d_reference_points, d_query_points[idx], &d_indices[idx], &d_distances[idx], idx, next_nodes, next_search_nodes, Judge); // every query points are independent |
430 | 430 | } |
431 | - template <typename T, int D> | |
432 | - __host__ void aaboundingboxing(stim::aabbn<T, D> &bb, T *reference_points, size_t reference_count) { | |
433 | - for(size_t i = 0; i < reference_count; i++) { | |
434 | - bb.insert(&reference_points[i]); | |
435 | - } | |
436 | - } | |
437 | 431 | |
438 | 432 | template <typename T, int D = 3> |
439 | 433 | class cuda_kdtree { |
... | ... | @@ -441,7 +435,7 @@ namespace stim { |
441 | 435 | cuda_kdnode<T> *d_nodes; |
442 | 436 | size_t *d_index; |
443 | 437 | kdtree::point<T, D>* d_reference_points; |
444 | - size_t d_reference_count; | |
438 | + size_t npts; | |
445 | 439 | int num_nodes; |
446 | 440 | public: |
447 | 441 | ~cuda_kdtree() { |
... | ... | @@ -449,13 +443,18 @@ namespace stim { |
449 | 443 | HANDLE_ERROR(cudaFree(d_index)); |
450 | 444 | HANDLE_ERROR(cudaFree(d_reference_points)); |
451 | 445 | } |
452 | - void create(T *h_reference_points, size_t reference_count, size_t max_levels, stim::aabbn<T, D> &bb) { | |
446 | + | |
447 | + /// Create a KD-tree given a pointer to an array of reference points and the number of reference points | |
448 | + /// @param h_reference_points is a host array containing the reference points in (x0, y0, z0, ...., ) order | |
449 | + /// @param reference_count is the number of reference point in the array | |
450 | + /// @param max_levels is the deepest number of tree levels allowed | |
451 | + void create(T *h_reference_points, size_t reference_count, size_t max_levels = 3) { | |
453 | 452 | if (max_levels > 10) { |
454 | 453 | std::cout<<"The max_tree_levels should be smaller!"<<std::endl; |
455 | 454 | exit(1); |
456 | 455 | } |
457 | - bb.init(&h_reference_points[0]); | |
458 | - aaboundingboxing<T, D>(bb, h_reference_points, reference_count); | |
456 | + //bb.init(&h_reference_points[0]); | |
457 | + //aaboundingboxing<T, D>(bb, h_reference_points, reference_count); | |
459 | 458 | |
460 | 459 | std::vector <kdtree::point<T, D>> reference_points(reference_count); // restore the reference points in particular way |
461 | 460 | for (size_t j = 0; j < reference_count; j++) |
... | ... | @@ -465,14 +464,14 @@ namespace stim { |
465 | 464 | tree.cpu_create(reference_points, max_levels); // building a tree on cpu |
466 | 465 | kdtree::kdnode<T> *d_root = tree.get_root(); |
467 | 466 | num_nodes = tree.get_num_nodes(); |
468 | - d_reference_count = reference_count; // also equals to reference_count | |
467 | + npts = reference_count; // also equals to reference_count | |
469 | 468 | |
470 | 469 | HANDLE_ERROR(cudaMalloc((void**)&d_nodes, sizeof(cuda_kdnode<T>) * num_nodes)); // copy data from host to device |
471 | - HANDLE_ERROR(cudaMalloc((void**)&d_index, sizeof(size_t) * d_reference_count)); | |
472 | - HANDLE_ERROR(cudaMalloc((void**)&d_reference_points, sizeof(kdtree::point<T, D>) * d_reference_count)); | |
470 | + HANDLE_ERROR(cudaMalloc((void**)&d_index, sizeof(size_t) * npts)); | |
471 | + HANDLE_ERROR(cudaMalloc((void**)&d_reference_points, sizeof(kdtree::point<T, D>) * npts)); | |
473 | 472 | |
474 | 473 | std::vector < cuda_kdnode<T> > tmp_nodes(num_nodes); |
475 | - std::vector <size_t> indices(d_reference_count); | |
474 | + std::vector <size_t> indices(npts); | |
476 | 475 | std::vector <kdtree::kdnode<T>*> next_nodes; |
477 | 476 | size_t cur_pos = 0; |
478 | 477 | next_nodes.push_back(d_root); |
... | ... | @@ -511,6 +510,12 @@ namespace stim { |
511 | 510 | HANDLE_ERROR(cudaMemcpy(d_index, &indices[0], sizeof(size_t) * indices.size(), cudaMemcpyHostToDevice)); |
512 | 511 | HANDLE_ERROR(cudaMemcpy(d_reference_points, &reference_points[0], sizeof(kdtree::point<T, D>) * reference_points.size(), cudaMemcpyHostToDevice)); |
513 | 512 | } |
513 | + | |
514 | + /// Search the KD tree for nearest neighbors to a set of specified query points | |
515 | + /// @param h_query_points an array of query points in (x0, y0, z0, ...) order | |
516 | + /// @param query_count is the number of query points | |
517 | + /// @param indices are the indices to the nearest reference point for each query points | |
518 | + /// @param distances is an array containing the distance between each query point and the nearest reference point | |
514 | 519 | void search(T *h_query_points, size_t query_count, size_t *indices, T *distances) { |
515 | 520 | std::vector < typename kdtree::point<T, D> > query_points(query_count); |
516 | 521 | for (size_t j = 0; j < query_count; j++) |
... | ... | @@ -536,7 +541,7 @@ namespace stim { |
536 | 541 | HANDLE_ERROR(cudaMalloc((void**)&next_search_nodes, threads * blocks * 50 * sizeof(int))); |
537 | 542 | HANDLE_ERROR(cudaMemcpy(d_query_points, &query_points[0], sizeof(T) * query_points.size() * D, cudaMemcpyHostToDevice)); |
538 | 543 | |
539 | - search_batch<<<threads, blocks>>> (d_nodes, d_index, d_reference_points, d_query_points, query_points.size(), d_indices, d_distances, next_nodes, next_search_nodes, Judge); | |
544 | + search_batch<<<blocks, threads>>> (d_nodes, d_index, d_reference_points, d_query_points, query_points.size(), d_indices, d_distances, next_nodes, next_search_nodes, Judge); | |
540 | 545 | |
541 | 546 | if (Judge == NULL) { // do the following work if the thread works safely |
542 | 547 | HANDLE_ERROR(cudaMemcpy(indices, d_indices, sizeof(size_t) * query_points.size(), cudaMemcpyDeviceToHost)); |
... | ... | @@ -549,6 +554,63 @@ namespace stim { |
549 | 554 | HANDLE_ERROR(cudaFree(d_indices)); |
550 | 555 | HANDLE_ERROR(cudaFree(d_distances)); |
551 | 556 | } |
557 | + | |
558 | + /// Return the number of points in the KD tree | |
559 | + size_t num_points() { | |
560 | + return npts; | |
561 | + } | |
562 | + | |
563 | + stim::aabbn<T, D> getbox() { | |
564 | + size_t N = npts; | |
565 | + //std::vector < typename kdtree::point<T, D> > cpu_ref(npts); //allocate space on the CPU for the reference points | |
566 | + T* cpu_ref = (T*)malloc(N * D * sizeof(T)); //allocate space on the CPU for the reference points | |
567 | + HANDLE_ERROR(cudaMemcpy(cpu_ref, d_reference_points, N * D * sizeof(T), cudaMemcpyDeviceToHost)); //copy from GPU to CPU | |
568 | + | |
569 | + stim::aabbn<T, D> bb(cpu_ref); | |
570 | + | |
571 | + for (size_t i = 1; i < N; i++) { //for each reference point | |
572 | + //std::cout << "( " << cpu_ref[i * D + 0] << ", " << cpu_ref[i * D + 1] << ", " << cpu_ref[i * D + 2] << ")" << std::endl; | |
573 | + bb.insert(&cpu_ref[i * D]); | |
574 | + } | |
575 | + return bb; | |
576 | + } | |
577 | + | |
578 | + //generate an implicit distance field for the KD-tree | |
579 | + void dist_field3(T* dist, size_t* dims, stim::aabbn<T, 3> bb) { | |
580 | + size_t N = 1; //number of query points that make up the distance field | |
581 | + for (size_t d = 0; d < 3; d++) N *= dims[d]; //calculate the total number of query points | |
582 | + | |
583 | + //calculate the grid spatial parameters | |
584 | + T dx = 0; | |
585 | + if (dims[0] > 1) dx = bb.length(0) / dims[0]; | |
586 | + T dy = 0; | |
587 | + if (dims[1] > 1) dy = bb.length(1) / dims[1]; | |
588 | + T dz = 0; | |
589 | + if (dims[2] > 1) dz = bb.length(2) / dims[2]; | |
590 | + | |
591 | + T* Q = (T*)malloc(N * 3 * sizeof(T)); //allocate space for the query points | |
592 | + size_t i; | |
593 | + for (size_t z = 0; z < dims[2]; z++) { //for each query point (which is a point in the grid) | |
594 | + for (size_t y = 0; y < dims[1]; y++) { | |
595 | + for (size_t x = 0; x < dims[0]; x++) { | |
596 | + i = z * dims[1] * dims[0] + y * dims[0] + x; | |
597 | + Q[i * 3 + 0] = bb.low[0] + x * dx + dx / 2; | |
598 | + Q[i * 3 + 1] = bb.low[1] + y * dy + dy / 2; | |
599 | + Q[i * 3 + 2] = bb.low[2] + z * dz + dz / 2; | |
600 | + //std::cout << i<<" "<<Q[i * 3 + 0] << " " << Q[i * 3 + 1] << " " << Q[i * 3 + 2] << std::endl; | |
601 | + } | |
602 | + } | |
603 | + } | |
604 | + size_t* temp = (size_t*)malloc(N * sizeof(size_t)); //allocate space to store the indices (unused) | |
605 | + search(Q, N, temp, dist); | |
606 | + } | |
607 | + | |
608 | + //generate an implicit distance field for the KD-tree | |
609 | + void dist_field3(T* dist, size_t* dims) { | |
610 | + stim::aabbn<T, D> bb = getbox(); //get a bounding box around the tree | |
611 | + dist_field3(dist, dims, bb); | |
612 | + } | |
613 | + | |
552 | 614 | }; |
553 | 615 | } //end namespace stim |
554 | 616 | #endif |
555 | 617 | \ No newline at end of file | ... | ... |
stim/visualization/aabbn.h
... | ... | @@ -15,17 +15,22 @@ struct aabbn{ |
15 | 15 | T low[D]; //top left corner position |
16 | 16 | T high[D]; //dimensions along x and y and z |
17 | 17 | |
18 | -//public: | |
19 | - | |
20 | - CUDA_CALLABLE aabbn(){ //initialize an axis aligned bounding box of size 0 at the given position | |
21 | - for(size_t d = 0; d < D; d++) | |
22 | - low[d] = high[d] = 0; | |
18 | + CUDA_CALLABLE void init(T* i) { | |
19 | + for (size_t d = 0; d < D; d++) | |
20 | + low[d] = high[d] = i[d]; | |
23 | 21 | } |
24 | 22 | |
25 | - CUDA_CALLABLE void init(T* init){ | |
26 | - for(size_t d = 0; d < D; d++) | |
27 | - low[d] = high[d] = init[d]; | |
23 | +//public: | |
24 | + | |
25 | + //CUDA_CALLABLE aabbn(){ //initialize an axis aligned bounding box of size 0 at the given position | |
26 | + // for(size_t d = 0; d < D; d++) | |
27 | + // low[d] = high[d] = 0; | |
28 | + //} | |
29 | + CUDA_CALLABLE aabbn() {} | |
30 | + CUDA_CALLABLE aabbn(T* i) { | |
31 | + init(i); | |
28 | 32 | } |
33 | + | |
29 | 34 | |
30 | 35 | //insert a point into the bounding box, growing the box appropriately |
31 | 36 | CUDA_CALLABLE void insert(T* p){ |
... | ... | @@ -46,6 +51,21 @@ struct aabbn{ |
46 | 51 | if(low[d] > b[d]) low[d] = b[d]; |
47 | 52 | } |
48 | 53 | |
54 | + CUDA_CALLABLE T length(size_t d) { | |
55 | + return high[d] - low[d]; | |
56 | + } | |
57 | + | |
58 | + CUDA_CALLABLE aabbn<T, D> operator*(T s) { | |
59 | + aabbn<T, D> newbox; | |
60 | + for (size_t d = 0; d < D; d++) { | |
61 | + T c = (low[d] + high[d]) / 2; | |
62 | + T l = high[d] - low[d]; | |
63 | + newbox.low[d] = c - l * s / 2; | |
64 | + newbox.high[d] = c + l * s / 2; | |
65 | + } | |
66 | + return newbox; | |
67 | + } | |
68 | + | |
49 | 69 | }; |
50 | 70 | |
51 | 71 | } | ... | ... |