我有两个矩阵,每个都是 MxN
,其中 M = 16
和 N
大得多(比如 n = 262144
,例如)。我的目标是生成一个长度为 N
的 vector ,其中每个元素对应于每个矩阵中的 nth
vector 的点积。
我尝试了以下方法,其中 cIdx
对应于每个矩阵中列 vector 的列索引。毫不奇怪,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-by-vector 的版本点积,我也尝试将其并行化。不幸的是,这比上面的实现慢了 20%(也许是因为我的 M = 16
?)。有没有更好的方法来解决我在这里缺少的这个问题?
最佳答案
这不是一个容易优化的案例,因为缺乏数据重用和小 vector 大小(小于扭曲)。代码应受内存限制,缓存性能至关重要。
可以尝试的一件事是通过对负载使用 vector 类型来减少容量并提高内存事务的效率。我希望是这样的:
__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 架构上的长度 vector 的 float4
情况下大约快四倍16. 这假设您在编译时知道 vector 的长度并且它们是 2 或 4 的整数倍,尽管我怀疑循环展开与 vector 类型本身一样重要。
关于c++ - 在 CUDA 中用小 M 对两个 MxN 矩阵执行逐 vector 点积的最快方法是什么?,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/44253005/