c++ - 矩阵矩形部分转置Cuda

Question

我正在编写 Cuda 程序来转置方阵，其想法是根据矩阵的大小分两部分进行；使用 Tile 将矩阵大小切成均匀大小，并保留矩形部分，我将其单独转置 例如:67 x 67 矩阵和 Tile :32，第一部分是 64x64 转置，然后第二部分是 3x67。

我的问题是在矩形部分，下面的第一个代码显示了具有定义值的主要代码:

const int TILE_DIM = 32;
const int BLOCK_ROWS = 8;
const int NUM_REPS = 100;

const int Nx = 2024; //size of the matrix
const int Ny = 2024;

int main(int argc, char **argv)
{
const int nx = Nx;
const int ny = Ny; // Size of the Arrays
const int mem_size = nx*ny*sizeof(int);// Size of the Orig.Arr

int *h_idata = (int*)malloc(mem_size); // original Host Arr.

int *d_idata; //device Arr.
checkCuda(cudaMalloc(&d_idata, mem_size));

dim3 dimGridX(nx / TILE_DIM, 1, 1); //grid dimension used
dim3 dimBlockX(TILE_DIM, 1, 1); // number of threads used

// the Kernel Function for only the rectangle
EdgeTransposeX << < dimGrid, dimBlock >> >(d_idata);
cudaEventRecord(startEvent, 0);
cudaEventRecord(stopEvent, 0);
cudaEventSynchronize(stopEvent);
cudaEventElapsedTime(&ms, startEvent, stopEvent);
cudaMemcpy(h_idata, d_idata, mem_size, cudaMemcpyDeviceToHost);

我被建议不要使用共享的内核代码，所以下面是我的做法:

__global__ void EdgeTransposeX(int *idata)
{

    int tile_C[Edge][Nx];
    int tile_V[Nx][Edge];

    int x = blockIdx.x * TILE_DIM + threadIdx.x;

    if (x == (nEven - 1))
    {

        for (int j = 0; j < Nx; j++)
            for (int i = 1; i <= Edge; i++)
            {

            tile_V[j][i - 1] = idata[j*Nx + (x + i)];
             tile_C[i - 1][j] = idata[(x + i)*Nx + j];}

         __syncthreads();

        for (int j = 0; j < Nx; j++)
          for (int i = 1; i <= Edge; i++)
         {
            idata[j*Nx + (x + i)] = tile_C[i - 1][j];
            idata[(x + i)*Nx + j] = tile_V[j][i - 1];}

       } }

代码在矩阵大小达到 1025 之前正常工作，之后它停止工作，知道为什么吗？我在这里遗漏了什么吗？

Answer 1

您的二维数组 tile_C 和 tile_V 在物理上存储在 GPU 的本地内存中。每个线程的本地内存量为 512KB。确认您没有为每个线程使用超过 512KB 的本地内存。

在设备代码中声明的自动变量，没有任何设备，本节中描述的shared 和constant 限定符通常位于寄存器中。但是在某些情况下，编译器可能会选择将它放在本地内存中。此片段摘自“CUDA C PROGRAMMING GUIDE 2015”第 89 页。

我的建议是您使用可视化分析器来检查占用、寄存器和本地内存使用情况。

此链接可能对您有帮助:link .

我在 2D 中使用 cuda 曲面实现了方矩阵的转置，它适用于从 2 到 16384 的大小，增量为 2。如果您不介意实现无平铺版本，我推荐这种方法。