diff --git a/torchsparse/backend/convolution/convolution_gather_scatter_cuda.cu b/torchsparse/backend/convolution/convolution_gather_scatter_cuda.cu index 6a6bd63..233a4cb 100644 --- a/torchsparse/backend/convolution/convolution_gather_scatter_cuda.cu +++ b/torchsparse/backend/convolution/convolution_gather_scatter_cuda.cu @@ -547,7 +547,7 @@ at::Tensor conv_forward_gather_scatter_cuda_latest( // all gather AT_DISPATCH_FLOATING_TYPES_AND_HALF( - in_feat.type(), "conv_forward_gather_scatter_cuda", ([&] { + in_feat.scalar_type(), "conv_forward_gather_scatter_cuda", [&] { gather_all_kernel_pad_sep_with_mask <<(), input_mask.data_ptr(), output_mask.data_ptr(), transpose, precompute_mid); - })); + }); at::Tensor in_buffer_activated, out_buffer_activated, kernel_buffer; int buffer_st; @@ -779,14 +779,14 @@ at::Tensor conv_forward_gather_scatter_cuda_fallback( // gather n_active_feats dense features from N sparse input features with c // feature dimensions AT_DISPATCH_FLOATING_TYPES_AND_HALF( - in_feat.type(), "conv_forward_gather_scatter_cuda", ([&] { + in_feat.scalar_type(), "conv_forward_gather_scatter_cuda", [&] { gather_kernel <<>>( n_active_feats, n_in_feats, n_in_channels, in_feat.data_ptr(), in_buffer_activated.data_ptr(), neighbor_map.data_ptr() + cur_offset, transpose); - })); + }); // gemm: (i, c) X (c, o) = (i, o) int kmap_idx = i; if (conv_mode == 2) { @@ -796,14 +796,14 @@ at::Tensor conv_forward_gather_scatter_cuda_fallback( // scatter n_active_feats dense features into n_out_feats output features of // dimension n_out_channels AT_DISPATCH_FLOATING_TYPES_AND_HALF( - in_feat.type(), "conv_forward_gather_scatter_cuda", ([&] { + in_feat.scalar_type(), "conv_forward_gather_scatter_cuda", [&] { scatter_kernel <<>>( n_active_feats, n_out_feats, n_out_channels, out_buffer_activated.data_ptr(), out_feat.data_ptr(), neighbor_map.data_ptr() + cur_offset, transpose); - })); + }); cur_offset += 2 * n_active_feats; } @@ -877,23 +877,23 @@ void conv_backward_gather_scatter_cuda(at::Tensor in_feat, at::Tensor grad_in_fe } // gather AT_DISPATCH_FLOATING_TYPES_AND_HALF( - in_feat.type(), "conv_forward_gather_scatter_cuda", ([&] { + in_feat.scalar_type(), "conv_forward_gather_scatter_cuda", [&] { gather_kernel <<>>( n_active_feats, n_out_feats, n_out_channels, grad_out_feat.data_ptr(), out_grad_buffer_activated.data_ptr(), neighbor_map.data_ptr() + cur_offset, !transpose); - })); + }); AT_DISPATCH_FLOATING_TYPES_AND_HALF( - in_feat.type(), "conv_forward_gather_scatter_cuda", ([&] { + in_feat.type(), "conv_forward_gather_scatter_cuda", [&] { gather_kernel <<>>( n_active_feats, n_in_feats, n_in_channels, in_feat.data_ptr(), in_buffer_activated.data_ptr(), neighbor_map.data_ptr() + cur_offset, transpose); - })); + }); // gemm torch::mm_out(in_grad_buffer_activated, out_grad_buffer_activated, torch::transpose(kernel[i], 0, 1)); @@ -902,14 +902,14 @@ void conv_backward_gather_scatter_cuda(at::Tensor in_feat, at::Tensor grad_in_fe out_grad_buffer_activated); // scatter AT_DISPATCH_FLOATING_TYPES_AND_HALF( - in_feat.type(), "conv_forward_gather_scatter_cuda", ([&] { + in_feat.scalar_type(), "conv_forward_gather_scatter_cuda", [&] { scatter_kernel <<>>( n_active_feats, n_in_feats, n_in_channels, in_grad_buffer_activated.data_ptr(), grad_in_feat.data_ptr(), neighbor_map.data_ptr() + cur_offset, !transpose); - })); + }); cur_offset += 2 * n_active_feats; } } diff --git a/torchsparse/backend/devoxelize/devoxelize_cuda.cu b/torchsparse/backend/devoxelize/devoxelize_cuda.cu index c2c0423..1362a4e 100644 --- a/torchsparse/backend/devoxelize/devoxelize_cuda.cu +++ b/torchsparse/backend/devoxelize/devoxelize_cuda.cu @@ -68,11 +68,11 @@ at::Tensor devoxelize_forward_cuda(const at::Tensor feat, torch::zeros({N, c}, at::device(feat.device()).dtype(feat.dtype())); AT_DISPATCH_FLOATING_TYPES_AND_HALF( - feat.type(), "devoxelize_forward_cuda", ([&] { + feat.scalar_type(), "devoxelize_forward_cuda", [&] { devoxelize_forward_kernel<<>>( N, c, indices.data_ptr(), weight.data_ptr(), feat.data_ptr(), out.data_ptr()); - })); + }); return out; } @@ -88,11 +88,11 @@ at::Tensor devoxelize_backward_cuda(const at::Tensor top_grad, {n, c}, at::device(top_grad.device()).dtype(top_grad.dtype())); AT_DISPATCH_FLOATING_TYPES_AND_HALF( - top_grad.type(), "devoxelize_backward_cuda", ([&] { + top_grad.scalar_type(), "devoxelize_backward_cuda", [&] { devoxelize_backward_kernel<<>>( N, n, c, indices.data_ptr(), weight.data_ptr(), top_grad.data_ptr(), bottom_grad.data_ptr()); - })); + }); return bottom_grad; } diff --git a/torchsparse/backend/voxelize/voxelize_cuda.cu b/torchsparse/backend/voxelize/voxelize_cuda.cu index 3ae63f2..c4aad21 100644 --- a/torchsparse/backend/voxelize/voxelize_cuda.cu +++ b/torchsparse/backend/voxelize/voxelize_cuda.cu @@ -86,10 +86,10 @@ at::Tensor voxelize_forward_cuda(const at::Tensor inputs, const at::Tensor idx, torch::zeros({N1, c}, at::device(idx.device()).dtype(inputs.dtype())); AT_DISPATCH_FLOATING_TYPES_AND_HALF( - inputs.type(), "voxelize_forward_cuda", ([&] + inputs.scalar_type(), "voxelize_forward_cuda", [&] { voxelize_forward_kernel<<>>( N, c, N1, inputs.data_ptr(), idx.data_ptr(), - counts.data_ptr(), out.data_ptr()); })); + counts.data_ptr(), out.data_ptr()); }); return out; } @@ -105,10 +105,10 @@ at::Tensor voxelize_backward_cuda(const at::Tensor top_grad, torch::zeros({N, c}, at::device(idx.device()).dtype(top_grad.dtype())); AT_DISPATCH_FLOATING_TYPES_AND_HALF( - top_grad.type(), "voxelize_backward_cuda", ([&] + top_grad.scalar_type(), "voxelize_backward_cuda", [&] { voxelize_backward_kernel<<>>( N, c, N1, top_grad.data_ptr(), idx.data_ptr(), - counts.data_ptr(), bottom_grad.data_ptr()); })); + counts.data_ptr(), bottom_grad.data_ptr()); }); return bottom_grad; } @@ -120,10 +120,10 @@ void to_dense_forward_cuda(const at::Tensor inputs, const at::Tensor idx, int c = inputs.size(1); AT_DISPATCH_FLOATING_TYPES_AND_HALF( - inputs.type(), "to_dense_forward_cuda", ([&] + inputs.scalar_type(), "to_dense_forward_cuda", [&] { to_dense_forward_kernel<<<(N * c + 255) / 256, 256>>>( N, c, inputs.data_ptr(), idx.data_ptr(), - range.data_ptr(), outputs.data_ptr()); })); + range.data_ptr(), outputs.data_ptr()); }); } void to_dense_backward_cuda(const at::Tensor top_grad, @@ -134,8 +134,8 @@ void to_dense_backward_cuda(const at::Tensor top_grad, int c = bottom_grad.size(1); AT_DISPATCH_FLOATING_TYPES_AND_HALF( - top_grad.type(), "to_dense_backward_cuda", ([&] + top_grad.scalar_type(), "to_dense_backward_cuda", [&] { to_dense_backward_kernel<<<(N * c + 255) / 256, 256>>>( N, c, top_grad.data_ptr(), idx.data_ptr(), - range.data_ptr(), bottom_grad.data_ptr()); })); + range.data_ptr(), bottom_grad.data_ptr()); }); }