pytorch

55843321 - Wrap cub in its own namespace (#55292)

Wrap cub in its own namespace (#55292) Summary: Tentative fix for https://github.com/pytorch/pytorch/issues/55027. Wraps cub import in its name space so that static variables used by cub and thrust don't conflict if they end up in the different libraries when torch is built with BUILD_SPLIT_CUDA. cub variables end up in their own namespace, thrust variables are unwrapped, so they don't clash. This also allows extensions to use cub without wrapping it (thrust will still be problematic). The solution to allowing extensions to use thrust is to stop using thrust in pytorch completely. Now importing cub and importing thrust cannot coexist, so I had to move nonzero to its own file, and remove reliance on thrust functions for it. Nonzero now uses cub only. Also, we cannot selectively import just some of cub headers, we are forced to import `cub/cub.cuh`, which is not great. Caffe2 ops using cub are not touched (there are too many), so mixing caffe2 and torch will (can) still result in the same bug. We are moving towards disabling c2 ops, so I think this is fine. Still, even with that compiler (correctly) warns about redefinition of `CUB_NS_PREFIX` because including `ATen/ATen.h` transitively includes `thrust/complex.h` and that in turn includes original (empty) definition of `CUB_NS_PREFIX`. We probably can just ignore this warning. Here's an example warning: ``` In file included from /data/users/ngimel/pytorch/aten/src/ATen/native/cuda/Nonzero.cu:9: /data/users/ngimel/pytorch/aten/src/ATen/cuda/CubUtils.cuh:4: warning: "CUB_NS_PREFIX" redefined #define CUB_NS_PREFIX namespace at{ namespace native{ In file included from /home/ngimel/local/cuda/include/thrust/system/cuda/config.h:76, from /home/ngimel/local/cuda/include/thrust/system/cuda/detail/execution_policy.h:33, from /home/ngimel/local/cuda/include/thrust/iterator/detail/device_system_tag.h:23, from /home/ngimel/local/cuda/include/thrust/iterator/iterator_traits.h:111, from /home/ngimel/local/cuda/include/thrust/detail/type_traits/pointer_traits.h:23, from /home/ngimel/local/cuda/include/thrust/type_traits/is_contiguous_iterator.h:27, from /home/ngimel/local/cuda/include/thrust/type_traits/is_trivially_relocatable.h:19, from /home/ngimel/local/cuda/include/thrust/detail/complex/complex.inl:20, from /home/ngimel/local/cuda/include/thrust/complex.h:1031, from /data/users/ngimel/pytorch/c10/util/complex.h:9, from /data/users/ngimel/pytorch/c10/core/ScalarType.h:4, from /data/users/ngimel/pytorch/c10/core/Scalar.h:10, from /data/users/ngimel/pytorch/build/aten/src/ATen/core/TensorBody.h:8, from /data/users/ngimel/pytorch/aten/src/ATen/Tensor.h:3, from /data/users/ngimel/pytorch/aten/src/ATen/Context.h:4, from /data/users/ngimel/pytorch/aten/src/ATen/ATen.h:9, from /data/users/ngimel/pytorch/aten/src/ATen/native/cuda/Nonzero.cu:1: /home/ngimel/local/cuda/include/cub/util_namespace.cuh:43: note: this is the location of the previous definition #define CUB_NS_PREFIX ``` We will need a lint rule to prevent people from including `cub/cub.cuh`, because this will lead to https://github.com/pytorch/pytorch/issues/55027 reappearing again for some sequence of operations (and will lead to errors with cub code in extensions). Also, for this to work reliably we'll need to make sure that everything calling cub ends up in only one of libtorch_cuda_cu or libtorch_cuda_cpp, otherwise even namespace won't help (there still will be same symbols in 2 libraries). Upd: libtorch_cuda_cpp and libtorch_cuda_cu still contain the same symbols, which means that there exists a sequence of operations that will cause cache bug to reappear, so this is not a solution, we need to adjust file lists for BUILD_SPLITC_CUDA: ``` (pytorch) [ngimel@ ~/local/pytorch/build/lib] nm libtorch_cuda_cu.so | grep PerDeviceAttributeCache | c++filt 000000000c6bf808 u guard variable for at::native::cub::GetPerDeviceAttributeCache<at::native::cub::PtxVersionCacheTag>()::cache 000000000c600830 u guard variable for cub::GetPerDeviceAttributeCache<cub::PtxVersionCacheTag>()::cache 00000000018625e0 t at::native::cub::PerDeviceAttributeCache::DevicePayload at::native::cub::PerDeviceAttributeCache::operator()<at::native::cub::PtxVersion(int&)::{lambda(int&)https://github.com/pytorch/pytorch/issues/1}>(at::native::cub::PtxVersion(int&)::{lambda(int&)https://github.com/pytorch/pytorch/issues/1}&&, int) 00000000009ce630 t cub::PerDeviceAttributeCache::DevicePayload cub::PerDeviceAttributeCache::operator()<cub::PtxVersion(int&)::{lambda(int&)https://github.com/pytorch/pytorch/issues/1}>(cub::PtxVersion(int&)::{lambda(int&)https://github.com/pytorch/pytorch/issues/1}&&, int) 000000000c6bf820 u at::native::cub::GetPerDeviceAttributeCache<at::native::cub::PtxVersionCacheTag>()::cache 000000000c600840 u cub::GetPerDeviceAttributeCache<cub::PtxVersionCacheTag>()::cache (pytorch) [ngimel@ ~/local/pytorch/build/lib] nm libtorch_cuda_cpp.so | grep PerDeviceAttributeCache | c++filt 0000000000ad2d98 u guard variable for at::native::cub::GetPerDeviceAttributeCache<at::native::cub::PtxVersionCacheTag>()::cache 0000000000ad2da0 u at::native::cub::GetPerDeviceAttributeCache<at::native::cub::PtxVersionCacheTag>()::cache ``` Upd2: Moved TensorFactories.cu to torch_cuda_cu sources (see a change to caffe2/CMakeLists.txt), so now cub-related symbols are only in libtorch_cuda_cu. We'd need a test for that, any suggestions on how best to test it? cc zasdfgbnm malfet Pull Request resolved: https://github.com/pytorch/pytorch/pull/55292 Reviewed By: anjali411 Differential Revision: D27576442 Pulled By: ngimel fbshipit-source-id: 1ef29503a342bb214794d34a42a47052092a66c1