在CUDA中使用小M的两个MXN矩阵执行矢量点产品的最快方法是什么?

What is the fastest way to perform vector-by-vector dot products for two MxN matrices with small M in CUDA?

本文关键字：执行 是什么 方法 MXN CUDA 两个      更新时间：2023-10-16

我有两个矩阵，每个矩阵是 MxN，其中 M = 16和 N较大(例如，n = 262144(。我的目标是产生一个长度N的向量，其中每个元素都对应于每个矩阵中nth矢量的点产物。

我尝试了以下方法，其中cIdx对应于每个矩阵中的列向量的列索引。毫不奇怪，Nvidia Visual Profiler告诉我，这种方法主要是记忆带宽。

    public static void MatrixDotProduct(
        float* matrix1,
        float* matrix2,
        float* dotProduct,
        int2 matrixDimensions)
    {
        int i = blockIdx.x * blockDim.x + threadIdx.x;
        int stride = gridDim.x * blockDim.x;
        float sum;
        for (int cIdx = i; cIdx < matrixDimensions.y; cIdx += stride)
        {
            int ci = cIdx * matrixDimensions.x;
            sum = 0f;
            for (int j = 0; j < matrixDimensions.x; j++)
            {
                sum += matrix1[ci + j] * matrix2[ci + j];
            }
            dotProduct[cIdx] = sum;
        }
    }

我还找到了一个vector-vector-vector Dot产品的版本，我也尝试并行化。不幸的是，这比上述实现慢20％(也许是因为我的M = 16？(。有没有更好的方法来解决我在这里错过的这个问题？

这对于缺乏数据重用和小型向量大小(较小的扭曲(而进行优化并不容易优化。代码应是内存绑定的，缓存性能将是至关重要的。

要尝试的一件事是减少音量并通过为负载使用向量类型来提高内存交易的效率。我期望的是：

__device__ float vdot(float2 v1, float2 v2) {
    return (v1.x * v2.x) + (v1.y * v2.y);
}
__device__ float vdot(float4 v1, float4 v2) {
    return (v1.x * v2.x) + (v1.y * v2.y) + (v1.z * v2.z) + (v1.w * v2.w);
}
template<typename VT, int NT>
__device__ float vector_dotprod(const VT* v1, const VT* v2) {
    float sum = 0.f;
#pragma unroll
    for (int j = 0; j < NT; j++) {
        sum += vdot(v1[j], v2[j]);
    }
    return sum;
}
template<typename VT, int Nrows>
__global__
void MatrixDotProductPlus(float* matrix1, float* matrix2, float* dotProduct, int2 matrixDimensions)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = gridDim.x * blockDim.x;
    int stride2 = stride * Nrows;
    VT* m1 = reinterpret_cast<VT*>(matrix1) + i * Nrows;
    VT* m2 = reinterpret_cast<VT*>(matrix2) + i * Nrows;
    for (; i < matrixDimensions.y; i += stride, m1 += stride2, m2 += stride2) {
        dotProduct[i] = vector_dotprod<VT,Nrows>(m1, m2);
    }
}

[警告：仅经过非常轻微的测试 - 使用自身风险]

对于float2案例的代码的速度约为两倍，对于Maxwell或Pascal架构的float4案例，对于长度为16它们是2或4的圆形倍数，尽管我怀疑循环的展开与向量类型本身一样至关重要。

相关文章：