我对CUDA编程比较陌生，所以有一些未解决的问题，我希望我能在正确的方向上得到一些提示 .
所以情况是我希望将2D数组与其转置相乘，并且准确地说我想要执行操作ATA .
我已经使用了cublas Dgemm 函数，现在我尝试使用平铺算法执行相同的操作，非常类似于CUDA指南中的算法 .
这种情况是，虽然初始算法运行正常，但我只想计算产品的上三角矩阵，希望我能够获得更好的操作时间，而且我不确定如何提取将有的块/块各个要素 .
所以如果你可以在这方面给我启发，或者给予任何暗示我会感激，因为我已经坚持了一段时间 .
这是内核的代码
__shared__ double Ads1[TILE_WIDTH][TILE_WIDTH];
__shared__ double Ads2[TILE_WIDTH][TILE_WIDTH];
//block row and column
//we save in registers for faster access
int by = blockIdx.y;
int bx = blockIdx.x;
int ty = threadIdx.y;
int tx = threadIdx.x;
int row = by * TILE_WIDTH + ty;
int col = bx * TILE_WIDTH + tx;
double Rvalue = 0;
if(row >= width || col >= width) return;
//Each thread block computes one sub-matrix Rsub of result R
for (int i=0; i<(int) ceil(((double) height/TILE_WIDTH)); ++i)
{
Ads1[tx][ty] = Ad[(i * TILE_WIDTH + ty)*width + col];
Ads2[tx][ty] = Ad[(i * TILE_WIDTH + tx)*width + row];
__syncthreads();
for (int j = 0; j < TILE_WIDTH; ++j)
{
if ((i*TILE_WIDTH + j) > height ) break; //in order not to exceed the matrix's height
Rvalue+=Ads1[j][tx]*Ads2[ty][j];
}
__syncthreads();
}
Rd [row * width + col] = Rvalue;
1 回答
您可能希望使用here描述的批量dgemm API函数，将输出矩阵与块对角线和角点进行递归划分 . 您还希望在计算中 balancer 最小块大小与开销，以避免小调用 . 最后，请注意矩阵乘法在某个阶段转换内存绑定，这可能在现代GPU上有点大 .