非矩形矩阵CUDA的通用快速转置

SDK提供了一个解决方矩阵转置的例子和策略，但是在非方阵上执行转置还有一个好方法吗？我目前比较幼稚的实现如下这可能是可怕的：非矩形矩阵CUDA的通用快速转置

template<class S> 
__global__ void transpose(S *Source, S *Destination, int SizeX, int SizeY) { 
    int tid = threadIdx.x + blockIdx.x * blockDim.x; 

    if (tid<SizeX*SizeY) { 
     int X = tid % SizeX; 
     int Y = tid/SizeX; 

     //(x,y) => (y,x) 

     int newId = (SizeY*X) + Y; 
     Destination[newId] = Source[tid]; 
    } 
}

来源

2012-07-17 illumi

为什么“可能很糟糕” - 你有没有基准测试内核？它达到峰值存储带宽的多少部分？ – talonmies 2012-07-17 11:35:34

可能很糟糕，因为我以前没有测试过它。 Profiler报告27GB/s写入吞吐量和5.3GB/s读取。 – illumi 2012-07-17 12:44:37

可能最好在块中读入共享内存，进行转置并写入块。你在这里的写作完全没有合并。 – 2012-07-17 19:57:26

这里我的想法是，只有必要的线程/块（每个线程掉期转置矩阵的方形部分的方形子的两个条目矩阵），然后遍历并转置剩余的条目。

__global__ void kernelTranspuesta(float *a, float *c, int m, int n) { 
    int i = threadIdx.x + blockIdx.x*blockDim.x; 
    int j = threadIdx.y + blockIdx.y*blockDim.y; 
    int smallest = M < N ? M : N; 

    while(j < smallest){ 
     i = threadIdx.x + blockIdx.x*blockDim.x; 
     while(i < j){ 
     c[i*m+j] = a[j*n+i]; 
     c[j*m+i] = a[i*n+j]; 
     i+= blockDim.x*gridDim.x; 
     } 
     if(i == j) 
     c[j*m+i] = a[i*n+j]; 

    j+= blockDim.y*gridDim.y; 
    } 

    if(M > N) { 
     i = threadIdx.x + blockIdx.x*blockDim.x + N; 
     j = threadIdx.y + blockIdx.y*blockDim.y; 

     while(i < M){ 
     j = threadIdx.y + blockIdx.y*blockDim.y; 
     while(j < N){ 
      c[j*m+i] = a[i*n+j]; 
      j+= blockDim.y*gridDim.y; 
     } 
     i+= blockDim.x*gridDim.x; 
     } 
    }else{ 
     i = threadIdx.x + blockIdx.x*blockDim.x; 
     j = threadIdx.y + blockIdx.y*blockDim.y + M; 

    while(i < M){ 
     j = threadIdx.y + blockIdx.y*blockDim.y + M; 
     while(j < N){ 
      c[j*m+i] = a[i*n+j]; 
      j+= blockDim.y*gridDim.y; 
     } 
     i+= blockDim.x*gridDim.x; 
     } 
    } 

}

内核调用

dim3 hilos(16,16); // hilos(blockDim.x, blockDim.y) 
    dim3 bloques(8,8); // bloques(gridDim.x, gridDim.y) 

    kernelTranspuesta<<<bloques, hilos>>>(aD, cD, m, n);

我测试了512x256和256×512个矩阵，让我知道你在想什么。

来源

2013-04-09 03:50:14

看起来不错，我通常在3000x2000左右的矩阵上。在共享内存空间中进行转置不会更快吗？ – illumi 2013-04-23 09:56:42

非矩形矩阵CUDA的通用快速转置

回答

相关问题