decoder error use multiple GPUs and multiple threads, with cuda-10.2

[housebaby]

We derived a multi-thread multi-GPU version decoder, which can parallelly decode many utterances on multi-GPUs. for example, 40 threads evenly on 4 GPU .

In the main thread, we init our decoder by calling cudaSetDevice(gpu_id) for each gpu card, and loading a copy of am model on each gpu ( each copy will be used by many decoding threads)

////pthreadkey_gpuid will be used to index cu-device, which we modified in a way that
//// threads on each gpu will share one cu-device.

for (int i = 0; i < gpu_num_; i++) {
pthread_setspecific(::pthreadkey_gpuid, &i);
#ifdef HAVE_CUDA
CuDevice::Instantiate().SelectGpuId(gpu_option); // we will select pthreadkey_gpuid
CuDevice::Instantiate().AllowMultithreading();
#endif
am_nnet_[i] = new nnet3::AmNnetSimple;
.................
Input ki(fileName, &binary);
...............
am_nnet_[i]->Read(ki.Stream(), binary);
}

///in cu-device.h , we change the instantiate func , to support diffenrent cu-device on each GPU
static inline CuDevice& Instantiate() {
int index = *((int *)pthread_getspecific(::pthreadkey_gpuid));
//KALDI_LOG << "now device index is:" << index;
return global_device_[index];
}

//Then we create several decoding threads which are evenly allocated on multi-GPUs and use the //corresponding AM copy .

for (int i = 0; i < decoders_num_; i++) {
.................................
int tm = ((i) % gpu_num_)
pthread_setspecific(::pthreadkey_gpuid, &tm);
#ifdef HAVE_CUDA
CuDevice::Instantiate().SelectGpuId(config_->gpu_option_);
CuDevice::Instantiate().AllowMultithreading();
#endif
decoder_[i]->Init(config_, model_);
}

it works well with cuda-9.1.
But after we upgrade to cuda-10.2. IT can not work. Some problems arises,
if we just use on GPU card ( export CUDA_VISIBLE_DEVICES=0) , that is N threads on one GPU ,it works fine.
But if we make 2 or more GPUs visible, we encounter some confusing errors.
when CUDA_VISIBLE_DEVICES=0,1 or 0,1,2 , the progam will be locked somewhere
If CUDA_VISIBLE_DEVICES=0,1,2,3( that's all we have on a machine) , it will give "cublasStatus_t 1 : "CUBLAS_STATUS_NOT_INITIALIZED" returned from 'cublas_gemv(GetCublasHandle(), (trans==kTrans? CUBLAS_OP_N:CUBLAS_OP_T), M.NumCols(), M.NumRows(), alpha, M.Data(), M.Stride(), v.Data(), 1, beta, data_, 1)"

---

Can anyone give me a hint on how this happens?
What is the reasonable way to use M threads evenly allocated on N Gpu cards ( M = k * N ), while currently we use one cu-device for one gpu which many threads use the same cu-device.

[danpovey]

I don't believe Kaldi's code supports using multiple GPUs.

[housebaby]

@danpovey thanks for your quick reply.
Can you give more details on what limits kaldi to support multiple GPUs.
We make some modifacation on cu-device.cc and wrap a decoder with nnet3 am and lattice-faster-online-decoder in the above way. Actually it works fine with cuda9.1 on multiple GPUs. Problems accur with cuda-10.2

[danpovey]

It's about cu-device.cc. It only really supports a singleton object. It might be possible to have different ones per thread by changing the code, but I if there were >1 thread using the same GPU they might have different CUBLAS handles which might not be ideal. Possibly your logic for setting up the CUBLAS handle (and other similar handles) is wrong.

…

On Wed, May 27, 2020 at 6:46 PM housebaby ***@***.***> wrote: @danpovey <https://github.com/danpovey> thanks for your quick reply. Can you give more details on what limits kaldi to support multiple GPUs. We make some modifacation on cu-device.cc and wrap a decoder with nnet3 am and lattice-faster-online-decoder in the above way. Actually it works fine with cuda9.1 on multiple GPUs. Problems accur with cuda-10.2 — You are receiving this because you were mentioned. Reply to this email directly, view it on GitHub <#4084 (comment)>, or unsubscribe <https://github.com/notifications/unsubscribe-auth/AAZFLO2RHC7OOG5LRPOVWJ3RTTVQBANCNFSM4NL7WV7Q> .

[whaozl]

@danpovey thanks for your quick reply.
Can you give more details on what limits kaldi to support multiple GPUs.
We make some modifacation on cu-device.cc and wrap a decoder with nnet3 am and lattice-faster-online-decoder in the above way. Actually it works fine with cuda9.1 on multiple GPUs. Problems accur with cuda-10.2

I suggest you use a docker container and be independent.

[btiplitz]

I did share with nvidia changes to both Kaldi and a minor change to nvidia's code that allows for multiple GPU decoders. I think nvidia is working on doing changes for multiple GPU's and might or might not steal from my code

[housebaby]

@btiplitz thanks a lot. can you give more details on which part of nvidia code need to change? is the code of the lib like libcuda.* libcublas* ?

[btiplitz]

Now I’m reading the thread better, my mind was in a different place. Kaldi is a single gpu only so I don’t see how you had it work going back many versions of kaldi. Cuda is multi threaded but the cudevice code needs an instance of the cudevice stuff per gpu. My implementation added 4 of everything in cuda device but it could also have been done with a cudevice per gpu. I did not have a chance to try the alternate. Kaldi maintains a global for the active gpu. If the gpu’s were strapped then that is beyond my knowledge as I used each gpu independently

[btiplitz]

And you need a copy of Cublas cusolver and all the other objects per gpu.

[btiplitz]

@housebaby I am finally looking at the whole issue on the computer (vs the mobile device). The kaldi library itself assumes there is only a single GPU, thus changes to cu-device are required and if you use the kaldi decoder, a minor tweek is also required as that library actually spawns threads. @hugovbraun did nvidia make progress on the multi-gpu support.

[hugovbraun]

We're currently looking at sharing the allocator, but that's mostly for a single GPU configuration.
For multi-gpu, given the fact that no communication between devices is needed, why would we want to support multi-gpu directly into a single process (instead of one process/GPU)? Ease of use?

[btiplitz]

If a vast amount of data is being processed, it is easier to balance load within a single process. Looking at the entire reason why multi-threading is done today vs the old coding technique of multiple programs, the same logic can be applied to a GPU where there are more reasons to use multiple GPU's in a process, though there are some reasons why this is not always good.

[whaozl]

In fact, personally think that the multi-threaded piece should be independent. The model and feature module only need a single thread or in the main thread, just asTensorFlow.

[btiplitz]

tensor flow often has some library doing multi-threading for you. If the code is single threaded, only a single GPU is required as the GPU is exceptionally fast. The cudadecoder code is multi-threaded, so the programs making use of that code are actually multi-threaded but only using a single GPU. To need multiple GPU's requires a very high CPU count which most servers do not have. Multi-threaded calls to the API itself can be helpful for processing as a single object that takes a long time to process don't block everything else. The 2 solutions to this are to use async coding or multi-threading. Multi-threading is simpler but the number of threads is not so unlimited, where async code works with higher parallel tasks.

[hugovbraun]

I see, so for easier load balancing, and I guess for stuff such as sharing a cpu threadpool. We can look at bringing true multi-gpu support for Kaldi, but we would need @danpovey 's input first