Function launch_kernel_from_module

src/runtime/cuda/cuda_queue.cpp:181–211

};

result launch_kernel_from_module(CUmodule module,
                                 std::string_view kernel_name,
                                 const rt::range<3> &grid_size,
                                 const rt::range<3> &block_size,
                                 unsigned shared_memory, cudaStream_t stream,
                                 void **kernel_args) {
  CUfunction f;
  CUresult err = cuModuleGetFunction(&f, module, kernel_name.data());

  if (err != CUDA_SUCCESS) {
    return make_error(__acpp_here(),
                      error_info{"cuda_queue: could not extract kernel from module",
                                 error_code{"CU", static_cast<int>(err)}});
  }

  err = cuLaunchKernel(f, static_cast<unsigned>(grid_size.get(0)),
                       static_cast<unsigned>(grid_size.get(1)),
                       static_cast<unsigned>(grid_size.get(2)),
                       static_cast<unsigned>(block_size.get(0)),
                       static_cast<unsigned>(block_size.get(1)),
                       static_cast<unsigned>(block_size.get(2)),
                       shared_memory, stream, kernel_args, nullptr);

  if (err != CUDA_SUCCESS) {
    return make_error(__acpp_here(),
                      error_info{"cuda_queue: could not submit kernel from module",
                                 error_code{"CU", static_cast<int>(err)}});
  }
  
  return make_success();
}
}