关于“流”的问题

我的程序的步骤是这样的:
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+by
blockDim.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,请版主解惑。

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

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

这是对您的现象的描述。

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

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

感谢来访。

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)的矩阵加法
譬如说将上个流的两个矩阵乘法结果相加。四个流都是类似的操作,并且操作的矩阵我都特意放到同一个流中
对应我上面写的代码是否有什么不合理的地方吗?请版主解惑?

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

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

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

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

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

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

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

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

这难道不对么?

为何楼主不看解释,重新问了一遍?

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

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

你说呢?

嗯嗯,我再去检查检查