关于“流”的问题

加速计算 CUDA
1

我的程序的步骤是这样的：
1.创建4个流
2.四个流异步传输数据
3.内核函数执行矩阵乘法（每一个流执行两个矩阵乘法）
4.内核函数执行矩阵加法（将上一步骤每个流中的两个矩阵相加，最终每个流得到一个矩阵加结果，四个流总共得到四个矩阵）
5.四个流异步传输到主机端

cudaMemcpyAsync(d_a1,a1,size,cudaMemcpyHostToDevice,stream[0]);
cudaMemcpyAsync(d_b1,b1,size,cudaMemcpyHostToDevice,stream[0]);
cudaMemcpyAsync(d_a2,a2,size,cudaMemcpyHostToDevice,stream[1]);
cudaMemcpyAsync(d_b2,b2,size,cudaMemcpyHostToDevice,stream[1]);
cudaMemcpyAsync(d_a3,a3,size,cudaMemcpyHostToDevice,stream[2]);
cudaMemcpyAsync(d_b3,b3,size,cudaMemcpyHostToDevice,stream[2]);
cudaMemcpyAsync(d_a4,a4,size,cudaMemcpyHostToDevice,stream[3]);
cudaMemcpyAsync(d_b4,b4,size,cudaMemcpyHostToDevice,stream[3]);
dim3 threads=dim3(BLOCK_SIZE,BLOCK_SIZE);
dim3 blocks=dim3(N/(2BLOCK_SIZE),N/(2BLOCK_SIZE));

matrixMul<<<blocks,threads,0,stream[0]>>>(d_c1,d_a1,d_b1,N/2,N/2);
matrixMul<<<blocks,threads,0,stream[1]>>>(d_c3,d_a1,d_b2,N/2,N/2);
matrixMul<<<blocks,threads,0,stream[2]>>>(d_c5,d_a3,d_b1,N/2,N/2);
matrixMul<<<blocks,threads,0,stream[3]>>>(d_c7,d_a3,d_b2,N/2,N/2);
matrixMul<<<blocks,threads,0,stream[0]>>>(d_c2,d_a2,d_b3,N/2,N/2);
matrixMul<<<blocks,threads,0,stream[1]>>>(d_c4,d_a2,d_b4,N/2,N/2);
matrixMul<<<blocks,threads,0,stream[2]>>>(d_c6,d_a4,d_b3,N/2,N/2);
matrixMul<<<blocks,threads,0,stream[3]>>>(d_c8,d_a4,d_b4,N/2,N/2);

addMat<<<blocks,threads,0,stream[0]>>>(d_C1,d_c1,d_c2);
addMat<<<blocks,threads,0,stream[1]>>>(d_C2,d_c3,d_c4);
addMat<<<blocks,threads,0,stream[2]>>>(d_C3,d_c5,d_c6);
addMat<<<blocks,threads,0,stream[3]>>>(d_C1,d_c7,d_c8);

cudaMemcpyAsync(C1,d_C1,size,cudaMemcpyDeviceToHost,stream[0]);
cudaMemcpyAsync(C2,d_C2,size,cudaMemcpyDeviceToHost,stream[1]);
cudaMemcpyAsync(C3,d_C3,size,cudaMemcpyDeviceToHost,stream[2]);
cudaMemcpyAsync(C4,d_C4,size,cudaMemcpyDeviceToHost,stream[3]);

cudaMemcpyAsync(c1,d_c1,size,cudaMemcpyDeviceToHost,stream[0]);
cudaMemcpyAsync(c3,d_c3,size,cudaMemcpyDeviceToHost,stream[1]);
cudaMemcpyAsync(c5,d_c5,size,cudaMemcpyDeviceToHost,stream[2]);
cudaMemcpyAsync(c7,d_c7,size,cudaMemcpyDeviceToHost,stream[3]);
cudaMemcpyAsync(c2,d_c2,size,cudaMemcpyDeviceToHost,stream[0]);
cudaMemcpyAsync(c4,d_c4,size,cudaMemcpyDeviceToHost,stream[1]);
cudaMemcpyAsync(c6,d_c6,size,cudaMemcpyDeviceToHost,stream[2]);
cudaMemcpyAsync(c8,d_c8,size,cudaMemcpyDeviceToHost,stream[3]);
cudaStreamSynchronize(stream[0]);
cudaStreamSynchronize(stream[1]);
cudaStreamSynchronize(stream[2]);
cudaStreamSynchronize(stream[3]);

我的问题：
第三步骤每个流中内核函数执行的矩阵乘法结果正确。但是第四步骤矩阵加法结果步骤错误，经过nisight monitor单步调试，发现每一个线程（譬如说线程块0中线程号tx=0，ty=0）执行了四次内核函数（执行完了一遍，又从内核函数开头执行，而且每一次内核函数传进的参数A,B也发生改变，我也没对传入的A,B参数进行操作，怎么会改变，这个怎么回事?为什么执行了四次，难道和四个流有关？）
加法内核函数代码
global void addMat(double* C,double* A,double* B)
{
// Block index
int bx = blockIdx.x;
int by = blockIdx.y;

// Thread index
int tx = threadIdx.x;
int ty = threadIdx.y;

int idx=tx+bxblockDim.x;
int idy=ty+byblockDim.y;

C[idy*(N/2)+idx]=A[idy*(N/2)+idx]+B[idy*(N/2)+idx];

变化时这样的
第一次：传入的参数值A[0]=4,B[0]=52，C是一个随机数字
第二次：传入的参数值A[0]=6，B[0]=62,C[0]=随机数
第三次：传入的参数值A[0]=36,B[0]=212,C[0]=随机数
第四次：传入参数值A[0]=70,B[0]=254,C[0]=56(这是我要的数据，怎么到现在才写入）
最终得到C[0]=324,请版主解惑。

2

楼主您说的现象（启动一次kernel, 结果发现kernel被启动了4次）是不可能的。

但是如果您启动了4次此kernel，则这个现象是完全可能的。同样的条件下的断点（例如，全局线程编号是x=888,y=777) 将会命中4次。因为4次同样的条件将会被命中。

这是对您的现象的描述。

关于您的上来的addMat结果不正常的问题，我看到了您的addMat使用了您原来的matrixMul的线程形状配置（具体数据不可知，因为您没有给出你的N之类的定义），请问此形状，在您的kernel中是否适用？（我看到了您的加法kernel中的坐标计算是y * N / 2 + x).

如果我是你，我会选择重新安排加法的线程形状，使得每个点正好被一个线程覆盖，这样清晰很多。（当然，如果您的N/2是您的目标矩阵的宽度的话，请您无视此条；如果不是，则您可能有漏掉的点，请您思考此建议）。

感谢来访。

3

dim3 threads=dim3(BLOCK_SIZE,BLOCK_SIZE);
dim3 blocks=dim3(N/(2BLOCK_SIZE),N/(2BLOCK_SIZE));
我想做的是：
矩阵乘法是两个规模为(N/2)x(N/2)的矩阵乘法
譬如说（N=4）一个流执行0 1 x 2 3 与 2 3 X 4 5
4 5 6 7 6 7 8 9
矩阵加法是两个规模为（N/2)X(N/2)的矩阵加法
譬如说将上个流的两个矩阵乘法结果相加。四个流都是类似的操作，并且操作的矩阵我都特意放到同一个流中
对应我上面写的代码是否有什么不合理的地方吗？请版主解惑？

4

既然楼主您的矩阵是(N/2) * (N/2)大小的（话说这个很奇葩），那么我没有问题了。

认真肉眼观测了下，在不越界的情况下，您的加法kernel无问题。
（是的，无问题，它就一句话里面进行加法，能有问题么）

唯一可能值得怀疑的是，
（1）您越界了，建议在加法kernel里用if判断下您的idx和idy，例如：
if (idx < N/2 && idy < N/2) {…}

或者
（2）或者实际上加法kernel运行无问题，但是您的验证过程有问题（包括但不限于您看错了）

请您尝试此2点。
感谢来访。

5

我也比较纳闷，我在每个流中只启动了一次addMat但是单步调试，一个线程块里的线程单步调试做完了每个操作又跳到传参的那一行，来回了4次。

6

哦，我想起来了，是不是这样的。他的四次虽然线程块和线程号是一样的，但分别是四个流的线程块线程号，就像有4个平行世界（每个流对应一个平行世界，线程块对应不同平行世界的同一个城市代号0，线程号对应不同平行世界的同一个人代号0）？

7

刚才LS说过了，你一共有4个kernel，你的条件断点（例如：idx == 0 && idy == 0),
将会被4次启动加法kernel中的对应条件的线程分别命中。

这难道不对么？

为何楼主不看解释，重新问了一遍？

您现在应该做的是：
单独抽取你的加法kernel出来，用2个测试数据（例如元素都是1和元素都是2的2个矩阵），调用该kernel进行加法，看能否得到元素都是3的矩阵。

如果你得到了（这应该的，该kernel极度简单，除了上面说的越界，不应该有其他问题），那么证明你的问题不在这里，这样可以排除它，从而可以检查其他方面。

你说呢？

8

嗯嗯，我再去检查检查