cuda子矩阵答案

【问题标题】：cuda sub matrixcuda子矩阵
【发布时间】：2012-04-12 22:23:53
【问题描述】：

问题：

我有 4 个单精度数字矩阵 (64x64)。需要进行如下计算：

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

想法：

为了加快计算使用共享内存。因为每个线程块（在我的 GPU 的情况下）有 16KB 共享内存并且浮点数为 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 个子矩阵并对其进行计算。

【问题讨论】：

cos(D) 是指D 的矩阵余弦，还是D 的所有元素的余弦？这两件事非常不同。
sry，所有元素 D 的余弦

标签： cuda gpgpu

【解决方案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 的峰值吞吐量，如果数组的大小超过您拥有的最佳一维网格的大小，每个线程将处理多个输入/输出值推出。当然，如果您总共只处理 4096 个元素，这是相当学术性的——这可能太小了大约 2 个数量级，无法从使用 GPU 中获得任何好处。

【讨论】：

我原来的内核是：int idx = blockIdx.x * blockDim.x + threadIdx.x;int idx = blockIdx.x * blockDim.x + threadIdx.x;if ( idx>= 64 || idy >= 64 ) return;R[idx*64+idy] = A[idx*64+idy]*sin(B[idx*64+idy]) + C[idx*64+idy]*sin(D[idx*64+idy]);所以shared mem不能加速吗？
请不要在 cmets 中发布代码，将其编辑到您的原始问题中。但答案是否定的。在这种情况下使用共享内存不会带来任何性能优势。

【解决方案2】：

这里有个问题，你的操作/传输比是 1。由于线程和全局内存之间的带宽瓶颈，你可能很难从 GPU 上获得任何不错的速度，并且没有办法减少它。

当从全局内存中重复调用某些数据时，共享内存解决方案通常是最好的。您无需从低带宽、高延迟的全局内存中重复加载此数据，而是从那里加载一次，然后从更高带宽、更低延迟的共享内存中进行后续加载。请注意，这是 higher 和 lower，而不是 high 和 low。使用共享内存仍然会降低性能。

你的情况，因为元素不会从全局内存中多次调用，将它们存储在共享内存中只会增加共享内存使用带来的带宽限制和延迟。因此，实际上，此解决方案只会增加从共享内存访问数据到加载数据的延迟。

现在，如果您要执行多个计算，并且其中也使用了其中一些矩阵，那么将它们组合到一个内核中可能会提高速度，因为您可能能够为整个事情加载一次而不是每次操作一次。如果情况并非如此，并且您无法提高运算/传输率，那么您将很难获得不错的速度，最好在 CPU 上进行这些计算。

您甚至可以从 CPU 上的多线程中获得一些不错的结果。

【讨论】：