cuda子矩阵

Question

问题:

我有4个矩阵(64x64)的单精度数.需要做如下计算:

R = A * sin(B) + C * cos(D)

理念:

加快计算使用共享内存.由于每个线程块都有(在我的GPU的情况下)16KB共享内存并且float的大小为4,因此可以在共享内存中存储4000个浮点数.因此,对于每个矩阵,使用1000个元素,每个维度为31个元素.

所以每个矩阵应分成16个子矩阵(16x16).

dim3 dimBlock(16, 16, 1)
dim3 dimGrid(4, 4, 1)

核心:

int Tx = threadIdx.x;
int Ty = threadIdx.y;

int Bx = blockIdx.x;
int By = blockIdx.y;

int idx = Bx * blockDim.x + Tx;
int idy = By * blockDim.y + Ty;

__shared__ float s_A[16*16];
__shared__ float s_B[16*16];
__shared__ float s_C[16*16];
__shared__ float s_D[16*16];

// I am not sure how to write this part

s_A[(Tx * blockDim.x + Ty + By) + Bx] = A[idx * 64 + idy];
s_B[(Tx * blockDim.x + Ty + By) + Bx] = B[idx * 64 + idy];
s_C[(Tx * blockDim.x + Ty + By) + Bx] = C[idx * 64 + idy];
s_D[(Tx * blockDim.x + Ty + By) + Bx] = D[idx * 64 + idy];

R[idx * 64 + idy] = s_A[(Tx * blockDim.x + Ty + By) + Bx] * sin(s_B[(Tx * blockDim.x + Ty + By) + Bx]) + s_C[(Tx * blockDim.x + Ty + By) + Bx] * cos(s_D[(Tx * blockDim.x + Ty + By) + Bx]);

如何将原始矩阵分配给子矩阵,使每个块具有自己的4个子矩阵并对其进行计算.

Answer 1

除非我误解了您的问题,否则您不需要也不应该使用共享内存进行此操作.共享内存对于在同一块内的线程之间共享和重新生成数据以及促进合并的内存访问非常有用.您的操作似乎不需要这些东西正常工作.以你提出的方式使用共享内存可能比直接从全局内存中读取更慢.此外,因为你只担心元素的运算,内核的索引方案可以大大简化-的事实是A,B,C和D是"矩阵"是风马牛不相及的计算,因为我明白你的问题.

因此,内核的近乎最佳版本可以简单地编写

__global__ void kernel(const float *A, const float *B, const float *C, 
                        const float *D, const int n, float *R)
{
    int tidx = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;

    while(tidx < n) {
        R[tidx] = A[idx] * sinf(B[idx]) + C[idx]*cosf(D[idx]);
        tidx += stride
    }
}

在此代码中,您将启动尽可能多的块以达到GPU的峰值吞吐量,并且如果阵列的大小超过您启动的最佳1D网格的大小,则每个线程将处理多个输入/输出值.当然,如果你只处理4096个元素,那么这是非常学术的 - 可能大约2个数量级太小而无法从使用GPU中获得任何好处.