CUDA中使用共享内存进行非方阵转置

Question

我试图获取适用于所有大小的CUDA 矩阵转置示例 的变体。简而言之，我必须将输入数组（double *a ）写入较大矩阵（double *tab）的两个不同部分（您将注意到不同的偏移量）。我按行优先格式存储数据，因此我使用此宏进行索引：

#define IDX2L(i,j,ld) (((i)*ld))+(j)) // 0 based index +row-major format

这是我使用的简单代码。

__global__ void cuda_a_Coalesced(double *tab, int tab_rows, int a_rows, double *a)
{
    __shared__  double tile[16*(16+1)]; 
    int col = threadIdx.x + blockIdx.x * blockDim.x;
    int row = threadIdx.y + blockIdx.y * blockDim.y;

    int col_2, row_2;
    int a_cols=tab_rows-a_rows; // tab_rows-a_rows is the number of columns of a
    int tab_cols=2*tab_rows+2;  // 2*tab_rows+2 is the number of columns of tab

    if( (col<a_cols) && (row<a_rows) ) 
    {
        // Load the data into shared mem
        tile[threadIdx.x+threadIdx.y*(16+1)]=a[IDX2L(row,col,a_cols)];

        // Normal copy (+ offsets)
        tab[IDX2L(row,col+tab_rows+a_rows,tab_cols)]= tile[threadIdx.x+threadIdx.y*(16+1)];

        // New idx
        col_2 = blockIdx.y * blockDim.y + threadIdx.x;
        row_2 = blockIdx.x * blockDim.x + threadIdx.y;
    }
    __syncthreads();

    if( (row_2<a_cols) && (col_2<a_rows) )
        // Transpose (+ other offsets)
        tab[IDX2L(row_2+a_rows,col_2+tab_rows,tab_cols)]= -tile[threadIdx.y+threadIdx.x*(16+1)];

}

启动参数如下：

b1=(int)ceil((float)a_cols/16);
b2=(int)ceil((float)a_rows/16);
dim bck(b1,b2):dim th(16,16);

cuda_a_Coalesced<<<bck,th>>>(tab,tab_rows,a_rows,a);

无论大小，正常复制始终表现良好。转置复制仅适用于块大小的整数倍大小（如CUDA示例中所示）。当转置复制失败时，某些操作的部分正确，而其他部分不正确，我无法准确预测或跟踪。请注意，想法是更改共享内存中的索引，以便可以将转置以行主格式写入输出矩阵中（由于行主格式）。

有人能告诉我代码只能使用那种大小的原因吗？

有没有什么技巧可以解决这种情况？

- engineer H

---

你确定代码正确吗？乍一看，对于某些线程，col_2和row_2可能未定义。 - user1545642

在这个版本之前，我尝试将涉及col_2和row_2的操作放在第一个if语句内，显然是在__syncthreads()之后，但我遇到了同样的问题。那就是你所指的问题吗？ - engineer H

对于那些不满足条件“(col<a_cols) && (row<a_rows)” 的线程，col_2和row_2将未初始化。也许你应该初始化col_2 = MAX_INT，row_2 = MAX_INT，这样下一个条件“(row_2<a_cols) && (col_2<a_rows)”只对正确的线程成立。 - user1545642

似乎问题与您提到的有关。我尝试了您的选项，但问题并没有消失。然而，继续按照您的假设，我尝试了另一种初始化方式：int col_2=blockIdx.y * blockDim.y + threadIdx.x; int row_2=blockIdx.x * blockDim.x + threadIdx.y;通过这样做，并删除if语句中出现的同名计算，我最终让代码正常工作，因此问题得以解决！！ - engineer H

@engineerH：请将您解决问题的方法作为答案添加并接受它。这将标记问题已得到解答和完成，并且可能有助于下一个遇到相同问题的人。 - talonmies

Answer 1

问题是由于一些未定义的线程导致的，因为在一个 if() 语句中分配了 col_2 和 row_2 的值，而不是所有线程都访问。

为了解决这种情况，我们可以在声明这些变量时给出 col_2 和 row_2 的值，并删除在上述 if() 中发生的同名计算。

__shared__  double tile[16*(16+1)];

int col = threadIdx.x + blockIdx.x * blockDim.x;
int row = threadIdx.y + blockIdx.y * blockDim.y;

int col_2 = blockIdx.y * blockDim.y + threadIdx.x;
int row_2 = blockIdx.x * blockDim.x + threadIdx.y;

int a_cols=tab_rows-a_rows; 
int tab_cols=2*tab_rows+2;

因此，代码的其余部分如下：

if( (col<a_cols) && (row<a_rows) ) 
{
    // Load the data into shared mem
    tile[threadIdx.x+threadIdx.y*(16+1)]=a[IDX2L(row,col,a_cols)];
    // Normal copy (+ offsets)
    tab[IDX2L(row,col+tab_rows+a_rows,tab_cols)]= tile[threadIdx.x+threadIdx.y*(16+1)];
}
__syncthreads();

if( (row_2<a_cols) && (col_2<a_rows) )
    // Transpose (+ other offsets)
    tab[IDX2L(row_2+a_rows,col_2+tab_rows,tab_cols)]= -tile[threadIdx.y+threadIdx.x*(16+1)];

- engineer H