[SYCL] [CUDA] Invalid memory access in multithreaded application when multiple contexts used

[s-kanaev]

Describe the bug
Invalid memory access when using multiple contexts.

To Reproduce

Build commands:

clang++ -isystem $DPCPP_HOME/build/include/sycl -isystem $CUDADIR/include -fsycl -fsycl-targets=nvptx64-nvidia-cuda-sycldevice -fsycl-unnamed-lambda reduced_2q_repro.cpp -o reduced_2q_repro.bin -Wl,-rpath,$CUDADIR/lib64: $DPCPP_HOME/build/lib/libsycl.so $CUDADIR/lib64/libcurand.so /usr/lib/x86_64-linux-gnu/libcuda.so $CUDADIR/lib64/libcudart_static.a -lpthread -ldl /usr/lib/x86_64-linux-gnu/librt.so

clang++ -DUSE_SAME_CTX -isystem $DPCPP_HOME/build/include/sycl -isystem $CUDADIR/include -fsycl -fsycl-targets=nvptx64-nvidia-cuda-sycldevice -fsycl-unnamed-lambda reduced_2q_repro.cpp -o reduced_2q_repro_same_ctx.bin -Wl,-rpath,$CUDADIR/lib64: $DPCPP_HOME/build/lib/libsycl.so $CUDADIR/lib64/libcurand.so /usr/lib/x86_64-linux-gnu/libcuda.so $CUDADIR/lib64/libcudart_static.a -lpthread -ldl /usr/lib/x86_64-linux-gnu/librt.so

The first command builds an application which will produce invalid memory access. The second one will produce an application which won't produce invalid memory access. The only difference between these two applications is that the former will use different contexts for two queues. The latter application will use the same context for both queues.

Source code:

#include <CL/sycl.hpp>
#include <CL/sycl/backend/cuda.hpp>
#include <cuda.h>

#ifdef USE_SAME_CTX
# define CONSTRUCT_CTX(name, dev) sycl::context name{dev}
# define CONSTRUCT_Q(name, ctx, dev, ...) sycl::queue name{ctx, dev, __VA_ARGS__}
#else
# define CONSTRUCT_CTX(name, dev)
# define CONSTRUCT_Q(name, ctx, dev, ...) sycl::queue name{dev, __VA_ARGS__}
#endif

void generate(sycl::queue &Q, sycl::buffer<float> &R) {
  Q.submit([&](sycl::handler &H) {
    auto Acc = R.get_access<sycl::access::mode::read_write>(H);

    H.host_task([=](sycl::interop_handle IH) {
      auto Ptr = reinterpret_cast<float *>(IH.get_native_mem<sycl::backend::cuda>(Acc));
      (void)Ptr;
    });
  }).wait_and_throw();
}

void test(sycl::device &D) {
  std::int64_t n = 1000;
  std::vector<float> RHost(n);
  std::vector<float> R1Host(n);

  CONSTRUCT_CTX(C, D);

  try {
    CONSTRUCT_Q(Q, C, D, [](sycl::exception_list EL) {
          for (auto EP : EL)
            std::rethrow_exception(EP);
        });
    sycl::buffer<float> R{RHost.data(), RHost.size()};

    generate(Q, R);

    Q.wait_and_throw();
  } catch (const std::exception &E) {
    std::cout << "Exception: " << E.what() << std::endl;
  }

  try {
    CONSTRUCT_Q(Q, C, D, [](sycl::exception_list EL) {
          for (auto EP : EL)
            std::rethrow_exception(EP);
        });
    sycl::buffer<float> R1{R1Host.data(), R1Host.size()};

    generate(Q, R1);

    Q.wait_and_throw();
  } catch (const std::exception &E) {
    std::cout << "Exception: " << E.what() << std::endl;
  }
}

int main() {
  sycl::device D{sycl::gpu_selector{}};

  std::cout << "Running on " << D.get_info<sycl::info::device::name>() << std::endl;

  for (size_t Idx = 0; Idx < 1; ++Idx) {
    std::cout << "Iteration " << Idx << " starts" << std::endl;
    test(D);
    std::cout << "Iteration " << Idx << " ends" << std::endl;
  }
}

reduced_2q_repro.bin will not fail. Though valgrind will report invalid memory accesses.

Payload of the application (i.e. actions before, after and inside host-task) can be modified in such a way that the application will eventually hang due to the same invalid memory access.

If the for loop in main has at least two iterations, the application is killed due to SEGFAULT which is due to the same memory access violation.

Culprit of this invalid memory access is that cuCtxSetCurrent frees some memory when host-task attempts to wait for own dependencies.

Valgrind's log (with SYCL_PI_TRACE=-1) is attached for reduced_2q_repro.bin with only a single iteration of for loop in main.

Environment (please complete the following information):

• OS: Linux
• Target device: CUDA
• DPC++ version:

clang version 13.0.0 (git@github.com:intel/llvm.git 63ba1ceab4945c325485cb99b205558de6025f9a)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: ......

• CUDA version: 10.1

Additional context

reduced_2q_repro.bin.memcheck.log

[sergey-semenov]

It appears that enqueue-arg-order-buffer/image.cpp tests from llvm-test-suite are failing on CUDA due to the same issue of context management race condition. I've noticed that, for reasons that are unclear to me, replacing cuCtxSetCurrent with cuCtxPushCurrent/cuCtxPopCurrent in ScopedContext from the CUDA plugin fixes the issue (both for the tests and this reproducer), but I'm not sure if that's the correct fix for this since according to my understanding of CUDA driver spec cuCtxSetCurrent should work just fine there.

[npmiller]

It appears that enqueue-arg-order-buffer/image.cpp tests from llvm-test-suite are failing on CUDA due to the same issue of context management race condition. I've noticed that, for reasons that are unclear to me, replacing cuCtxSetCurrent with cuCtxPushCurrent/cuCtxPopCurrent in ScopedContext from the CUDA plugin fixes the issue (both for the tests and this reproducer), but I'm not sure if that's the correct fix for this since according to my understanding of CUDA driver spec cuCtxSetCurrent should work just fine there.

Okay so I had a look into this and found another work around, in the ScopedContext, removing the needToRecover boolean and always recovering the context also fixes the invalid memory access.

From what I can tell what's going on is that the ScopedContext is used in a different thread than piContextCreate and piContextDestroy, so for the first queue, the ScopedContext sets the context as current but then leaves it as current, there's a comment explaining that this is to avoid context switches:

      if (original_ == nullptr) {
        // No context is installed on the current thread
        // This is the most common case. We can activate the context in the
        // thread and leave it there until all the PI context referring to the
        // same underlying CUDA context are destroyed. This emulates
        // the behaviour of the CUDA runtime api, and avoids costly context
        // switches. No action is required on this side of the if.
      } else {
        needToRecover_ = true;
      }

But the problem is then that when the context is destroyed in piContextDestroy, it is destroyed and popped in the main thread but is still current the second thread where it was set by the ScopedContext, this is spelled out in the cuCtxDestroy spec:

If ctx is current to other threads, then ctx will remain current to those threads, and attempting to access ctx from those threads will result in the error CUDA_ERROR_CONTEXT_IS_DESTROYED.

And you can see it by debugging/printing the context addresses in ScopedContext.

First queue:

cuCtxCreate: 0x13b3210
ScopedContext (desired, original): 0x13b3210, (nil)
~ScopedContext
cuCtxDestroy: 0x13b3210

Second queue:

cuCtxCreate: 0x19a6580
ScopedContext (desired, original): 0x19a6580, 0x13b3210
~ScopedContext > cuCtxSetCurrent: 0x13b3210
cuCtxDestroy: 0x19a6580

In the second queue the ScopedContext has a current context which is the context from the previous queue which has since been destroyed.

And that makes sense as all the memcheck failures are in ~ScopedContext > cuCtxSetCurrent when trying to set the destroyed context as current. I'm not sure why we're not seeing CUDA_ERROR_CONTEXT_IS_DESTROYED here but maybe we're just "lucky" that the memory there still looks enough like a context.

This also explains why replacing cuCtxSetCurrent with cuCtxPushCurrent/cuCtxPopCurrent in the ScopedContext also fixes the issue because then the first context isn't left hanging.

Either solution seems okay however I'm a bit unsure about the performance implications since it seems it was originally done this way for performance.

[npmiller]

Here's a quick graph of the issue at play here:

[npmiller]

It is a little awkward to solve because the PI plugins have no knowledge of the thread pool or anything like that at the moment.

[s-kanaev]

It is a little awkward to solve because the PI plugins have no knowledge of the thread pool or anything like that at the moment.

Still, the API has to be somewhat thread-safe. As far as I understand PI API was "copied" from OpenCL and OpenCL guarantees are implied.

@bader , @romanovvlad , @smaslov-intel, WDYT?

[romanovvlad]

Still, the API has to be somewhat thread-safe. As far as I understand PI API was "copied" from OpenCL and OpenCL guarantees are implied.

+1

[npmiller]

Well it is somewhat thread-safe, this isn't really a threading problem.

It stems from an optimization in the CUDA PI plugin trying to limit the number of context switches by leaving a context set on the thread. On the main thread it works fine because the main thread releases the contexts properly but then on the other internal threads the deleted contexts may still be referenced.

But yeah it's definitely a bug and needs to be fixed, it's just tricky to fix without either losing performance or introducing potential bugs in other scenarios.

I believe there's three main ways of solving this:

• Don't leave contexts hanging: this is probably the most "correct" but will incur a performance hit that can be pretty significant
• Leave contexts hanging but never try to restore contexts that were set: this would work because the deleted contexts on the other threads would just get overridden silently when the second queue comes around, however it could cause issues in some cases with CUDA interop if the user makes assumptions on what contexts are set.
• Find a way to cleanup the CUDA contexts from the thread pool threads at context deletion

[s-kanaev]

this isn't really a threading problem

Could you, please, elaborate on this? Probably, I've missed something.

Off-topic:

however it could cause issues in some cases with CUDA interop if the user makes assumptions on what contexts are set.

As long as the user is using SYCL, he shouldn't be assuming on internals. I mean, SYCL is sort of black-box, isn't it?

[npmiller]

Could you, please, elaborate on this? Probably, I've missed something.

Yes, what I mean is that it's not a concurrency issue, there's no race condition or anything like that, everything happens in the correct order, the issue is that we're not cleaning up the CUDA contexts properly in the other internal threads.

The problem is that CUDA contexts are thread specific, and that cuCtxDestroy only cleans up the thread it is called from, see:

If ctx is current to other threads, then ctx will remain current to those threads, and attempting to access ctx from those threads will result in the error CUDA_ERROR_CONTEXT_IS_DESTROYED.

• https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__CTX.html#group__CUDA__CTX_1g27a365aebb0eb548166309f58a1e8b8e

That's why you can see on the graph I posted that the context A remains active even after q1 is deleted (which deletes the context).

The tasks that were run in the thread pool left the context set as current for the thread pool thread.

As long as the user is using SYCL, he shouldn't be assuming on internals. I mean, SYCL is sort of black-box, isn't it?

Well yes, but I'm talking of a use case where someone would use both CUDA and SYCL code. At the moment if you set the CUDA context with the CUDA API, then use SYCL, SYCL will restore your original CUDA context every time so you can just carry out CUDA work as normal after SYCL execution. With the solution I was describing in that bullet point, SYCL wouldn't restore the original context and the user would have to set his CUDA context manually again.