Cuml: [BUG] make_regression pytest segfault (CI)

[dantegd]

Describe the bug
Regular CI in Python 3.6, in Ubuntu 16.04 and 18.04, CUDA 10.0 ran into a segfault from running the following pytest of make_regression:

cuml.test.test_make_regression.test_make_blobs_scalar_parameters[3.5--4.0-None-None-True-True-1-7-10-100000-single] 

Output:

RuntimeError: Exception occured! file=/rapids/cuml/cpp/src/common/cumlHandle.cpp line=273: FAIL: call='cudaStreamCreate(&stream)'. Reason:an illegal memory access was encountered  Obtained 64 stack frames #0 in /opt/conda/envs/rapids/lib/libcuml++.so(_ZN8MLCommon9Exception16collectCallStackEv+0x3e) [0x7fe2edeac19e] #1 in /opt/conda/envs/rapids/lib/libcuml++.so(_ZN8MLCommon9ExceptionC2ERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE+0x80) [0x7fe2edeaccb0] #2 in /opt/conda/envs/rapids/lib/libcuml++.so(_ZN2ML15cumlHandle_impl15createResourcesEv+0x517) [0x7fe2edeab9a7] #3 in /opt/conda/envs/rapids/lib/libcuml++.so(_ZN2ML15cumlHandle_implC1Ei+0xf0) [0x7fe2edeabc40] #4 in 

Afterwards, all subsequent pytests fail. Given that this is essentially in cudaStreamCreate, this could be an error from a prior test or section of code being caught here.

Steps/Code to reproduce bug
Haven't reproduced it locally, but here is a CI log showing the error, and the full stack trace:

https://gpuci.gpuopenanalytics.com/job/docker/job/tests/job/docker-test-cuml/97/CUDA_VERSION=10.0,LINUX_VERSION=ubuntu16.04,PYTHON_VERSION=3.6/testReport/junit/cuml.test/test_make_regression/test_make_blobs_scalar_parameters_3_5__4_0_None_None_True_True_1_7_10_100000_single_/

Coming from this job: https://gpuci.gpuopenanalytics.com/job/docker/job/tests/job/docker-test-cuml/97/

NOTE: Test says make_blobs incorrectly (failed to catch that small typo in the PR), but it is the make_regressions file that calls cuml.make_regression

Expected behavior
Pytest should never seg fault in any supported OS/CUDA/Python

Environment details (please complete the following information):

• Environment location: CI/Container
• Linux Distro/Architecture: Ubuntu 16.04 and 18.04
• GPU Model/Driver: V100 32 GB,
• CUDA: 10.0
• Method of cuDF & cuML install: cuDF nightly, cuML from source, see attached links above for more details.

[teju85]

Ouch... can we try running this with CUDA_LAUNCH_BLOCKING env-var set to 1?
Also tagging @Nyrio if he knows of any known issues with make_regression (and probably to also update our unit-tests to plug this hole).

[danielhanchen]

I took a quick glance at the code. Since it's an Illegal Mem Access, it might be because of (purely speculation)

[Line 112 https://github.com/rapidsai/cuml/blob/branch-0.11/cpp/src_prims/random/make_regression.h]

template <typename DataT, typename IdxT>
static __global__ void _gather2d_kernel(DataT* out, const DataT* in,
                                        const IdxT* perms, IdxT n_rows,
                                        IdxT n_cols) {
  IdxT tid = blockIdx.x * blockDim.x + threadIdx.x;
  ****** const DataT* row_in = in + n_cols * perms[tid];  ******
  ****** DataT* row_out = out + n_cols * tid; ******

  if (tid < n_rows) {
    for (IdxT i = 0; i < n_cols; i++) {
      row_out[i] = row_in[i];
    }
  }
}

Should be:

template <typename DataT, typename IdxT>
static __global__ void _gather2d_kernel(DataT* out, const DataT* in,
                                        const IdxT* perms, IdxT n_rows,
                                        IdxT n_cols) {
  IdxT tid = blockIdx.x * blockDim.x + threadIdx.x;

  if (tid < n_rows) {
    ****** const DataT* row_in = in + n_cols * perms[tid];  ******
    ****** DataT* row_out = out + n_cols * tid; ******
    for (IdxT i = 0; i < n_cols; i++) {
      row_out[i] = row_in[i];
    }
  }
}

[Nyrio]

@danielhanchen is right, thanks for the quick reaction and sorry for any inconvenience caused.

@dantegd Hopefully fixed in https://github.com/rapidsai/cuml/pull/1326 (and the typo as well)

@teju85 I don't think that the unit tests need to be updated. They were already encountering cases where this was an issue, they just didn't segfault before...

[danielhanchen]

:)

[teju85]

I guess if we had an ability to run all unit-tests with cuda-memcheck maybe that'd have caught this earlier.

Probably we could enable #433 and/or #451, atleast as part of weekly regressions @dantegd ? @cjnolet ? @JohnZed ?

[dantegd]

@teju85 this was caught as part of the normal CI runner jobs, so a weekly CUDA memcheck wouldn’t have caught it earlier. Regardless, hopefully it can be a weekly job soon. Probably impossible as part of PR CI since our runtimes are already pretty high, but a nightly/weekly would be ideal. Haven’t ran CUDA memcheck in the full c++ tests lately, we will need to check if any issues rise with the latest commit.

@Nyrio thanks for tackling this so quickly!

[dantegd]

Will close issue now that the PR was merged, will reopen if needed.

[JohnZed]

I think CUDA memcheck is always pushed back on the agenda because there are so many false positives , especially in third-party code like faiss. Maybe with extensive filtering it could be useful.