有一个二维矩阵，存储方式是：
[第一行][第二行]…
用gpu写了个矩阵算法，测试了2种模式：
- 一个block处理一行的数据，
- 一个block处理一列的数据
结果：按行处理的方法比按列处理的方法速度高非常多！
用visual profile分析发现，按列处理时，global memory load/store的速度非常慢，overhead很高
而按行处理时，global memory load/store的速度很快，overhead很低。
请问：有没有办法提升按列处理的速度？
编写了一段测试代码，
addKernel将列方向的每个元素与上下两个元素做平均。
addKernel2将行方向的每个元素与左右两个元素做平均。
这2个方法性能差别了整整1000倍！
有什么办法可以改进吗？
#include <stdio.h>
#include <math.h>
#include <time.h>
#include “cuda_runtime.h”
#include “device_launch_parameters.h”
#include <stdio.h>
cudaError_t addWithCuda(int cnt);
global void addKernel(int *a, int w, int h)
{
int x = blockIdx.x;
int y = threadIdx.x;
if (x >= w) return;
if (y >= h - 1) return;
if (y == 0) return;
int a1 = a[y * w + x];
int a2 = a[(y - 1) * w + x];
int a3 = a[(y + 1)* w + x];
a[y * w + x] = (a1 + a2 + a3) / 3;
}
global void addKernel2(int *a, int w, int h)
{
int y = blockIdx.x;
int x = threadIdx.x;
if (x >= w - 1) return;
if (x == 0) return;
if (y >= h) return;
int a1 = a[y * w + x];
int a2 = a[y * w + x - 1];
int a3 = a[y * w + x + 1];
a[y * w + x] = (a1 + a2 + a3) / 3;
}
int s_type = 0;
int main(int argc, char *argv)
{
int cnt = 100;
if (argc > 1)
cnt = atoi(argv[1]);
if (argc > 2)
s_type = atoi(argv[2]);
// Choose which GPU to run on, change this on a multi-GPU system.
cudaError_t cudaStatus = cudaSetDevice(0);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, “cudaSetDevice failed! Do you have a CUDA-capable GPU installed?”);
return 0;
}
addWithCuda(cnt);
// cudaThreadExit must be called before exiting in order for profiling and
// tracing tools such as Nsight and Visual Profiler to show complete traces.
cudaStatus = cudaThreadExit();
if (cudaStatus != cudaSuccess) {
fprintf(stderr, “cudaThreadExit failed!”);
return 1;
}
return 0;
}
// Helper function for using CUDA to add vectors in parallel.
cudaError_t addWithCuda(int cnt)
{
const int size = 2048 * 1024;
int * a = new int[size];
cudaStream_t stream;
int *dev_a = 0;
cudaError_t cudaStatus;
clock_t st, ed;
memset(a, 0, sizeof(int) * size);
if ((cudaStatus = cudaStreamCreate(&stream)) != cudaSuccess)
{
fprintf(stderr, “error create stream!\n”);
return cudaStatus;
}
// Allocate GPU buffers for three vectors (two input, one output) .
cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
if (cudaStatus != cudaSuccess) {
fprintf(stderr, “cudaMalloc failed!”);
goto Error;
}
// Copy input vectors from host memory to GPU buffers.
cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, “cudaMemcpy failed!”);
goto Error;
}
printf(“begin\n”);
st = clock();
for (int i = 0; i < cnt; i ++)
{
// Launch a kernel on the GPU with one thread for each element.
if (s_type == 0)
addKernel<<<2048, 1024, 0, stream>>>(dev_a, 2048, 1024);
else
addKernel2<<<1024, 2048, 0, stream>>>(dev_a, 2048, 1024);
}
// cudaThreadSynchronize waits for the kernel to finish, and returns
// any errors encountered during the launch.
cudaStatus = cudaStreamSynchronize(stream);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, “cudaThreadSynchronize returned error code %d after launching addKernel!\n”, cudaStatus);
goto Error;
}
ed = clock();
printf(“done:%d\n”, ed - st);
cudaStreamDestroy(stream);
// Copy output vector from GPU buffer to host memory.
cudaStatus = cudaMemcpy(a, dev_a, size * sizeof(int), cudaMemcpyDeviceToHost);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, “cudaMemcpy failed!”);
goto Error;
}
Error:
cudaFree(dev_a);
return cudaStatus;
}
是否按列访问就一定很慢，这个只能说一般情况下是的，
因为一般情况下我们的行存储优，这样在按行访问的是否是合并的，但是按列访问，2个元素之间相距了每行元素数*每元素大小这么长（假设每行末尾无空白元素），这样就导致严重的访问不合并的，从而性能下降。
但是如果你可以将数据按照列存储的，这样即无问题的（但按行访问却又会导致性能下降的）。
如果实在需要，您可以考虑将数据复制成2个副本，一个是另外一个的转置。这样可以同时按2种方式访问，满足您的需要。
感谢来访，周末愉快。
谢谢回复！
如果数据必须使用行存储方式，进行按列处理，有没有办法提升速度呢？
数据必须用行存储，不代表您在您的kernel前不可以转换一下存储方式。您觉得呢？
您也可以重新看下3#的最后一段。
就当是纯技术讨论吧。。。
行存储方式按列访问的问题无解？
[
不太明白您的意思。
一个block处理3行，在上面的例子2048x1024的矩阵，有1024-2个block，or，1024/3个block？
请重新阅读3#, 5#, 谢谢。
刚才在ICE重新提醒下，我再次阅读了您的代码：
当您的代码和您的文字描述不符的情况下，我将优先以你的文字描述为准。
因为我认为您的文字描述可能更能符合您的本意，毕竟您使用中文比使用CUDA的经验更为丰富。
谢谢您的回复！
我想实现的是：一个block处理一列，每个thread处理一列中的一个元素。
这个内核函数的代码，与我的文字描述内容不一致吗？
global void addKernel(int *a, int w, int h)
{
int x = blockIdx.x;
int y = threadIdx.x;
if (x >= w) return;
if (y >= h - 1) return;
if (y == 0) return;
int a1 = a[y * w + x];
int a2 = a[(y - 1) * w + x];
int a3 = a[(y + 1)* w + x];
__syncthreads(); /// 增加一个同步，以防止不同线程间读写的冲突。
a[y * w + x] = (a1 + a2 + a3) / 3;
}
你的文字是“按列访问”，而实际上您是“按行访问”，完全相反了。
实际上您的2个代码性能应该几乎一样的，别说差1000倍了（您原话），连1倍都差不上的。
（请确保您的2个kernel是同样的数据量启动的，以及，请确保时间正确，以及，却确保它们实现上都正确。）（只有实现上正确了，才能讨论优化）。
请三思上述三个询问。以及，我坚持我的它们不会差1000倍的观点，甚至1倍都差不上。（这个也是我和ICE的共同意见）。
谢谢回复！
请帮我看看下面代码的理解对不对。
按行存储的数据，指针为data。
行宽2048，列高1024。设w=2048, h=1024。
位于(x,y)的元素是指在第y行上的第x列的元素，它的位置应该是data + y * w + x。
例如下面的图示，按c语言习惯以0为基数，用()标示的元素(2,1)：
x=2
y=1 * * (*) * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * *
global void addKernel(int *a, int w, int h)
{
int x = blockIdx.x; //在同一个block里，每个线程处理的元素的x是一样的，因此也就是说一个block里所有线程处理的元素都在同一个列位置。
int y = threadIdx.x; //在线程里，处理的元素位于第y行。
if (x >= w) return;
if (y >= h - 1) return;
if (y == 0) return;
int a1 = a[y * w + x]; //(x, y) 元素
int a2 = a[(y - 1) * w + x]; //(x, y-1)
int a3 = a[(y + 1)* w + x]; //(x, y+1)
a[y * w + x] = (a1 + a2 + a3) / 3;
}
global void addKernel2(int *a, int w, int h)
{
int y = blockIdx.x; //在同一个block里，每个线程处理的元素的y是一样的，因此也就是说一个block里所有线程处理的元素都在同一个行位置。
int x = threadIdx.x;//在线程里，处理的元素位于第x列。
if (x >= w - 1) return;
if (x == 0) return;
if (y >= h) return;
int a1 = a[y * w + x]; //(x, y)
int a2 = a[y * w + x - 1]; //(x-1, y)
int a3 = a[y * w + x + 1]; //(x+1, y)
a[y * w + x] = (a1 + a2 + a3) / 3;
}
LZ您好：
看了您12#的回帖，又重新翻看了之前的各楼，现统一答复如下：
1：您的行优先存储的数据，行宽，列高，寻址各方面理解是正确的。
2：您的第一个kernel确实是按列访问的（即，相邻的threads访问同一列里面连续的位置）。以及您使用了int x=blockIdx.x的用法，这个用法虽然是正确的，但完全不符合一般习惯，具有极大的迷惑性。一般习惯上将threadIdx.x作为x方向上的索引，并和行优先存储联用。
以及您这样使用会造成强烈的合并访问问题，大为降低您的访存效率。
3：您的第二个kernel是按行访问的（即，相邻的threads访问同一行里面的连续数据）。以及您的索引用法是较为符合一般习惯的。此时基本能实现合并访问。
4：因为您使用了多个block，而在计算a[y * w + x] = (a1 + a2 + a3) / 3;的时候会读取到其他block所更新的值，但是block之间的顺序并不得到保证，因此您读到的可能是更新过的值，也可能是没有更新的值，所以您的这种写法结果具有随机性，可能无法实现您的意图。
5：您在2#给出的
if (s_type == 0)
addKernel<<<2048, 1024, 0, stream>>>(dev_a, 2048, 1024);
else
addKernel2<<<1024, 2048, 0, stream>>>(dev_a, 2048, 1024);
其中addKernel2启动时，您指定了一个block是2048个线程，这在当前的硬件上是无法执行的。
所以这可能也影响了您之前的测试。
6：对于按行优先排列的数据，如果您需要按列访问，那么请您合理安排您的算法实现，比如如果您要实现一列内部数据的相加之后取平均，您可以安排单一的线程在列方向读取3次数据，而此时相邻的线程实际上是访问同一行的数据，这就满足了合并访问。
或者您可以考虑先将部分数据按照合并访问的方式读入shared memory，之后再做使用。
或者其他什么解决方法。
总而言之，对于算法中要求的数据的按列获取，您不能拘泥于连续线程去访问这一列的低效方式，而应该合理安排具体实现中的线程行为，争取实现合并访问的高效访存方式。
大致如此，供您参考。
祝您编码顺利~
谢谢回复！非常感谢您指出的错误！
我修改了宽高，w=1024, h =1024.
并且在代码里读取a数组后，加入了__syncthread(),保证读写不冲突。测试程序里，这个错误对结果影响不大。
分别用了2个显卡测试，分别是660ti, 520mx，结论是相似的：“每个block处理一行”的速度为"每个block处理一列"的速度的10倍甚至更多一些！(测试多次，排除了其他应用的影响。)
问题：
LZ您好：
1：访存引起一个量级的速度差别还是基本靠谱的，以及您的“__syncthreads()”也许无法解决问题，请您给出代码。
2：您的问题留白了，请尽快补上。
祝您编码顺利~
我。。。不能说什么了。。。
楼主用threadIdx.x作x, 然后另外一个kernel却用blockIdx.x做x…
然后我只看到了您[(y + 1) * w + x]和[y *w + (x+1)]的区别，没有看到您将x和y也颠倒了。。。
这个真心无语中。。。
请让ICE服务您吧。
-
kernel函数里加了syn，防止读写冲突。您看看这么加对么？
global void addKernel(int *a, int w, int h)
{
int x = blockIdx.x;
int y = threadIdx.x;if (x >= w) return;
if (y >= h - 1) return;
if (y == 0) return;int a1 = a[y * w + x];
int a2 = a[(y - 1) * w + x];
int a3 = a[(y + 1)* w + x];__syncthreads();
a[y * w + x] = (a1 + a2 + a3) / 3;
}
global void addKernel2(int *a, int w, int h)
{
int y = blockIdx.x;
int x = threadIdx.x;
if (x >= w - 1) return;
if (x == 0) return;
if (y >= h) return;
int a1 = a[y * w + x];
int a2 = a[y * w + x - 1];
int a3 = a[y * w + x + 1];
__syncthreads();
a[y * w + x] = (a1 + a2 + a3) / 3;
}
- “一般习惯上将threadIdx.x作为x方向上的索引，并和行优先存储联用。” 我已经清楚的理解了这句话。
我的问题还是：按逐行存储数据方式，“一个block里处理一列数据”的条件下，有无办法提高效率呢？
改进的方法，“一个block处理相邻的4列数据”，可以提高速度40%左右，但还是与“一个block处理一行”的速度差距很大。
您看到的没错，我没写错，讨论的是在"一个block里处理一列数据"的问题。
我在寻求一些trick的方法来提升速度。
LZ您好：
1：按照您当前的计算规模，这个同步是可以的。您的第一个kernel，一个block就完成了一列的计算，以及可以让block内部的所有线程都先读取并等待，最后再一并写入。您的第二个kernel情况类似。
以及，这样的写法，您的计算规模是有限制的，如果一行的长度/一列的长度大于1024的时候，您就不能一行/一列只上一个block了。此时假如一列是两个block分管，那么在两个block交界的地方您无法确定是原始数据还是被改写过的。
不过，您依然可以通过每个线程干两倍的活来解决此问题。
2：就您的具体kernel，完全有在计算列平均的时候依然保持合并访问的写法。
您可以 ：
int x=threadIdx.x;
int y=blockIdx.x;
// 判断边界…
…
int a1 = a[y * w + x];
int a2 = a[(y - 1) * w + x];
int a3 = a[(y + 1)* w + x];
a_out[y * w + x] = (a1 + a2 + a3) / 3;
这样即可。
这样，您既保证了算法上是列方向做平均，也保证了访存效率，同时通过将结果写到a_out中，完全不修改原有的a，这样完全避免了任何同步问题。
总之，写代码是实现方法迁就硬件特性的。我们应该在硬件结构和算法逻辑两个约束下寻找高效的解决途径，而不是先定下一个不符合硬件结构的实现框架，然后试图寻找各种trick来弥补，这样显得事倍功半。这也是13#第6点的主旨所在。
祝您编码顺利~
LZ您好：
书写的逻辑性正确一般地只是我们的追求之一。
更高地追求是符合常规习惯，逻辑清晰，这样能提高效率，减少差错率，便于维护和交流。
因为据研究表明，很多差错出现在大脑按照固有模式自动推导的过程中，以及这个过程是不可控的。那么按照一般习惯书写，无论大脑是有意识地思考，还是按照固有模式自动推导，都不会产生问题。
这将减少人为差错率，提高效率。
同时，在考虑某个问题的时候，如果这个问题里面大部分因素是熟悉的用法或者符合习惯的，那么大脑只需要使用较少的资源便能掌控这些信息，并将大多数资源投入到核心的逻辑分析和思考中。
但如果反之，多数环节都是生疏的或者怪异的，那么大脑需要不断地反复重复和强调这些定义，以免思维落入常规习惯的自动推导过程中，这将极大地分散精力和注意力，难以解决最核心的问题。
工欲善其事，必先利其器——这在思维环节也是成立的。
大致如上，祝您编码顺利~