9ca33cf4d5
Sources (all CUDA 10.2 compatible, no CUTLASS/Triton dependency): - leimao/CUDA-GEMM-Optimization: v00-v07, fp16 WMMA variant, double buffered - siboehm/SGEMM_CUDA: kernel 1-12, warp tiling + double buffering - wangzyon/NVIDIA_SGEMM_PRACTICE: kernel 1-7 - edtallison/sgemm-cuda: kernel 1-12 (reimplementation with notes) Key porting issue: ALL kernels hardcode WARPSIZE=32. BI-V100 has warp_size=64. Need to: 1. Replace all 32U / WARPSIZE constants with 64 2. Adjust warp subtile decomposition (WMITER, WNITER, WSUBM, WSUBN) 3. Adjust shared memory bank conflict avoidance (may have different bank count) 4. Test __shfl_down_sync with mask=0xFFFFFFFFFFFFFFFF (64-bit)
80 lines
2.8 KiB
Plaintext
80 lines
2.8 KiB
Plaintext
#pragma once
|
|
|
|
#include <algorithm>
|
|
#include <cassert>
|
|
#include <cstdio>
|
|
#include <cstdlib>
|
|
#include <cublas_v2.h>
|
|
#include <cuda_runtime.h>
|
|
|
|
#define CEIL_DIV(M, N) (((M) + (N)-1) / (N))
|
|
|
|
template <const int BM, const int BN, const int BK, const int TM>
|
|
__global__ void sgemm1DBlocktiling(int M, int N, int K, float alpha,
|
|
const float *A, const float *B, float beta,
|
|
float *C) {
|
|
// If we flip x and y here we get ~30% less performance for large matrices.
|
|
// The current, 30% faster configuration ensures that blocks with sequential
|
|
// blockIDs access columns of B sequentially, while sharing the same row of A.
|
|
// The slower configuration would share columns of A, but access into B would
|
|
// be non-sequential. So the faster configuration has better spatial locality
|
|
// and hence a greater L2 hit rate.
|
|
const uint cRow = blockIdx.y;
|
|
const uint cCol = blockIdx.x;
|
|
|
|
// each warp will calculate 32*TM elements, with 32 being the columnar dim.
|
|
const int threadCol = threadIdx.x % BN;
|
|
const int threadRow = threadIdx.x / BN;
|
|
|
|
// allocate space for the current blocktile in SMEM
|
|
__shared__ float As[BM * BK];
|
|
__shared__ float Bs[BK * BN];
|
|
|
|
// Move blocktile to beginning of A's row and B's column
|
|
A += cRow * BM * K;
|
|
B += cCol * BN;
|
|
C += cRow * BM * N + cCol * BN;
|
|
|
|
// todo: adjust this to each thread to load multiple entries and
|
|
// better exploit the cache sizes
|
|
assert(BM * BK == blockDim.x);
|
|
assert(BN * BK == blockDim.x);
|
|
const uint innerColA = threadIdx.x % BK; // warp-level GMEM coalescing
|
|
const uint innerRowA = threadIdx.x / BK;
|
|
const uint innerColB = threadIdx.x % BN; // warp-level GMEM coalescing
|
|
const uint innerRowB = threadIdx.x / BN;
|
|
|
|
// allocate thread-local cache for results in registerfile
|
|
float threadResults[TM] = {0.0};
|
|
|
|
// outer loop over block tiles
|
|
for (uint bkIdx = 0; bkIdx < K; bkIdx += BK) {
|
|
// populate the SMEM caches
|
|
As[innerRowA * BK + innerColA] = A[innerRowA * K + innerColA];
|
|
Bs[innerRowB * BN + innerColB] = B[innerRowB * N + innerColB];
|
|
__syncthreads();
|
|
|
|
// advance blocktile
|
|
A += BK;
|
|
B += BK * N;
|
|
|
|
// calculate per-thread results
|
|
for (uint dotIdx = 0; dotIdx < BK; ++dotIdx) {
|
|
// we make the dotproduct loop the outside loop, which facilitates
|
|
// reuse of the Bs entry, which we can cache in a tmp var.
|
|
float tmpB = Bs[dotIdx * BN + threadCol];
|
|
for (uint resIdx = 0; resIdx < TM; ++resIdx) {
|
|
threadResults[resIdx] +=
|
|
As[(threadRow * TM + resIdx) * BK + dotIdx] * tmpB;
|
|
}
|
|
}
|
|
__syncthreads();
|
|
}
|
|
|
|
// write out the results
|
|
for (uint resIdx = 0; resIdx < TM; ++resIdx) {
|
|
C[(threadRow * TM + resIdx) * N + threadCol] =
|
|
alpha * threadResults[resIdx] +
|
|
beta * C[(threadRow * TM + resIdx) * N + threadCol];
|
|
}
|
|
} |