复数矩阵加法cuda实现问题

[system]

如下两个kernel完成的复数矩阵加法运算

#define ROW 256
#define COL 256
#define TILE_WIDTH 16

kernel1:
//复数矩阵加法
global void GPU_MatAdd(double2 A,double2 B,double2 result)
{
int i=blockIdx.xblockDim.x+threadIdx.x;
int j=blockIdx.yblockDim.y+threadIdx.y;
int idx = jCOL+i;
if(i<ROW&&j<COL)
{
result[idx].x = A[idx].x + B[idx].x ;
result[idx].y = A[idx].y + B[idx].y ;
}
}

kernel1配置：
dim3 dimBlock(TILE_WIDTH,TILE_WIDTH);
dim3 dimGrid(ROW/TILE_WIDTH,COL/TILE_WIDTH);

kernel2：
//复数矩阵加法
global void GPU_MatAdd2(double2 A,double2 B,double2 result)
{
int i=blockIdx.xblockDim.x+threadIdx.x;
int j=blockIdx.yblockDim.y+threadIdx.y;
if(i<ROW/4 && j<COL/4)
{
int idx = jCOL+i;
result[idx].x = A[idx].x + B[idx].x ;
result[idx].y = A[idx].y + B[idx].y ;

	result[idx+16384].x = A[idx+16384].x + B[idx+16384].x;
	result[idx+16384].y = A[idx+16384].y + B[idx+16384].y;

	result[idx+16384*2].x = A[idx+16384*2].x + B[idx+16384*2].x;
	result[idx+16384*2].y = A[idx+16384*2].y + B[idx+16384*2].y;

	result[idx+16384*3].x = A[idx+16384*3].x + B[idx+16384*3].x;
	result[idx+16384*3].y = A[idx+16384*3].y + B[idx+16384*3].y;

}

}

kernel2配置：
dim3 dimBlock_x(TILE_WIDTH,TILE_WIDTH);
dim3 dimGrid_x(ROW/(TILE_WIDTH2),COL/(TILE_WIDTH2));

问题：
kernel1的结果正确，kernel2结果错误，麻烦帮我看看问题出在哪了？谢谢！

[system]

楼主您好：
您这样安排形状，有的元素被覆盖多次，有的元素没有线程去处理啊。

例如反例：您看下标为(100,100)的元素，有线程负责它么？果断没有的。

您应当这样设置线程形状：
从：
dim3 dimBlock_x(TILE_WIDTH,TILE_WIDTH);
dim3 dimGrid_x(ROW/(TILE_WIDTH2),COL/(TILE_WIDTH2));

改成：
dim3 dimBlock_x(TILE_WIDTH,TILE_WIDTH);
dim3 dimGrid_x(ROW / TILE_WIDTH,COL/(TILE_WIDTH * 4));
（请您注意我将一个2移动到后面了，变成了4）

这样可以不多不少的覆盖到的。您觉得呢？
（以及，需要说明的是，您的row和col是反的么？当然，您这里都是256, 无碍，但以后可能是个隐患）

[system]

@"system"

楼主您好：
您这样安排形状，有的元素被覆盖多次，有的元素没有线程去处理啊。

修改完毕后，别忘记该if中的条件，改成：
if (i < ROW&& j < COL / 4)

（这里的col依然是您的行的意思）
（建议以后不要使用反意的词汇，避免给自己麻烦）

[system]

@"system"

修改完毕后，别忘记该if中的条件，改成：
if (i < ROW&& j < COL / 4)

玫瑰版主您好
（1）
dim3 dimBlock_x(TILE_WIDTH,TILE_WIDTH);
dim3 dimGrid_x(ROW / TILE_WIDTH,COL/(TILE_WIDTH * 4));
这里的ROW 和 COL 应该没反吧，只是if (i < ROW&& j < COL / 4) 这里反了 对吗？
（2）修改Grid大小不会影响占用率吧，只有修改block大小才会影响占用率，我的理解对吗？
（3）如果把grid的大小缩小4倍 与 把block大小缩小4倍 本质的区别是什么？

[system]

@"system"

玫瑰版主您好
（1）
dim3 dimBlock_x(TILE_WIDTH,TILE_WIDTH);

楼主您好：

（1）一般情况下，我们写变量名称，各种符号，使用的是词语的本意。
所以一般的，我们将row作为“行数”的标记，而将“column"作为列的标记。

但是词语是死的，人是活的，您也可以反过来用，例如将“row”记作列数或者编号的。
这个是您的自由。版主是为您服务的，有权为您建议，但最终的决定权在您自身。

所以您完全可以不采纳建议的。
（玫瑰完全不会生气的）

（2）将grid里的block规模扩大4倍和将block里的线程数目数目缩小4倍没有什么本质区别，但请您注意，如果您的线程总数要确保一样，请您同时进行此两者。
否则您的线程数量将发生变化。

请您思考此两点。

[system]

@"system"

楼主您好：

（1）一般情况下，我们写变量名称，各种符号，使用的是词语的本意。

版主没有理解我的意思，我再补充一下

（1）这里我写错了 if(i<ROW && j<COL/4)
{
int idx = jCOL+i;
…
}
修改为
if(i<ROW && j<COL/4)
{
int idx = jROW+i;
…
}

就可以了把。
（2）我所说的grid和block大小的改变是说，如果我把grid缩小4倍或者block缩小4倍，我在kernel中的线程就要多做4倍的计算量，我比较的是grid缩小4倍（block大小不变）与 block缩小4倍（grid大小不变） 它们对于占用率的影响，我测试的结果是grid缩小4倍之后占用率没有受影响，而block缩小4倍占用率从100%变为了50%。

[system]

@"system"

版主没有理解我的意思，我再补充一下，（1）这里我写错了 if(i

…

楼主啊！你的数组是qz[50][60], 这代表了它有60列(column)，50行(row)啊！！
（不是你认为的60行，50列！）

所以你的任何一个坐标(x,y)实际上线性化后是y * col + x, 即y * 60 + x
而不是您想象的y * row + x…（即：y * 50 + x是错误的）

请您三思吧。

[system]

@"system"

版主没有理解我的意思，我再补充一下

（1）这里我写错了 if(i

补充一下7#玫瑰斑竹，关于LZ 6#的第二个问题：

如果LZ保证总计算量不变，而更改单个线程的计算量，从而改变block的大小或者grid的大小。

根据您前述的算法，此时 一个线程所完成的工作依然是简单的，并不会在这里产生瓶颈。

那么如果您的threads per block数量在缩小4倍以后，依然有128/192/256这样的适当规模；或者您的blocks per grid数量在缩小4倍以后，依然有足够多的数量以便充分给各个SM分配任务，那么两者的实现效率将不会明显差别。

否则，如果有任何一项变得过小，将影响您的程序效率。

总之，这是一个多因素制约的问题，您需要考虑其间的制约关系，寻找最优实现，而不能简单说改变某个参数就一定会怎样。

祝您编码顺利~

[system]

@"system"

补充一下7#玫瑰斑竹，关于LZ 6#的第二个问题：

如果LZ保证总计算量不变，而更改单个线程的计算量，从而 …

ice版主你好，我看了Volkov的关于 “低occupancy达到高performance” 的一个文档，我就想做个实验，发现对于矩阵加法和点乘，执行效率没什么效果。

[system]

@"system"

ice版主你好，我看了Volkov的关于 “低occupancy达到高performance” 的一个文档，我就想做个实验，发现 …

LZ您好：

对于有些问题和精心设计的kernel确实可以使用较少的线程达到良好的效果。
但是对于一般情况，请尽量提升occupancy，以利用线程的自动切换以达到较好的效果。

祝您好运~

[system]

@"system"

…

楼主啊！你的数组是qz[50][60], 这代表了它有60列(column)，50行(row)啊！！

int i=blockIdx.xblockDim.x+threadIdx.x;
int j=blockIdx.yblockDim.y+threadIdx.y;
if(i<COL && j<ROW/4)
{
int idx = j*COL+i;
。。。
}
那就是上面这样了？
我对于blockIdx.x表示行还是列混乱了，我的程序中的 i 计算得到的是列，j 是行吧?

[system]

@"system"

LZ您好：

对于有些问题和精心设计的kernel确实可以使用较少的线程达到良好的效果。

好的 谢谢ice版主

[system]

@"system"

int i=blockIdx.xblockDim.x+threadIdx.x;
int j=blockIdx.yblockDim.y+threadIdx.y;
…

LZ您好：

我来就行，列，行元素个数，列元素个数等问题简要说下：

以一个三行五列的矩阵为例。

按照数学上的一般说法，row=3，column=5。
同时有，行的长度为列数，列的长度为行数，即row length=column=5，column length=row=3。

现在考虑c风格的行优先存储，按照此时后面的指标先变化的规律，应写为qz[3][5]。(即前面是行数，后面是列数；前面是列长度，后面是行长度。按行存储，后面的变量先变化。)

此时考虑用CUDA，并考虑使用二维的block，grid等对应上述存储规律。
仅以block为例，blockDim.x表示block在x方向的元素数/长度，blockDim.y表示block在y方向的元素/长度。
同时，block内部的threadIdx是blockIdx.x先变化，然后是blockIdx.y后变化。

那么此时，一般习惯让线程编号递增规律和数据存储规律相一致，即线程编号的x方向对应于数据存储的一行，一行线程处理一行矩阵的数据，（blockDim.x对应于row length），这样相邻的线程总是访问相邻的数据，适于合并访问。

大致如此，请您在思考的时候，尽量单一以row length和column length考虑并明确参数含义，而不是混合考虑row length，column length，row，column这四个概念，而产生各种困扰。

祝您编码顺利~

[system]

@"system"

LZ您好：

我来就行，列，行元素个数，列元素个数等问题简要说下：

这次完全明白了 非常感谢

[system]

@"system"

这次完全明白了 非常感谢

不客气的，欢迎您常来论坛~

祝您编码顺利~