I’ve been playing around with cuBLAS’s strided batch matrix-matrix multiplication and I was suprised to find that for small matrix sizes, enabling the transposition flag drastically reduces performance for small matrix sizes.

I’ve been playing around with cuBLAS’s strided batch matrix-matrix multiplication and I was suprised to find that for small matrix sizes, enabling the transposition flag drastically reduces performance for small matrix sizes.

I’ve been playing around with cuBLAS’s strided batch matrix-matrix multiplication and I was suprised to find that for small matrix sizes, enabling the transposition flag drastically reduces performance for small matrix sizes.

I’ve been playing around with cuBLAS’s strided batch matrix-matrix multiplication and I was suprised to find that for small matrix sizes, enabling the transposition flag drastically reduces performance for small matrix sizes.

The full code is in the book CUDA by Example. Mainly the kernel function reads:

Here is the code, if you compile it with nvcc, it gives me as error:

I’m working with CUDA and have implemented a 2D array in shared memory, but I’m encountering bank conflicts that I’m struggling to understand. Here’s my setup:

I want to know how to make this cuda kernel more precise or atleast consistent.

Error checking (and adding a print) after each CUDA call makes the code difficult to read. For example, here is the code to initialize 4 variables on the GPU with error handling:

I am trying to optimize a CUDA program. In this program, thread i needs to wait for thread i-1 to store data in shared memory before it can proceed. Is there a better synchronization method than __syncthreads()?