How to Optimize a CUDA Matmul Kernel for CuBLAS-Like Performance: A Worklog

Re:How to Optimize a CUDA Matmul Kernel for CuBLAS-Like Perf... by WithinReason on Thursday January 05, 2023 @08:12AM (#34259140)

He mentions having bank conflicts in shared memory, there is a simple trick to avoid that. You allocate an extra column for your shared buffers:

    __shared__ float As[(CHUNKSIZE+1) * CHUNKSIZE];

If CHUNKSIZE is a power of 2 this will eliminate bank conflicts.

[ Parent | on HN ]

Re:How to Optimize a CUDA Matmul Kernel for CuBLAS-Like Perf... by siboehm on Thursday January 05, 2023 @09:15AM (#34259800)

Author here: Seems like a good trick! Though won't this affect shared memory alignment and make me loose those LDS.128 instructions? Or do these not require alignment? There's so little good docs on SASS.

In general I'm still confused about whether vectorized load instructions (LDS.128) necessarily lead to bank conflicts or not. My impression was that consecutive 32b floats get mapped to different banks, so to avoid conflicts I'd want the warp to load 32*32b consecutive elements at each step.

[ Parent | on HN ]

Re:How to Optimize a CUDA Matmul Kernel for CuBLAS-Like Perf... by WithinReason on Thursday January 05, 2023 @09:24AM (#34259928)

Hmm, I think you might have to adjust the padding to be 128 bits then:

    __shared__ float As[(CHUNKSIZE+4) \* CHUNKSIZE]

Ultimately it's down to trial and error, like always with GPGPU.

[ Parent | on HN ]