cuda子矩阵

Question

问题:

我有 4 个单精度数字矩阵 (64x64)。需要像这样进行计算:

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

想法:

为了加速计算使用共享内存。由于每个线程 block (在我的 GPU 的情况下)有 16KB 共享内存，并且 float 的大小为 4，因此可以在共享内存中存储 4000 个 float 。因此对于每个矩阵使用 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]);

如何将原始矩阵分解为子矩阵，使每个 block 都有自己的 4 个子矩阵并对其进行计算。

Answer 1

除非我误解了你的问题，否则你不需要也不应该为此操作使用共享内存。共享内存对于在同一 block 内的线程之间共享和恢复数据以及促进联合内存访问非常有用。您的操作似乎不需要这些东西即可正常工作。按照您建议的方式使用共享内存可能比直接从全局内存中读取慢。此外，因为您只担心元素操作，内核的索引方案可以大大简化——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
    }
}

在此代码中，您将启动尽可能多的 block 以达到 GPU 的峰值吞吐量，并且如果数组的大小超过您拥有的最佳一维网格的大小，则每个线程将处理多个输入/输出值推出。当然，如果您总共只处理 4096 个元素，这是非常学术的——这可能太小了 2 个数量级，无法从使用 GPU 中获得任何好处。