有一个二维矩阵,存储方式是:
[第一行][第二行]…
用gpu写了个矩阵算法,测试了2种模式:
- 一个block处理一行的数据,
- 一个block处理一列的数据
结果:按行处理的方法比按列处理的方法速度高非常多!
用visual profile分析发现,按列处理时,global memory load/store的速度非常慢,overhead很高
而按行处理时,global memory load/store的速度很快,overhead很低。
请问:有没有办法提升按列处理的速度?
有一个二维矩阵,存储方式是:
[第一行][第二行]…
用gpu写了个矩阵算法,测试了2种模式:
结果:按行处理的方法比按列处理的方法速度高非常多!
用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;
}
您看到的没错,我没写错,讨论的是在"一个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您好:
书写的逻辑性正确一般地只是我们的追求之一。
更高地追求是符合常规习惯,逻辑清晰,这样能提高效率,减少差错率,便于维护和交流。
因为据研究表明,很多差错出现在大脑按照固有模式自动推导的过程中,以及这个过程是不可控的。那么按照一般习惯书写,无论大脑是有意识地思考,还是按照固有模式自动推导,都不会产生问题。
这将减少人为差错率,提高效率。
同时,在考虑某个问题的时候,如果这个问题里面大部分因素是熟悉的用法或者符合习惯的,那么大脑只需要使用较少的资源便能掌控这些信息,并将大多数资源投入到核心的逻辑分析和思考中。
但如果反之,多数环节都是生疏的或者怪异的,那么大脑需要不断地反复重复和强调这些定义,以免思维落入常规习惯的自动推导过程中,这将极大地分散精力和注意力,难以解决最核心的问题。
工欲善其事,必先利其器——这在思维环节也是成立的。
大致如上,祝您编码顺利~