用于NxN矩阵的CUDA矩阵转置到位实现，该实现仅在上部traingular条目上迭代

Question

我正在研究矩阵转置就地实现，其想法如下：让索引只取与上三角条目相对应的值(因为对角条目在转换期间保持不变(，并通过临时整数在一个线程内交换值A[i，j]和A[j，i]。

我很难将正确的下一个索引对A[I_next，j_next]分配给当前处理

A[I，j]示例：用于N=5; total-threads=3;

这是一个5x5矩阵，其中的条目包含它们将被访问的威胁的thread_id。属于第一个线程的条目也是斜体的(只是为了显示我希望为该特定线程采用row_idx、col_idx的索引。

1

Col0	Col1	Col2	Col3	Col4
/	1
			2	3	1
					1

我应该提到，我希望这能适用于N=10^8，所以不幸的是，创建一个有足够线程的网格不是一个选项。关于性能，你当然是对的，这个想法不是很快。。然而，在转到更优化的版本之前，我仍然想让它发挥作用

• 我怀疑规模意味着你不能创建足够的线程，但我们可以把它放在一边
• 我同意这些评论，这些评论建议你可能需要仔细考虑你的算法起点，以高效使用GPU，但你似乎已经承认了这一点

有了这篇序言，我相信你想要的是如何将线性索引转换为二维三角形索引。

这个问题可以在不使用迭代、递归、浮点运算、平方根或其他高级函数的情况下解决，只需使用加法、减法、乘法和除法。

一旦我们将问题简化为可以使用线性索引来解决的问题，那么我们就可以在线性索引上应用网格步长循环等方法，并将其转换为2D三角形"；在飞行中"；以允许任意数量的线程处理并完成问题。

基于这些想法，这里有一种可能的方法：

$ cat t2137.cu
#include <iostream>
// conversion of linear index to triangular matrix (row,col) index
__host__ __device__ void linear_to_triangular(const unsigned i, const unsigned n, unsigned &trow, unsigned &tcol)
{
int c = i;
int row = c/(n-1);
int col = c - (row*(n-1)); // effectively modulo
if (col < row) {
col = (n-2)-col;
row = (n-1)-row;
}
tcol = col+1;
trow = row;
}
const size_t mdim = 547;
using mt = int;
const size_t tdim = 136;
const unsigned bdim = 64;
template <typename T>
__global__ void t(T *m, const size_t n){
for (int idx = blockIdx.x*blockDim.x+threadIdx.x; idx < ((n*(n-1)) >> 1); idx += gridDim.x*blockDim.x){ // 1D grid-stride loop
unsigned i,j;
linear_to_triangular(idx, n, i, j);
T e1 = m[i*n+j];
T e2 = m[j*n+i];
m[i*n+j] = e2;
m[j*n+i] = e1;
}
}
int main(){
mt *d, *h;
h = new mt[mdim*mdim];
cudaMalloc(&d, mdim*mdim*sizeof(mt));
for (int i = 0; i < mdim*mdim; i++) h[i] = i;
cudaMemcpy(d, h, mdim*mdim*sizeof(mt), cudaMemcpyHostToDevice);
t<<<(tdim+bdim-1)/bdim, bdim>>>(d, mdim);
cudaMemcpy(h, d, mdim*mdim*sizeof(mt), cudaMemcpyDeviceToHost);
for (int i = 0; i < mdim*mdim; i++) {
int row = i/mdim;
int col = i - row*mdim;
if (h[col*mdim+row] != i) {std::cout << "Mismatch at: " << i << " was: " << h[col*mdim+row] << " should be: " << i << std::endl; return 0;}
}
}
$ nvcc -o t2137 t2137.cu
$ compute-sanitizer ./t2137
========= COMPUTE-SANITIZER
========= ERROR SUMMARY: 0 errors
$

我认为深入表演没有多大意义。一个天真的转置，比如你所展示的，可以有一个"转置"；合并的"；一半，但会有另一半(糟糕的("；未康复"；。就这一点而言，这种算法没有什么不同。线性到三角形的转换仍然倾向于创建索引，使得相邻线程通常会加载(或存储，取决于您的写入方式(相邻元素，但当然，三角形数据补丁的交错模式会在相邻模式中引入不连续性。如果你的负载部分合并，你的商店就不会了。