Fix race condition in CUDA stream creation

[rafbiels]

Do not increment NumComputeStreams / NumTransferStreams before cuStreamCreateWithPriority returns. Too early increment caused other threads to read the incremented count before a CUDA stream was created and try to use an invalid stream handle, causing crashes.

The construction:

if (condition) {
  lock_this_scope
  if (condition) {
    create_object
    update_condition
  }
}
use_object

is only thread-safe if update_condition happens after create_object is completed. This PR ensures the ordering.

intel/llvm PR: intel/llvm#15100

[rafbiels]

Reproducer for the crashes:

#include <future>
#include <sycl/sycl.hpp>

int main() {
  sycl::queue q{};

  constexpr static unsigned int numThreads{128};
  std::vector<std::future<void>> futures;
  std::array<unsigned int, numThreads> hostData{};
  std::array<unsigned int*, numThreads> devicePointers{};
  futures.reserve(numThreads);

  for (unsigned int i{0}; i<numThreads; ++i)  {
    devicePointers[i] = sycl::malloc_device<unsigned int>(1, q);
  }

  for (unsigned int i{0}; i<numThreads; ++i)  {
    futures.push_back(std::async([&q, &devicePointers, &hostData, i](){
      q.copy(hostData.data()+i, devicePointers[i], 1);
    }));
  }

  for (unsigned int i{0}; i<numThreads; ++i)  {
    futures[i].wait();
  }
  futures.clear();

  q.wait_and_throw();

  for (unsigned int i{0}; i<numThreads; ++i)  {
    sycl::free(devicePointers[i], q);
  }

}

This crashes for me in ~50% of runs before this PR with:

UR CUDA ERROR:
        Value:           400
        Name:            CUDA_ERROR_INVALID_HANDLE
        Description:     invalid resource handle
        Function:        urEnqueueUSMMemcpy

[hdelan]

Nice catch! Tricky one

[npmiller]

Should we add this to v0.10.0?