2

我在 CUDA 中用于矩阵乘法的代码让我可以将方阵和非方阵相乘，但是，宽度和高度都必须是块大小的倍数。

因此，例如，我可以乘以 [3][6] * [6][3]（使用 blocksize=3），但不能乘以 [3][2]*[2][3]。

有谁知道这样做的方法？这是我的内核：

#include <stdio.h>

#include <limits.h>

#include <stdlib.h>
#define blocksize 3
#define HM (1*blocksize) 
#define WM (2*blocksize) 
#define WN (1*blocksize)
#define HN WM 
#define WP WN   
#define HP HM  
#define PTH WM
#define PTW HM

__global__ void nonsquare(float*M, float*N, float*P, int uWM,int uWN)

{
__shared__ float MS[blocksize][blocksize];
__shared__ float NS[blocksize][blocksize];


int tx=threadIdx.x, ty=threadIdx.y, bx=blockIdx.x, by=blockIdx.y;
int rowM=ty+by*blocksize;
int colN=tx+bx*blocksize;
float Pvalue=0;


for(int m=0; m< uWM/blocksize;++m){
    MS[ty][tx]=M[rowM*uWM+(m*blocksize+tx)];
    NS[ty][tx]=M[colN + uWN*(m*blocksize+ty)];
    __syncthreads();

    for(int k=0;k<blocksize;k++)
        Pvalue+=MS[ty][k]*NS[k][tx];
    __syncthreads();
    P[rowM*WP+colN]=Pvalue;
     }
    }

提前致谢！

4

2 回答 2

5

我认为最简单的方法就是用零填充最后的块：

for(int m=0; m< uWM/blocksize;++m){
    colM = m*blocksize+tx;
    rowN = m*blocksize+ty;
    if (rowM > uWN || rowN > uWM || colM > uWM || colN > uWN) {
        MS[ty][tx]=0.;
        NS[ty][tx]=0.;
    } else {
        MS[ty][tx]=M[rowM*uWM+colM];
        NS[ty][tx]=N[colN + uWN*rowN];
    }

加号或减号。（那条 NS 行应该引用 N，而不是 M，对吗？）

但是，既然我似乎是这里唯一一个提倡尽可能使用现有调整库的人——为什么不使用CUBLAS或MAGMA而不是自己滚动呢？它们速度很快，并且经过数百名用户的测试。

于 2011-03-28T15:28:07.467 回答
2

这里的基本性能要求是共享内存“切片”的第一维或第二维是 16 的整数倍 - 从历史上看，这是实现最佳全局内存带宽（即半扭曲合并事务）所必需的。它应该是图块的第一个维度还是第二个维度取决于矩阵是按列主要顺序还是行主要顺序存储的。没有什么说共享内存块需要是正方形的，只是存储的前导维度（BLAS 表示法中的 LDA）是 16 的整数倍。

您可以轻松地使用瓦片尺寸作为模板参数对内核进行模板化，并根据矩阵尺寸实例化多个版本。对于给定的架构，可能存在平衡占用率和指令级并行度的最佳切片尺寸。解决这个问题的“聪明”方法可能是将矩阵乘法分解为两个操作 - 第一个以最佳平铺大小完成大部分工作，第二个以不同大小处理剩余的列。如果结果是在产品完成后直接返回到主机内存，则最好使用优化的 BLAS 在主机上完成第二个操作，并与 GPU 内核重叠。这是 UTK Magma 库中的许多例程使用的方法。

于 2011-03-29T10:29:21.673 回答