非矩形矩陣CUDA的通用快速轉置

Question

SDK提供了一個解決方矩陣轉置的例子和策略，但是在非方陣上執行轉置還有一個好方法嗎？我目前比較幼稚的實現如下這可能是可怕的：

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]; 
    } 
} 

Answer 1

這裏我的想法是，只有必要的線程/塊（每個線程掉期轉置矩陣的方形部分的方形子的兩個條目矩陣），然後遍歷並轉置剩餘的條目。

__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個矩陣，讓我知道你在想什麼。