请教个问题, 在kernel函数里按列方式访问矩阵很慢吗?

有一个二维矩阵,存储方式是:
[第一行][第二行]…

用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服务您吧。

  1. 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;
}

  1. “一般习惯上将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您好:

书写的逻辑性正确一般地只是我们的追求之一。
更高地追求是符合常规习惯,逻辑清晰,这样能提高效率,减少差错率,便于维护和交流。

因为据研究表明,很多差错出现在大脑按照固有模式自动推导的过程中,以及这个过程是不可控的。那么按照一般习惯书写,无论大脑是有意识地思考,还是按照固有模式自动推导,都不会产生问题。

这将减少人为差错率,提高效率。

同时,在考虑某个问题的时候,如果这个问题里面大部分因素是熟悉的用法或者符合习惯的,那么大脑只需要使用较少的资源便能掌控这些信息,并将大多数资源投入到核心的逻辑分析和思考中。
但如果反之,多数环节都是生疏的或者怪异的,那么大脑需要不断地反复重复和强调这些定义,以免思维落入常规习惯的自动推导过程中,这将极大地分散精力和注意力,难以解决最核心的问题。

工欲善其事,必先利其器——这在思维环节也是成立的。

大致如上,祝您编码顺利~