last day (16 days later) » 

8:37 PM
0
Q: Matrix-vector multiplication in CUDA: benchmarking & performance

Pantelis SopasakisI'm updating my question with some new benchmarking results (I also reformulated the question to be more specific and I updated the code)... I implemented a kernel for matrix-vector multiplication in CUDA C following the CUDA C Programming Guide using shared memory. Let me first present some ben...

 
Remember that your Jetson has only one Streaming Multiprocessor. In a such "elementary" approach, above a certain small value, further blocks can be executed only sequentially. Perhaps therefore your linear increase.
 
@JackOLantern OK, this is true, but still I'm running cuBLAS on the same platform. I'd like to know whether I can somehow get rid of divergent threads (i.e., to get rid of these if blocks that account for inexact tiling...) Is there some standard practice?
 
Could you please post a full version of your code so that perhaps someone could play with it?
 
@JackOLantern I updated my question with new benchmarking results and I also posted the new kernel and uploaded all the source code on github (see link in the question). I'm playing with the profiler trying to understand what can be improved in the code.
 
In your timing toc() function, you have a cudaDeviceSynchronize() call that is causing, on my GT540M, unstable timing. Also, I have the feeling that you are compiling in Debug mode since in your Matlab plotting code you are searching for fetch_this in a Debug directory. If you are requesting performance, you should compile in Release mode.
 
8:37 PM
@RobertCrovella OK, I see. I just didn't want to overload the question.
@JackOLantern Thanks for the hints! I'm re-running some benchmarks and I will update the question if I see some difference.
 
Should the beta parameter of cublasSgemv() be equal to 0, instead of 1? Furthermore, have you checked that cublasSgemv() and matvec return the same results? Have you accounted that cuBLAS requires column-major ordering of the matrix? The result of the cublasSgemv() function in your code returns all NaNs to me.
 
@JackOLantern You're absolutely right! Actually I should have put a cudaMemset(dev_y_cublas, 0, n_rows_max * sizeof(real_t)); (I updated my question now). I compare with compare_results(dev_y_cublas,dev_y, nrows);
 
Your code works correctly now for square matrices, namely, when CONSTANT_ROWS = CONSTANT_COLS, but not yet when CONSTANT_ROWS != CONSTANT_COLS. Please, correctness before optimization. Also, there is no need to using cudaMemset for the cuBLAS case. Finally, as pointed out by @NicolasBertJohnson, you seem that you are assuming column-major ordering in your hand-written kernel, but this will lead to highly uncoalesced global memory reads. A more consistent comparison would be to consider row-major ordering and CUBLAS_OP_N for cuBLAS.
I think I fixed the above problem. if (idx_x + row < nx) { x_shared[row] = xsub[row]; } requires the else branch else x_shared[row] = 0.f;.
 
@JackOLantern Are you using the latest version of the code (in the question)? I've tested it for row and column sizes 1 to 512 and it works fine for all their combinations. I can post the code I used to test it if you like.
I moved the discussion here following the suggestion of Stackoverflow
I created a git repository where I've put the latest version of the code: github.com/alphaville/mv
 
Yes, I'm using your latest version. Often uninitialized shared memory values are 0s. If this happens, your code works, but unintentionally. The else statement will cover the case of inexact tiling.
 
8:48 PM
Yes, I understand, but then I do for (unsigned int e = 0; e < n_star; ++e) {..., i.e., uninitialized values of x_shared are not used
I'm using a different type of if-else block for inexact tiling. That's what I'm using n_star for
This part of the code:
if ((nRows % block_size == 0 && nx % block_size == 0)
|| (m < num_hor_blocks -1 && bid < gridDim.x - 1) ) {
covers cases of total exact tiling (both dimensions of the matrix are multiples of 16), or the tile totally covers the matrix.
 
9:04 PM
I see what do you say. However, without the else statement I receive wrong results. I will check it further. Have you though about the non-coalesced memory reads?
 
Yes, I did a quick benchmark and I'm getting better (but wrong) results. So, I first need to see what is wrong and come up with a correct implementation first. But, I think it's an excellent idea!
 

  last day (16 days later) »