add origin code

LICENSE

+Copyright (c) 2018 Matthias Fey <matthias.fey@tu-dortmund.de>
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.

MANIFEST.in

+include LICENSE
+recursive-include cuda *

README.md

-# torch-bincount
+[pypi-image]: https://badge.fury.io/py/torch-bincount.svg
+[pypi-url]: https://pypi.python.org/pypi/torch-bincount
+[build-image]: https://travis-ci.org/rusty1s/pytorch_bincount.svg?branch=master
+[build-url]: https://travis-ci.org/rusty1s/pytorch_bincount
+[coverage-image]: https://codecov.io/gh/rusty1s/pytorch_bincount/branch/master/graph/badge.svg
+[coverage-url]: https://codecov.io/github/rusty1s/pytorch_bincount?branch=master
 
+# PyTorch BinCount
+
+[![PyPI Version][pypi-image]][pypi-url]
+[![Build Status][build-image]][build-url]
+[![Code Coverage][coverage-image]][coverage-url]
+
+--------------------------------------------------------------------------------
+
+**PyTorch 0.4.1 now supports [`bincount`](https://pytorch.org/docs/stable/torch.html#torch.bincount) both for CPU and GPU.
+Therefore, this package is no longer needed and will not be updated.**
+
+--------------------------------------------------------------------------------
+
+This package consists of a small extension library of a highly optimized `bincount` operation for the use in [PyTorch](http://pytorch.org/), which is missing in the main package.
+The operation works on varying data types and is implemented both for CPU and GPU.
+
+## Installation
+
+Ensure that at least PyTorch 0.4.1 is installed and verify that `cuda/bin` and `cuda/include` are in your `$PATH` and `$CPATH` respectively, *e.g.*:
+
+```
+$ python -c "import torch; print(torch.__version__)"
+>>> 0.4.1
+
+$ echo $PATH
+>>> /usr/local/cuda/bin:...
+
+$ echo $CPATH
+>>> /usr/local/cuda/include:...
+```
+
+Then run:
+
+```
+pip install torch-bincount
+```
+
+If you are running into any installation problems, please create an [issue](https://github.com/rusty1s/pytorch_bincount/issues).
+Be sure to import `torch` first before using this package to resolve symbols the dynamic linker must see.
+
+## Usage
+
+```
+torch_bincount.bincount(src, size=None) -> LongTensor
+```
+
+Counts the number of occurrences of each value in a non-negative tensor.
+
+### Parameters
+
+* **src** *(Tensor)* - The input tensor.
+* **size** *(int, optional)* - The maximum number of bins for the output array. (default: `None`)
+
+### Returns
+
+* **out** *(LongTensor)* - The result of binning the input tensor.
+
+### Example
+
+```py
+import torch
+from torch_bincount import bincount
+
+src = torch.tensor([2, 1, 1, 2, 4, 4, 2])
+out = bincount(src)
+```
+
+```
+print(out)
+tensor([ 0,  2,  3,  0,  2])
+```
+
+## Running tests
+
+```
+python setup.py test
+```

cuda/bincount.cpp

+#include <torch/torch.h>
+
+#define CHECK_CUDA(x) AT_ASSERTM(x.type().is_cuda(), #x " must be CUDA tensor")
+
+at::Tensor bincount_cuda(at::Tensor src, int64_t size);
+
+at::Tensor bincount(at::Tensor src, int64_t size) {
+  CHECK_CUDA(src);
+  return bincount_cuda(src, size);
+}
+
+PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
+  m.def("bincount", &bincount, "BinCount (CUDA)");
+}

cuda/bincount_kernel.cu

+#include <ATen/ATen.h>
+#include <THC/THCAtomics.cuh>
+
+#define THREADS 1024
+#define BLOCKS(N) (N + THREADS - 1) / THREADS
+
+template <typename scalar_t>
+__global__ void bincount_cuda_kernel(scalar_t *__restrict__ src, int64_t *out,
+                                     size_t numel) {
+  const size_t index = blockIdx.x * blockDim.x + threadIdx.x;
+  const size_t stride = blockDim.x * gridDim.x;
+  for (ptrdiff_t i = index; i < numel; i += stride) {
+    atomicAdd(out + (ptrdiff_t)src[i], 1);
+  }
+}
+
+at::Tensor bincount_cuda(at::Tensor src, int64_t size) {
+  auto out = at::zeros(size, src.type().toScalarType(at::kLong));
+
+  AT_DISPATCH_ALL_TYPES(src.type(), "bincount_cuda_kernel", [&] {
+    bincount_cuda_kernel<scalar_t><<<BLOCKS(src.numel()), THREADS>>>(
+        src.data<scalar_t>(), out.data<int64_t>(), src.numel());
+  });
+
+  return out;
+}

setup.cfg

+[metadata]
+description-file = README.md
+
+[aliases]
+test = pytest
+
+[tool:pytest]
+addopts = --capture=no --cov

setup.py

+from setuptools import setup, find_packages
+import torch
+from torch.utils.cpp_extension import BuildExtension, CUDAExtension
+
+__version__ = '0.1.1'
+url = 'https://github.com/rusty1s/pytorch_bincount'
+
+install_requires = ['numpy']
+setup_requires = ['pytest-runner']
+tests_require = ['pytest', 'pytest-cov']
+ext_modules = []
+cmdclass = {}
+
+if torch.cuda.is_available():
+    ext_modules += [
+        CUDAExtension('bincount_cuda',
+                      ['cuda/bincount.cpp', 'cuda/bincount_kernel.cu'])
+    ]
+    cmdclass['build_ext'] = BuildExtension
+
+setup(
+    name='torch_bincount',
+    version=__version__,
+    description='Optimized PyTorch BinCount Operation',
+    author='Matthias Fey',
+    author_email='matthias.fey@tu-dortmund.de',
+    url=url,
+    download_url='{}/archive/{}.tar.gz'.format(url, __version__),
+    keywords=['pytorch', 'bincount', 'python'],
+    install_requires=install_requires,
+    setup_requires=setup_requires,
+    tests_require=tests_require,
+    ext_modules=ext_modules,
+    cmdclass=cmdclass,
+    packages=find_packages(),
+)

test/test_bincount.py

+from itertools import product
+
+import pytest
+from torch_bincount import bincount
+
+from .utils import dtypes, devices, tensor
+
+
+@pytest.mark.parametrize('dtype,device', product(dtypes, devices))
+def test_bincount(dtype, device):
+    src = tensor([2, 1, 1, 2, 4, 4, 2], dtype, device)
+    out = bincount(src)
+
+    assert out.tolist() == [0, 2, 3, 0, 2]

test/utils.py

+import torch
+from torch.testing import get_all_dtypes
+
+dtypes = get_all_dtypes()
+dtypes.remove(torch.half)
+
+devices = [torch.device('cpu')]
+if torch.cuda.is_available():
+    devices += [torch.device('cuda:{}'.format(torch.cuda.current_device()))]
+
+
+def tensor(x, dtype, device):
+    return None if x is None else torch.tensor(x, dtype=dtype, device=device)

torch_bincount/__init__.py

+from .bincount import bincount
+
+__version__ = '0.1.1'
+
+__all__ = ['bincount', '__version__']

torch_bincount/bincount.py

+import torch
+import numpy as np
+
+if torch.cuda.is_available():
+    import bincount_cuda
+
+
+def bincount(src, size=None):
+    """Counts the number of occurrences of each value in a non-negative tensor.
+
+    Args:
+        src (Tensor): The input tensor.
+        size (int, optional): The maximum number of bins for the output array.
+            (default: `None`)
+
+    :rtype: :class:`LongTensor`
+    """
+
+    src = src.contiguous().view(-1)
+
+    if src.is_cuda:
+        size = src.max() + 1 if size is None else size
+        return bincount_cuda.bincount(src, size)
+    else:
+        out = np.bincount(src.long().numpy(), minlength=size)
+        return torch.from_numpy(out)