fixed a crucial bug in ptr2ind

csrc/cuda/convert_cuda.cu

@@ -57,7 +57,7 @@ torch::Tensor ptr2ind_cuda(torch::Tensor ptr, int64_t E) {
   auto ptr_data = ptr.data_ptr<int64_t>();
   auto out_data = out.data_ptr<int64_t>();
   auto stream = at::cuda::getCurrentCUDAStream();
-  ptr2ind_kernel<<<(ptr.numel() + THREADS - 1) / THREADS, THREADS, 0, stream>>>(
-      ptr_data, out_data, E, ptr.numel());
+  ptr2ind_kernel<<<(ptr.numel() - 1 + THREADS - 1) / THREADS, THREADS, 0,
+                   stream>>>(ptr_data, out_data, E, ptr.numel() - 1);
   return out;
 }

test/test_cat.py

@@ -39,7 +39,8 @@ def test_cat(device):
     assert out.storage.has_rowptr()
     assert out.storage.num_cached_keys() == 5
 
-    mat1 = mat1.set_value_(torch.randn((mat1.nnz(), 4), device=device))
+    value = torch.randn((mat1.nnz(), 4), device=device)
+    mat1 = mat1.set_value_(value, layout='coo')
     out = cat([mat1, mat1], dim=-1)
     assert out.storage.value().size() == (mat1.nnz(), 8)
     assert out.storage.has_row()

test/test_matmul.py

@@ -40,7 +40,7 @@ def test_spmm(dtype, device, reduce):
     out = matmul(src, other, reduce)
     out.backward(grad_out)
 
-    assert torch.allclose(expected, out)
+    assert torch.allclose(expected, out, atol=1e-6)
     assert torch.allclose(expected_grad_value, value.grad, atol=1e-6)
     assert torch.allclose(expected_grad_other, other.grad, atol=1e-6)