矩阵矩形部分转置Cuda

Matrix the Rectangle Part transpose Cuda

本文关键字：Cuda 转置 形部      更新时间：2023-10-16

im编写Cuda程序来转换平方矩阵，其想法是根据矩阵的大小分为两部分；矩阵大小用Tile切割成均匀大小，并保留矩形部分左侧i将其单独转置Ex: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，之后它就停止工作了，知道为什么吗？我是不是遗漏了什么？

二维数组tile_C和tile_V实际存储在GPU的本地内存中。每个线程的本地内存量为512KB。请验证每个线程使用的本地内存是否不超过512KB。

设备代码中声明的一个自动变量，没有任何设备，本节中描述的共享和常量限定符通常驻留在寄存器中。然而，在某些情况下，编译器可能会选择将其放置在本地内存中。该片段摘自《CUDA C编程指南2015》第89页。

我的建议是使用可视化探查器来检查占用率、注册和本地内存使用情况。

此链接可能对您有所帮助：链接。

我在2D中使用cuda曲面实现了方形矩阵的Transpose，它适用于从2到16384的大小，增量为2的幂。如果你不介意实现无平铺版本，我建议你采用这种方法。

相关文章：