[gpu] Replace volatile shared memory with shfl_sync in KNNSearch #4306

haritha-j · 2020-08-05T19:01:57Z

Fix for distance calculation issue caused by the use of volatile shared memory in GPU K nearest neighbour search.

larshg · 2020-08-06T10:06:10Z

These lines here:

Lines 181 to 210 in b0ff2df

    
           //broadcast beg 
        
           if (active_lane == laneId) 
        
               *warp_buffer = batch.octree.begs[node_idx];                     
        
           int beg = *warp_buffer; 
        
           //broadcast end 
        
           if (active_lane == laneId) 
        
               *warp_buffer = batch.octree.ends[node_idx]; 
        
           int end = *warp_buffer; 
        
           float3 active_query; 
        
           volatile float* warp_buffer_float = (float*)&per_warp_buffer[warpId]; 
        
           //broadcast warp_query 
        
           if (active_lane == laneId) 
        
               *warp_buffer_float = query.x; 
        
           active_query.x = *warp_buffer_float; 
        
           if (active_lane == laneId) 
        
               *warp_buffer_float = query.y; 
        
           active_query.y = *warp_buffer_float; 
        
           if (active_lane == laneId) 
        
               *warp_buffer_float = query.z; 
        
           active_query.z = *warp_buffer_float;                             
        
           //broadcast query_index 
        
           if (active_lane == laneId) 
        
               *warp_buffer = query_index; 
        
           float active_query_index = *warp_buffer;

Can also be shared using the shfl instructions.

SergioRAgostinho · 2020-08-06T12:53:06Z

gpu/octree/src/cuda/knn_search.cu

@@ -222,87 +222,50 @@ namespace pcl { namespace device { namespace knn_search
        }

        template<int CTA_SIZE>
-		__device__ __forceinline__ int NearestWarpKernel(const float* points, int points_step, int length, const float3& active_query, float& dist)
+		__device__ __forceinline__ int NearestWarpKernel(const int beg, const int points_step, const int length, const float3& active_query, float& dist)


no out parameters if possible. use return variables. also remove the __forceinline__ from here. this method/function is huge.

Can we tidy up indentation in this function. it's looking slightly chaotic. run this function through clang-format.

should point_step be renamed to field_step it seems to measuring byte spacing between x, y, and z fields.

SergioRAgostinho · 2020-08-06T13:00:52Z

gpu/octree/src/cuda/knn_search.cu

-			float mind2 = dist2[tid];
+			//find minimum distance among warp threads
+            constexpr unsigned FULL_MASK = 0xFFFFFFFF;
+            static_assert(sizeof(KernelPolicy::WARP_SIZE) <= sizeof(unsigned int));


this check is wrong. WARP_SIZE represents the number of bits in the mask. sizeof gives you the number of bytes. you also need an error message here explaining the user what is wrong if the static_assert is violated.

Suggested change

static_assert(sizeof(KernelPolicy::WARP_SIZE) <= sizeof(unsigned int));

static_assert(KernelPolicy::WARP_SIZE <= 8*sizeof(unsigned int));

SergioRAgostinho · 2020-08-06T13:02:56Z

gpu/octree/src/cuda/knn_search.cu

+				float next = __shfl_down_sync(FULL_MASK, dist2, bit_offset);
+				int next_index = __shfl_down_sync(FULL_MASK, index, bit_offset);


SergioRAgostinho · 2020-08-06T13:03:49Z

gpu/octree/src/cuda/knn_search.cu

-			}
+			//retrieve index and dist2
+            index = __shfl_sync(FULL_MASK, index, 0);
+            dist2 = __shfl_sync(FULL_MASK, dist2, 0);


no need for this temp.

SergioRAgostinho · 2020-08-06T13:04:04Z

gpu/octree/src/cuda/knn_search.cu

+			//retrieve index and dist2
+            index = __shfl_sync(FULL_MASK, index, 0);
+            dist2 = __shfl_sync(FULL_MASK, dist2, 0);
+            dist = sqrt(dist2);


const (in case it stops being an out parameter)

SergioRAgostinho · 2020-08-06T13:09:33Z

gpu/octree/src/cuda/knn_search.cu

-            dist = sqrt(dist2[tid - lane]);
-			return index[tid - lane];
-		}          
+			return index;


Suggested change

return index;

return std::make_pair(index, dist);

…qualify variables

haritha-j · 2020-08-08T11:23:59Z

I'm recieving the pair returned from NearestWarpKernel using auto here, which makes the next bit of code a bit less readable, but std::tie doesn't seem to work properly with cuda __device__ functions, and auto [offset, index] is only supported from c++17 from what I read.

                const auto nearestPoint = NearestWarpKernel<KernelPolicy::CTA_SIZE>(beg, batch.points_step, end - beg, active_query);

                if (active_lane == laneId)
                    if (min_distance > nearestPoint.second)
                    {
                       min_distance = nearestPoint.second;
                       min_idx = beg + nearestPoint.first;
                    }

kunaltyagi · 2020-08-09T01:39:29Z

If the return type doesn't matter to the API/ABI (or explicitly needs to be given no API/ABI guarantee), you can use the following super-legal code

auto myFunction(int a, double b) -> decltype(auto) {
  struct GuessTheReturnType {
    bool correct = false;
    float magic = 4.2e1;
  };
  return GuessTheReturnType{a, b};
}

// usage
auto& val = myFunction(9, 33.);
if (val.correct) { return val.magic; }

SergioRAgostinho · 2020-08-10T11:27:26Z

@larshg please merge in case you're ok with the changes.

gpu/octree/src/cuda/knn_search.cu

larshg · 2020-08-10T19:22:46Z

https://github.com/PointCloudLibrary/pcl/blob/master/cmake/pcl_find_cuda.cmake

This one should probably also be revised soon. Not sure what cuda version is required for __shfl_sync instructions etc.

Right now it adds many cuda arch versions and with a new configure, it added 3.0 etch which failed to build.

SergioRAgostinho · 2020-08-10T20:09:37Z

https://github.com/PointCloudLibrary/pcl/blob/master/cmake/pcl_find_cuda.cmake

This one should probably also be revised soon. Not sure what cuda version is required for __shfl_sync instructions etc.

Right now it adds many cuda arch versions and with a new configure, it added 3.0 etch which failed to build.

Good catch. They were introduced in CUDA 9. CUDA C+14 support was also only introduced in CUDA 9.

@PointCloudLibrary/maintainers any opposition to bump CUDA's minimum version to 9?

larshg · 2020-08-11T06:21:13Z

Please go ahead and merge @SergioRAgostinho. I don't have those powers 😄

replace volatile shared memory with shfl_sync

49b80de

haritha-j added priority: gsoc Reason for prioritization module: gpu changelog: fix Meta-information for changelog generation labels Aug 5, 2020

haritha-j marked this pull request as draft August 6, 2020 07:06

const qualify and other fixes

b0ff2df

haritha-j marked this pull request as ready for review August 6, 2020 07:46

SergioRAgostinho requested changes Aug 6, 2020

View reviewed changes

haritha-j added 2 commits August 8, 2020 16:45

Replace read/writes to volatile storage by __shfl primitives & const …

33a8622

…qualify variables

fix static assertion

bf22852

haritha-j mentioned this pull request Aug 9, 2020

[gpu] Replace volatile shared memory with shfl_sync in ApproxNearestSearch #4313

Closed

SergioRAgostinho approved these changes Aug 10, 2020

View reviewed changes

larshg requested changes Aug 10, 2020

View reviewed changes

gpu/octree/src/cuda/knn_search.cu Outdated Show resolved Hide resolved

add warning for assertion

66fa90b

larshg approved these changes Aug 11, 2020

View reviewed changes

SergioRAgostinho merged commit c0c0cb2 into PointCloudLibrary:master Aug 11, 2020

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[gpu] Replace volatile shared memory with shfl_sync in KNNSearch #4306

[gpu] Replace volatile shared memory with shfl_sync in KNNSearch #4306

haritha-j commented Aug 5, 2020

larshg commented Aug 6, 2020 •

edited

Loading

SergioRAgostinho Aug 6, 2020

SergioRAgostinho Aug 6, 2020

SergioRAgostinho Aug 6, 2020 •

edited

Loading

SergioRAgostinho Aug 6, 2020

SergioRAgostinho Aug 6, 2020

SergioRAgostinho Aug 6, 2020

SergioRAgostinho Aug 6, 2020

SergioRAgostinho Aug 6, 2020

haritha-j commented Aug 8, 2020 •

edited

Loading

kunaltyagi commented Aug 9, 2020 •

edited

Loading

SergioRAgostinho commented Aug 10, 2020

larshg commented Aug 10, 2020

SergioRAgostinho commented Aug 10, 2020

larshg commented Aug 11, 2020

	static_assert(sizeof(KernelPolicy::WARP_SIZE) <= sizeof(unsigned int));
	static_assert(KernelPolicy::WARP_SIZE <= 8*sizeof(unsigned int));

		float next = __shfl_down_sync(FULL_MASK, dist2, bit_offset);
		int next_index = __shfl_down_sync(FULL_MASK, index, bit_offset);

[gpu] Replace volatile shared memory with shfl_sync in KNNSearch #4306

[gpu] Replace volatile shared memory with shfl_sync in KNNSearch #4306

Conversation

haritha-j commented Aug 5, 2020

larshg commented Aug 6, 2020 • edited Loading

SergioRAgostinho Aug 6, 2020

Choose a reason for hiding this comment

SergioRAgostinho Aug 6, 2020

Choose a reason for hiding this comment

SergioRAgostinho Aug 6, 2020 • edited Loading

Choose a reason for hiding this comment

SergioRAgostinho Aug 6, 2020

Choose a reason for hiding this comment

SergioRAgostinho Aug 6, 2020

Choose a reason for hiding this comment

SergioRAgostinho Aug 6, 2020

Choose a reason for hiding this comment

SergioRAgostinho Aug 6, 2020

Choose a reason for hiding this comment

SergioRAgostinho Aug 6, 2020

Choose a reason for hiding this comment

haritha-j commented Aug 8, 2020 • edited Loading

kunaltyagi commented Aug 9, 2020 • edited Loading

SergioRAgostinho commented Aug 10, 2020

larshg commented Aug 10, 2020

SergioRAgostinho commented Aug 10, 2020

larshg commented Aug 11, 2020

larshg commented Aug 6, 2020 •

edited

Loading

SergioRAgostinho Aug 6, 2020 •

edited

Loading

haritha-j commented Aug 8, 2020 •

edited

Loading

kunaltyagi commented Aug 9, 2020 •

edited

Loading