Disallow programmatic stream serialization in JIT kernel launches (#1932)

This option allows the driver to overlap kernels on the same stream. But this is not really used by the cuVS kernel model and we expect no overlap between kernels on the same stream. This may cause some downstream issues in mixed stream usage.

Authors:
  - Divye Gala (https://github.com/divyegala)

Approvers:
  - Jinsol Park (https://github.com/jinsolp)

URL: #1932

Author: divyegala

cpp/src/detail/jit_lto/AlgorithmLauncher.cpp

@@ -37,17 +37,12 @@ AlgorithmLauncher& AlgorithmLauncher::operator=(AlgorithmLauncher&& other) noexc
 void AlgorithmLauncher::call(
   cudaStream_t stream, dim3 grid, dim3 block, std::size_t shared_mem, void** kernel_args)
 {
-  cudaLaunchAttribute attribute[1];
-  attribute[0].id = cudaLaunchAttributeProgrammaticStreamSerialization;
-  attribute[0].val.programmaticStreamSerializationAllowed = 1;
-
   cudaLaunchConfig_t config;
   config.gridDim          = grid;
   config.blockDim         = block;
   config.stream           = stream;
-  config.attrs            = attribute;
-  config.numAttrs         = 1;
   config.dynamicSmemBytes = shared_mem;
+  config.numAttrs         = 0;
 
   RAFT_CUDA_TRY(cudaLaunchKernelExC(&config, kernel, kernel_args));
 }