'still a bug'

Author: tchaton

cuda/src/ball_query_gpu.cu

@@ -68,15 +68,17 @@ __global__ void query_ball_point_kernel_partial_dense(int size_x,
 	const ptrdiff_t end_idx_y = batch_y[batch_idx + 1];
 	float radius2 = radius * radius;
 
-	for (ptrdiff_t n_x = start_idx_x + idx; n_x < end_idx_x + 1; n_x += THREADS) {
+	for (ptrdiff_t n_x = start_idx_x + idx; n_x < end_idx_x; n_x += THREADS) {
+		printf("n_x: %d \n", n_x);
+
 		int64_t count = 0;
-		for (ptrdiff_t n_y = start_idx_y; n_y < end_idx_y + 1; n_y++) {
+		for (ptrdiff_t n_y = start_idx_y; n_y < end_idx_y; n_y++) {
 			float dist = 0;
 			for (ptrdiff_t d = 0; d < 3; d++) {
 				dist += (x[n_x * 3 + d] - y[n_y * 3 + d]) *
 					(x[n_x * 3 + d] - y[n_y * 3 + d]);
 			}
-			printf("Hello from (%d, %d) block, %d, thread %d\n", n_x, n_y, blockIdx.x, threadIdx.x);
+			printf("n_x: %d, n_y: %d \n", n_x, n_y);
 			if(dist <= radius2){
 				idx_out[n_x * nsample + count] = n_y;
 				dist_out[n_x * nsample + count] = dist;

test/test_ballquerry_partial.py

@@ -18,7 +18,7 @@ def test_simple_gpu(self):
         batch_x = torch.from_numpy(np.asarray([0, 0, 1, 1])).long().cuda()
         batch_y = torch.from_numpy(np.asarray([0])).long().cuda()
 
-        idx, dist2 = ball_query(1, 2, x, y, batch_x, batch_y, mode="PARTIAL_DENSE")
+        idx, dist2 = ball_query(1., 2, x, y, batch_x, batch_y, mode="PARTIAL_DENSE")
 
         idx = idx.detach().cpu().numpy()
         dist2 = dist2.detach().cpu().numpy()

torch_points/torchpoints.py

@@ -330,7 +330,7 @@ def ball_query_partial_dense(radius, nsample, x, y, batch_x, batch_y):
     """
     return BallQueryPartialDense.apply(radius, nsample, x, y, batch_x, batch_y)
 
-def ball_query(radius, nsample, x, y, batch_x=None, batch_y=None, mode=None):
+def ball_query(radius: float, nsample: int, x, y, batch_x=None, batch_y=None, mode=None):
     if mode is None:
         raise Exception('The mode should be defined within ["PARTIAL_DENSE | DENSE"]')