MAGMA and CUDA streams

[vish108]

Hi,

I am trying to write code in which multiple CPU threads will perform GPU based operations in parallel. The matrix operations are listed below:
1. magma_sgemm/magma_dgemm
2. magmablas_sgeadd
3. cublasSaxpy

Based on what I have read, I can use the function magmablas_setkernelstream(a_cuda_stream) to set a kernel stream in which the operations will be performed. Can you please tell if all the above functions support cuda streams? I have also read that modern GPUs can allow upto 16 concurrent kernel launches. The code I am trying to write has a vector of vctors. Each vector will be of size N, and the number of vectorss will be 2N. N can be upto 10,000.

The way I am planning to do this (pseudocode):

Code: Select all

#pragma omp parallel for schedule(static,1) num_threads(16)
for(int i =0; i < 2N - 1; i++)
{
    int thread_id = get_omp_thread_num();
    create cuda stream using thread_id
    set magam stream using the cuda stream created above
    call magma functions to operate on vector[i] (listed above)
}

Please let me know if this approach will yield good results. Is it very naive to just have the max. possible CUDA streams? Should I take into account the vector size to determine the number of streams to use?

Sincerely,

Vishal

[mgates3]

Yes, those functions all support using a stream. The magma_sgemm/magma_dgemm are simply wrappers around cublasSgemm/cublasDgemm. Similarly, there is a wrapper magma_saxpy around cublasSaxpy. (This helps porting to OpenCL and other architectures.)

Unfortunately, magmablas_setkernelstream is not currently thread local. If you set a stream in one thread, all other threads see the same stream. You will need to use mutex locks or some equivalent mechanism, e.g.,

lock
magma_setkernelstream
magma_sgemm
unlock

We are working to fix this thread safety issue, but it changes the MAGMA API significantly.
-mark

[vish108]

Thank you. I will try using mutexes. Meanwhile, can you please point me to some code samples which use magma along with openmp/pthreads.

[vish108]

Hi,

I am a little unsure about using locks. Won't that slow things down a lot? Would it be faster to not use openmp at all, and just process the loop serially?

Thank you,

Vishal

[mgates3]

There is some overhead in applying locks. However, remember that the gemm calls are async, so it won't wait for each gemm to finish before unlocking -- just until the gemm is queued.
-mark

[vish108]

Excellent! Thank you. Mark ,from your post "However, remember that the gemm calls are async, so it won't wait for each gemm to finish before unlocking -- just until the gemm is queued." So are there magma functions which are not asynchronous? Or, were you referring to my particular use case.

[mgates3]

BLAS calls are asynchronous, in general. More specifically, all the magma_xxxx BLAS functions are async -- those are just wrappers around CUBLAS. There are a couple exceptions for functions beginning with "magmablas_" -- trsm, hemv, symv without workspace; multi-GPU hemm, symm.

Most of the other functions, providing LAPACK functionality, are synchronous. For instance, getrf, getrf_gpu, potrf, potrf_gpu, etc. These routines are hybrid, doing computation on both CPU and GPU, so they cannot be async.

-mark