我正在学习如何为 GPU 优化代码。我读到了内存局部性的重要性。我也看到了一些 tutorials和 examples GPU卷积。基于此，我编写并测试了几个自己的内核。令人惊讶的是我发现最简单的天真的内核是最快的!？ 它比 CPU 快 <10 倍。 (是的，我通过运行 kenrnel 64x 来分摊上传/下载时间)。
我做错了什么？ 我希望卷积就是优化 GPU 的那种操作。如果我能得到100x speed-up on matrix multiplication ，为什么卷积这么慢？
性能 [CPU 滴答数/像素](越低越好):
编辑:GPU-scanline_async 我是在阅读了关于
async_work_group_copy 的建议后制作的我想知道两件事:
测试是在 my laptop 上完成的使用 CPU Intel Core i7 6700HQ Skylake 和 GPU nVidia 960M，通过在 256x256 像素的浮点阵列上以 64x/帧的速度运行内核。 code full can be seen here .
============ 内核代码 ============
内核 GPU-Naive 二维全局=(256,256)局部=(16,16)
__kernel void blur2D_naive(
__global float* I,
__global float* O
){
const int ix = get_global_id (0)+1;
const int iy = get_global_id (1)+1;
const int nx = get_global_size(0)+2;
int i = iy * nx + ix;
// 1.6 ticks/pixel
O[i] =( I[i-nx-1] + I[i-nx] + I[i-nx+1] +
I[i -1] + I[i ] + I[i +1] +
I[i+nx-1] + I[i+nx] + I[i+nx+1] ) * 0.11111111111;
// modified with gaussian mask 4.9 ticks/pixel
//O[i] =( 0.0625*I[i-nx-1] + 0.125*I[i-nx] + 0.0625*I[i-nx+1] +
// 0.125 *I[i -1] + 0.25 *I[i ] + 0.125 *I[i +1] +
// 0.0625*I[i+nx-1] + 0.125*I[i+nx] + 0.0625*I[i+nx+1] );
}
内核 GPU-local 二维全局=(256,256)局部=(16,16)
#define NBx 18 // tile size including borders [halo] 16+2
#define NBy 18
// seems to be slower than naive method
__kernel void blur2D_local(
__global float* I,
__global float* O
){
__local float L[NBx*NBy];
const int2 iG = (int2)(get_global_id (0)+1 , get_global_id (1)+1 );
const int2 nG = (int2)(get_global_size(0)+2 , get_global_size(1)+2 );
const int2 iL = (int2)(get_local_id (0)+1 , get_local_id (1)+1 );
const int2 nL = (int2)(get_local_size (0)+2 , get_local_size (1)+2 );
const int2 iGR = (int2)(get_group_id (0) , get_group_id (1) );
// copy boundary pixels to local memory
switch( get_local_id(1) ){ // some threads copy one more of boundary (halo) pixels
case 4:
switch( get_local_id(0) ){ // copy corner points
case 0: L[ 0 ] = I[ nG.x* get_group_id(1)*get_local_size(1) + get_group_id(0)*get_local_size(0) ]; break; // upper-left
case 1: L[ NBx-1 ] = I[ nG.x* get_group_id(1)*get_local_size(1) + get_group_id(0)*get_local_size(0)+(NBx-1) ]; break; // upper-right
case 2: L[ (NBy-1)*NBx ] = I[ nG.x*(get_group_id(1)*get_local_size(1)+(NBy-1)) + get_group_id(0)*get_local_size(0) ]; break; // lower-left
case 3: L[ NBy* NBx-1 ] = I[ nG.x*(get_group_id(1)*get_local_size(1)+(NBy-1)) + get_group_id(0)*get_local_size(0)+(NBx-1) ]; break; // lower-rigth
}
// copy border lines
case 0: L[ iL.x ] = I[ nG.x* get_group_id(1)*get_local_size(1) + iG.x ]; break; // top line
case 1: L[ NBx*(NBy-1) + iL.x ] = I[ nG.x*(get_group_id(1)*get_local_size(1)+(NBy-1) ) + iG.x ]; break; // botton line
case 2: L[ NBx*iL.x ] = I[ nG.x*(get_group_id(1)*get_local_size(1)+get_local_id(0) ) + get_group_id(0)*get_local_size(0) ]; break; // left line
case 3: L[ NBx*iL.x + (NBx-1) ] = I[ nG.x*(get_group_id(1)*get_local_size(1)+get_local_id(0) ) + (get_group_id(0)*get_local_size(0)+(NBx-1)) ]; break; // right line
} // each thread coppied at max. 1 border pixels
int ig = iG.y*nG.x + iG.x;
int il = iL.y*nL.x + iL.x;
L[il] = I[ig]; // each thread copy his pixel to local memory
barrier(CLK_LOCAL_MEM_FENCE);
const float renorm = 1.0/9.0;
O[ig] =( L[il-NBx-1] + L[il-NBx] + L[il-NBx+1] +
L[il -1] + L[il ] + L[il +1] +
L[il+NBx-1] + L[il+NBx] + L[il+NBx+1] ) / 9.0;
}
内核 GPU-local_async 二维全局=(256,16) 局部=(16,16)
#define nTiles 16
#define NBx 18
#define NBy 18
#define copy_tile(event,ig0,I,L) { int ig_=ig0; int il_=0; for(int i=0; i<NBy; i++){ event = async_work_group_copy( L+il_, I+ig_, NBx, event ); ig_+=nx; il_+=NBx; } }
// https://streamcomputing.eu/blog/2014-06-19/using-async_work_group_copy-on-2d-data/
__kernel void blur2D_local_async(
__global float* I,
__global float* O
){
const int nx = get_global_size(0)+2;
__local float LI[NBx*NBy*2];
int iL0 = 0;
int iL1 = NBx*NBy;
event_t event = 0;
int ig0 = get_group_id(0)*get_local_size(0);
copy_tile(event,ig0,I,LI);
for( int it=0; it<nTiles; it++ ){
int ig = ig0 + (get_local_id(1)+1)*nx + get_local_id(0)+1;
int il = (get_local_id(1)+1)*NBx + get_local_id(0) + iL0;
ig0 += get_local_size(1)*nx;
event_t event_ = 0;
copy_tile(event_,ig0,I,LI+iL1);
wait_group_events(1, &event);
//barrier(CLK_LOCAL_MEM_FENCE);
O[ig] =( LI[il-NBx] + LI[il-NBx+1] + LI[il-NBx+2] +
LI[il ] + LI[il +1] + LI[il +2] +
LI[il+NBx] + LI[il+NBx+1] + LI[il+NBx+2] ) * 0.11111111111;
int iLtmp=iL0; iL0=iL1; iL1=iLtmp;
event = event_;
}
}
内核 GPU-scanline_private 一维全局=(256) 局部=(32)
__kernel void blur2D_scanline_priv(
int nx, int ny,
__global float* I,
__global float* O
){
int ig = get_global_id(0)+1;
float3 Lm = (float3)( I[ig-1], I[ig], I[ig+1] ); ig += nx;
float3 L0 = (float3)( I[ig-1], I[ig], I[ig+1] );
for(int iy=1; iy<(ny-1); iy++ ){
ig += nx;
float3 Lp= (float3)( I[ig-1], I[ig], I[ig+1] );
O[ig-nx] =
( Lm.x + Lm.y + Lm.z +
L0.x + L0.y + L0.z +
Lp.x + Lp.y + Lp.z ) * 0.11111111111;
Lm=L0; L0=Lp;
}
}
内核 GPU-scanline_async 一维全局=(256) 局部=(32)
#define NB 34
__kernel void blur2D_scanline_async(
int nx, int ny,
__global float* I,
__global float* O
){
__local float L[NB*4];
int i0=0;
int i1=NB;
int i2=NB*2;
int i3=NB*3;
event_t event = 0;
int ig0 = get_group_id(0)*get_local_size(0);
event = async_work_group_copy( L , I+ig0, NB, event ); ig0 += nx;
event = async_work_group_copy( L+NB , I+ig0, NB, event ); ig0 += nx;
event = async_work_group_copy( L+NB*2, I+ig0, NB, event ); ig0 += nx;
const int il = get_local_id(0);
int ig = get_global_id(0)+1;
for(int iy=1; iy<(ny-2); iy++ ){
wait_group_events(1, &event);
event = async_work_group_copy( L+i3, I+ig0, NB, event ); ig0 += nx;
ig += nx;
O[ig] =
( L[i0+il] + L[i0+il+1] + L[i0+il+2] +
L[i1+il] + L[i1+il+1] + L[i1+il+2] +
L[i2+il] + L[i2+il+1] + L[i2+il+2] ) * 0.11111111111;
__local float *Ltmp;
int itmp=i0; i0=i1; i1=i2; i2=i3; i3=itmp;
}
}
内核 CPU-naive
void blur(int nx, int ny, float * I, float * O ){
float renorm = 1.0/9.0;
for(int iy=1;iy<ny-1;iy++){ for(int ix=1;ix<nx-1;ix++){
int i = iy*nx+ix;
O[i] =( I[i-nx-1] + I[i-nx] + I[i-nx+1] +
I[i -1] + I[i ] + I[i +1] +
I[i+nx-1] + I[i+nx] + I[i+nx+1] ) * renorm;
} }
}
最佳答案
在矩阵乘法中，每个子矩阵(补丁)用于另一个矩阵中所有行中的所有补丁。如果补丁中有 2x2 子矩阵，如果主矩阵是 20x20，那么每个子矩阵被用于乘法 10 次。 GPU 通常使用 16x16 或 32x32 大小的补丁，这意味着，对于 2kx2k 乘法，每个 16x16 补丁至少重复使用 128 次。
MM reuse = 128
并添加子矩阵-子矩阵乘法重用，足以将gpu推到极限。
在 3x3 卷积中，3x3 patch 不用于整个扫描线或整个图片。只有它的像素被重用。
3x3 模板:每个像素都被相邻的 8 个模板重复使用。
5x5 模板:每个像素都被相邻的 24 个模板重复使用。
要 catch 矩阵乘法，它需要
11x11 stencil to have a reuse of 120
这也比矩阵乘法更局部，应该得到比它更多的 gflops，但它没有做等量的乘法和加法。
它正在做 9 次加法 + 1 次乘法。
8 个潜在的乘法丢失。 GFLOPS 限制的近一半丢失。
您应该尝试异步工作组副本。
矩阵乘法/具有 16x16 子矩阵)与卷积(17x17 画笔大小):
拥有一个加法重的内核为另一个乘法重的内核同时工作或在同一个内核中异步工作留下了空间。也许如果您将它用于图像处理，也许您可以在其中添加一些“混合”或“调整大小”内核，以便它们一起工作？
扫描线版本正在加载 3 个元素，执行 9 个 add + 1 mul 然后重复，加载的元素停留 3 圈，这意味着它们仅被重复使用 3 次，其邻居(x 或 y 方向)可能不会落入邻居线程甚至邻居工作组。此外，3 个负载与 1 个存储是不平衡的。如果内存带宽为 100 GB/s，那么它将使用 50 GB/s 的负载，15 GB/s 的存储，除非它们来自 L1。
您可以使用累加器减少加/乘不平衡。
store = (accumulator) * 0.1111111
accumulator+=new vector // 3 adds
accumulator-=old vecotr // 3 adds
所以现在是 6 个 add + 1 muls，所以更平衡:1Tflops GPU 将有 500Gflops 用于添加，90 Gflops 用于 muls。
Naive 版本不使用本地内存，为更多的波前飞行留出更多空间。本地内存版本实际上打破了 L1 访问模式，并减少了飞行中的波前。这减少了VALU占用。
您可以通过在工作组级别而不是线程级别执行 scanline 来减少本地内存使用量。我的意思是这样的:
从内存加载:x x x x x x x x x x
为它做扫描线:(从左到右，1-D)a b c d e f g h i j
现在将它用于工作组级别的扫描线:a c c um ul a t or r (+new)
(从上到下)z x z x z x z x z x(-旧)
calculate frontline 1-d scanline: 30 additions for each new row
calculate wide vector 2-d scanline:30*30 additions
each pixel get 1 value instead of adding 3 values
storing: 16x16 multiplications
much less local memory used, more balanced (~8 add 1 mul)
这具有一维扫描线，它是 N 个周期的单线程或 LogN 个周期的多线程减少(考虑到计算单元中的足够线程)。
关于c - 不应该在 GPU (OpenCL) 上更快地进行 3x3 卷积，我们在Stack Overflow上找到一个类似的问题： https://stackoverflow.com/questions/43217449/