trivialfis
commented
3 years ago cc @hcho3
Closed trivialfis closed 3 years ago
trivialfis
commented
3 years ago cc @hcho3
kkraus14
commented
3 years ago @trivialfis any chance you can run the workload through cuda-memcheck and see if it pops out anything actionable? Right now there's not really any information for us to act on here.
trivialfis
commented
3 years ago @kkraus14 Thanks for the reply, I will try to run it next week.
trivialfis
commented
3 years ago Em .. running it locally with memcheck/sanitizer on my 2 GPUs system works fine, but it's somehow easier to reproduce on CI (another log). But even on CI, it's still not deterministically reproducible. I will try to run it on other clusters.
kkraus14
commented
3 years ago Hmm, well it looks to be concatenation related:
[2021-03-29T12:36:59.078Z] Current thread 0x00007f0381574740 (most recent call first):
[2021-03-29T12:36:59.078Z] File "/opt/python/envs/gpu_test/lib/python3.7/site-packages/cudf/core/column/column.py", line 278 in _concat
[2021-03-29T12:36:59.078Z] File "/opt/python/envs/gpu_test/lib/python3.7/site-packages/cudf/core/series.py", line 1734 in _concat
[2021-03-29T12:36:59.078Z] File "/opt/python/envs/gpu_test/lib/python3.7/site-packages/cudf/core/reshape.py", line 378 in concatLet's see if it's caused by https://github.com/dmlc/xgboost/pull/6798 .
trivialfis
commented
3 years ago Let's see if it's caused by dmlc/xgboost#6798 .
Nope.
trivialfis
commented
3 years ago Hi @kkraus14 . I got a memcheck log by running it on our CI. Hopefully that can clear things up a little bit. log.txt
Full log in here: https://xgboost-ci.net/blue/rest/organizations/jenkins/pipelines/xgboost/branches/PR-6815/runs/3/nodes/226/log/?start=0
CUAD: 10.2 Conda env:
conda create -n gpu_test -c rapidsai-nightly -c rapidsai -c nvidia -c conda-forge -c defaults \
python=3.7 cudf=0.18* rmm=0.18* cudatoolkit=10.2 dask dask-cuda dask-cudf cupy \
numpy pytest scipy scikit-learn pandas matplotlib wheel python-kubernetes urllib3 graphviz hypothesisRelevant piece:
2021-03-31T19:09:52.917Z] ========= Invalid __global__ read of size 8
[2021-03-31T19:09:52.917Z] ========= at 0x00000460 in void cudf::detail::fused_concatenate_kernel<long, int=256, bool=0>(cudf::column_device_view const *, unsigned long const *, int, cudf::mutable_column_device_view, int*)
[2021-03-31T19:09:52.917Z] ========= by thread (19,0,0) in block (2,0,0)
[2021-03-31T19:09:52.917Z] ========= Address 0x7f471c0006f8 is out of bounds
[2021-03-31T19:09:52.917Z] ========= Device Frame:void cudf::detail::fused_concatenate_kernel<long, int=256, bool=0>(cudf::column_device_view const *, unsigned long const *, int, cudf::mutable_column_device_view, int*) (void cudf::detail::fused_concatenate_kernel<long, int=256, bool=0>(cudf::column_device_view const *, unsigned long const *, int, cudf::mutable_column_device_view, int*) : 0x460)The used input is a small dataset from skl:
def test_boost_from_prediction(local_cuda_cluster: LocalCUDACluster) -> None:
import cudf
from sklearn.datasets import load_breast_cancer
with Client(local_cuda_cluster) as client:
X_, y_ = load_breast_cancer(return_X_y=True)
X = dd.from_array(X_, chunksize=100).map_partitions(cudf.from_pandas)
y = dd.from_array(y_, chunksize=100).map_partitions(cudf.from_pandas)Concatenation is used to get all data parts within each worker into 1.
kkraus14
commented
3 years ago @trivialfis how many workers is CI running in this situation? Could you point us to the code that xgboost uses to concatenate the parts on each worker?
trivialfis
commented
3 years ago Thanks.
how many workers is CI running in this situation?
8 GPUs I believe.
Could you point us to the code that xgboost uses to concatenate the parts on each worker?
This is the concat function: https://github.com/dmlc/xgboost/blob/905fdd3e08d91077aada776346c7e49e4ff69334/python-package/xgboost/dask.py#L170 called by https://github.com/dmlc/xgboost/blob/905fdd3e08d91077aada776346c7e49e4ff69334/python-package/xgboost/dask.py#L743 .
The test starts from here: https://github.com/dmlc/xgboost/blob/905fdd3e08d91077aada776346c7e49e4ff69334/tests/python-gpu/test_gpu_with_dask.py#L176 and the actual logic of this test is defined in here: https://github.com/dmlc/xgboost/blob/905fdd3e08d91077aada776346c7e49e4ff69334/tests/python/test_with_dask.py#L165 (for sharing between CPU and GPU)
Docker file: https://github.com/dmlc/xgboost/blob/master/tests/ci_build/Dockerfile.gpu
trivialfis
commented
3 years ago @hcho3 tried to reenable the test by upgrading the nightly https://github.com/dmlc/xgboost/pull/6825. A new error trace appears and cupy is also complaining of memory error. I'm suspecting something new in arrow that's not handled in downstream projects since both errors come from construction (map_blocks/map_partitions). Or is it somewhere in rmm?
```
[2021-04-06T04:11:56.613Z] self = Dask DataFrame Structure:
[2021-04-06T04:11:56.613Z] 0 1 2 3 4 5 6 7 ... ... ... ... ... ... ... ... ... ... ...
[2021-04-06T04:11:56.613Z] Dask Name: from-dask, 100 tasks
[2021-04-06T04:11:56.613Z] func =
```
[2021-04-06T04:11:56.618Z] def map_blocks(
[2021-04-06T04:11:56.618Z] func,
[2021-04-06T04:11:56.618Z] *args,
[2021-04-06T04:11:56.618Z] name=None,
[2021-04-06T04:11:56.618Z] token=None,
[2021-04-06T04:11:56.618Z] dtype=None,
[2021-04-06T04:11:56.618Z] chunks=None,
[2021-04-06T04:11:56.618Z] drop_axis=[],
[2021-04-06T04:11:56.618Z] new_axis=None,
[2021-04-06T04:11:56.618Z] meta=None,
[2021-04-06T04:11:56.618Z] **kwargs,
[2021-04-06T04:11:56.618Z] ):
[2021-04-06T04:11:56.618Z] """Map a function across all blocks of a dask array.
[2021-04-06T04:11:56.618Z]
[2021-04-06T04:11:56.618Z] Note that ``map_blocks`` will attempt to automatically determine the output
[2021-04-06T04:11:56.618Z] array type by calling ``func`` on 0-d versions of the inputs. Please refer to
[2021-04-06T04:11:56.618Z] the ``meta`` keyword argument below if you expect that the function will not
[2021-04-06T04:11:56.618Z] succeed when operating on 0-d arrays.
[2021-04-06T04:11:56.618Z]
[2021-04-06T04:11:56.618Z] Parameters
[2021-04-06T04:11:56.618Z] ----------
[2021-04-06T04:11:56.618Z] func : callable
[2021-04-06T04:11:56.618Z] Function to apply to every block in the array.
[2021-04-06T04:11:56.618Z] args : dask arrays or other objects
[2021-04-06T04:11:56.618Z] dtype : np.dtype, optional
[2021-04-06T04:11:56.618Z] The ``dtype`` of the output array. It is recommended to provide this.
[2021-04-06T04:11:56.618Z] If not provided, will be inferred by applying the function to a small
[2021-04-06T04:11:56.618Z] set of fake data.
[2021-04-06T04:11:56.618Z] chunks : tuple, optional
[2021-04-06T04:11:56.618Z] Chunk shape of resulting blocks if the function does not preserve
[2021-04-06T04:11:56.618Z] shape. If not provided, the resulting array is assumed to have the same
[2021-04-06T04:11:56.618Z] block structure as the first input array.
[2021-04-06T04:11:56.618Z] drop_axis : number or iterable, optional
[2021-04-06T04:11:56.618Z] Dimensions lost by the function.
[2021-04-06T04:11:56.618Z] new_axis : number or iterable, optional
[2021-04-06T04:11:56.618Z] New dimensions created by the function. Note that these are applied
[2021-04-06T04:11:56.618Z] after ``drop_axis`` (if present).
[2021-04-06T04:11:56.618Z] token : string, optional
[2021-04-06T04:11:56.618Z] The key prefix to use for the output array. If not provided, will be
[2021-04-06T04:11:56.618Z] determined from the function name.
[2021-04-06T04:11:56.618Z] name : string, optional
[2021-04-06T04:11:56.618Z] The key name to use for the output array. Note that this fully
[2021-04-06T04:11:56.618Z] specifies the output key name, and must be unique. If not provided,
[2021-04-06T04:11:56.618Z] will be determined by a hash of the arguments.
[2021-04-06T04:11:56.618Z] meta : array-like, optional
[2021-04-06T04:11:56.618Z] The ``meta`` of the output array, when specified is expected to be an
[2021-04-06T04:11:56.618Z] array of the same type and dtype of that returned when calling ``.compute()``
[2021-04-06T04:11:56.618Z] on the array returned by this function. When not provided, ``meta`` will be
[2021-04-06T04:11:56.618Z] inferred by applying the function to a small set of fake data, usually a
[2021-04-06T04:11:56.618Z] 0-d array. It's important to ensure that ``func`` can successfully complete
[2021-04-06T04:11:56.618Z] computation without raising exceptions when 0-d is passed to it, providing
[2021-04-06T04:11:56.618Z] ``meta`` will be required otherwise. If the output type is known beforehand
[2021-04-06T04:11:56.618Z] (e.g., ``np.ndarray``, ``cupy.ndarray``), an empty array of such type dtype
[2021-04-06T04:11:56.618Z] can be passed, for example: ``meta=np.array((), dtype=np.int32)``.
[2021-04-06T04:11:56.618Z] **kwargs :
[2021-04-06T04:11:56.618Z] Other keyword arguments to pass to function. Values must be constants
[2021-04-06T04:11:56.618Z] (not dask.arrays)
[2021-04-06T04:11:56.618Z]
[2021-04-06T04:11:56.618Z] See Also
[2021-04-06T04:11:56.618Z] --------
[2021-04-06T04:11:56.618Z] dask.array.blockwise : Generalized operation with control over block alignment.
[2021-04-06T04:11:56.618Z]
[2021-04-06T04:11:56.618Z] Examples
[2021-04-06T04:11:56.618Z] --------
[2021-04-06T04:11:56.618Z] >>> import dask.array as da
[2021-04-06T04:11:56.618Z] >>> x = da.arange(6, chunks=3)
[2021-04-06T04:11:56.618Z]
[2021-04-06T04:11:56.618Z] >>> x.map_blocks(lambda x: x * 2).compute()
[2021-04-06T04:11:56.618Z] array([ 0, 2, 4, 6, 8, 10])
[2021-04-06T04:11:56.618Z]
[2021-04-06T04:11:56.618Z] The ``da.map_blocks`` function can also accept multiple arrays.
[2021-04-06T04:11:56.618Z]
[2021-04-06T04:11:56.618Z] >>> d = da.arange(5, chunks=2)
[2021-04-06T04:11:56.618Z] >>> e = da.arange(5, chunks=2)
[2021-04-06T04:11:56.618Z]
[2021-04-06T04:11:56.618Z] >>> f = da.map_blocks(lambda a, b: a + b**2, d, e)
[2021-04-06T04:11:56.618Z] >>> f.compute()
[2021-04-06T04:11:56.618Z] array([ 0, 2, 6, 12, 20])
[2021-04-06T04:11:56.618Z]
[2021-04-06T04:11:56.618Z] If the function changes shape of the blocks then you must provide chunks
[2021-04-06T04:11:56.618Z] explicitly.
[2021-04-06T04:11:56.618Z]
[2021-04-06T04:11:56.618Z] >>> y = x.map_blocks(lambda x: x[::2], chunks=((2, 2),))
[2021-04-06T04:11:56.618Z]
[2021-04-06T04:11:56.618Z] You have a bit of freedom in specifying chunks. If all of the output chunk
[2021-04-06T04:11:56.618Z] sizes are the same, you can provide just that chunk size as a single tuple.
[2021-04-06T04:11:56.618Z]
[2021-04-06T04:11:56.618Z] >>> a = da.arange(18, chunks=(6,))
[2021-04-06T04:11:56.618Z] >>> b = a.map_blocks(lambda x: x[:3], chunks=(3,))
[2021-04-06T04:11:56.618Z]
[2021-04-06T04:11:56.618Z] If the function changes the dimension of the blocks you must specify the
[2021-04-06T04:11:56.618Z] created or destroyed dimensions.
[2021-04-06T04:11:56.618Z]
[2021-04-06T04:11:56.618Z] >>> b = a.map_blocks(lambda x: x[None, :, None], chunks=(1, 6, 1),
[2021-04-06T04:11:56.618Z] ... new_axis=[0, 2])
[2021-04-06T04:11:56.618Z]
[2021-04-06T04:11:56.618Z] If ``chunks`` is specified but ``new_axis`` is not, then it is inferred to
[2021-04-06T04:11:56.619Z] add the necessary number of axes on the left.
[2021-04-06T04:11:56.619Z]
[2021-04-06T04:11:56.619Z] Map_blocks aligns blocks by block positions without regard to shape. In the
[2021-04-06T04:11:56.619Z] following example we have two arrays with the same number of blocks but
[2021-04-06T04:11:56.619Z] with different shape and chunk sizes.
[2021-04-06T04:11:56.619Z]
[2021-04-06T04:11:56.619Z] >>> x = da.arange(1000, chunks=(100,))
[2021-04-06T04:11:56.619Z] >>> y = da.arange(100, chunks=(10,))
[2021-04-06T04:11:56.619Z]
[2021-04-06T04:11:56.619Z] The relevant attribute to match is numblocks.
[2021-04-06T04:11:56.619Z]
[2021-04-06T04:11:56.619Z] >>> x.numblocks
[2021-04-06T04:11:56.619Z] (10,)
[2021-04-06T04:11:56.619Z] >>> y.numblocks
[2021-04-06T04:11:56.619Z] (10,)
[2021-04-06T04:11:56.619Z]
[2021-04-06T04:11:56.619Z] If these match (up to broadcasting rules) then we can map arbitrary
[2021-04-06T04:11:56.619Z] functions across blocks
[2021-04-06T04:11:56.619Z]
[2021-04-06T04:11:56.619Z] >>> def func(a, b):
[2021-04-06T04:11:56.619Z] ... return np.array([a.max(), b.max()])
[2021-04-06T04:11:56.619Z]
[2021-04-06T04:11:56.619Z] >>> da.map_blocks(func, x, y, chunks=(2,), dtype='i8')
[2021-04-06T04:11:56.619Z] dask.array
kkraus14
commented
3 years ago What version of Arrow is being used? cuDF is pinned to 1.0.1 as of now so I doubt it's something in Arrow.
In the CuPy test are you setting the CuPy allocator to RMM?
How are you creating the host memory in this case?
trivialfis
commented
3 years ago I don't think we have done anything special in this case. It's still in the test data construction phase. The data is either from sklearn dataset or constructed using da.random.
trivialfis
commented
3 years ago cuDF is pinned to 1.0.1 as of now so I doubt it's something in Arrow.
Thanks for sharing.
Looking at the rmm configuration in tests, this should be the only parameter rmm_pool_size=2GB for LocalCUDACluster.
kkraus14
commented
3 years ago cuDF is pinned to 1.0.1 as of now so I doubt it's something in Arrow.
Thanks for sharing.
Looking at the rmm configuration in tests, this should be the only parameter
rmm_pool_size=2GBforLocalCUDACluster.
Do you always import cudf? It implicitly sets the CuPy allocator under the hood.
trivialfis
commented
3 years ago Do you always import cudf?
I assume it's yes. Since it's used in pytest.mark.skipif( no_cudf() ). The no_cudf function try to import cudf.
kkraus14
commented
3 years ago Any chance you could try not importing cudf to see if the cupy related failures go away?
trivialfis
commented
3 years ago Yup. I will try to sort it out tomorrow. Thanks for the suggestion.
trivialfis
commented
3 years ago I will need to try harder on reproducing it locally. Debugging on CI with this size of changes is not easy.
trivialfis
commented
3 years ago I think it's caused by xgboost setting the device to something else than 0 in another test.
trivialfis
commented
3 years ago I resolved the issue by ensuring xgboost doesn't change the device ordinal. But still, it might be better for cuDF to have some guards on GPU ID. Feel free to close this one if I should open a new issue on this topic. ;-) Thanks for all the replies!
This has happened in our CI a few times but I haven't been able to reproduce it deterministically yet. At the end is the error log copied from one of the failing logs from Jenkins. Opening an issue and see if there's someone else who can reproduce it more deterministically. Sorry for the noise.
The used cudf version is 0.18.