首先我要说的是,我仔细阅读了关于SO的所有类似问题:
我的目的是尝试动态计算(而不是硬编码值)我正在开发的前馈神经网络库.
我的数据不是正方形点阵(矩阵),正如我所看到的大多数例子一样,而是两个向量产生矩阵,行到行的行数不等:
float x[6] {1.f, 1.f, 0.f, 1.f, 1.f, 0.f};
thrust::device_vector<float> in_vec( x, x+6 );
float y[9] {1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f};
thrust::device_vector<float> w_vec( y, y+9 );
thrust::device_vector<float> o_wec(9);
thrust::device_vector<float> mtx_vec( 9 * 6 );
float * i_ptr = thrust::raw_pointer_cast( in_vec.data() );
float * w_ptr = thrust::raw_pointer_cast( w_vec.data() );
float * out_ptr = thrust::raw_pointer_cast( mtx_vec.data() );
dim3 threadsPerBlock(9,6);
dim3 numBlocks(1,1);
prop_mtx<<<numBlocks,threadsPerBlock>>>( w_ptr, i_ptr, out_ptr, 6 );
和内核:
__global__ void prop_mtx( float * w, float * i, float * o, int s )
{
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
o[y + x * s] = w[x] * i[y];
}
我采用这种方法的原因是因为在矢量/矩阵计算方面,它在ANN计算中是有意义的.我想保持这一点,AFAIK使用2D网格进行重量*输入计算是合理的.
我必须将每个块的线程计算为2D,并在网格中使用不等数量的线程.
我正在使用GTX 660,它具有:
CUDA Capability Major/Minor version number: 3.0
Total amount of global memory: 2047 MBytes
( 5) Multiprocessors, (192) CUDA Cores/MP: 960 CUDA Cores
Maximum Texture Dimension Size (x,y,z) 1D=(65536), 2D=(65536, 65536), 3D=(4096, 4096, 4096)
Warp size: 32
Maximum number of threads per multiprocessor: 2048
Maximum number of threads per block: 1024
Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
Max dimension size of a grid size (x,y,z): (2147483647, 65535, 65535)
我试图了解如何推断/计算网格大小,每个块的线程数和块数.
我们假设我有一个800项的权重向量,以及6500项的输入向量.
我知道每个块的最大线程数是1024,但由于它是一个2D网格,它更可能是:
dim3 threadPerBlock(X,Y);
由于我的网格不是方形矩阵,我需要以不同的方式计算每个块的X,Y线程？
或者我需要先推断出所需的块数？
最后,因为我的线程warp大小是32,
任何伪代码,或解释我应该怎么做,将不胜感激.
我试过的是通过将我的数据除以32个包裹大小来计算我的2D网格大小.然后我考虑使用可用的SM计算网格线程.例如
800 weights / 5 SM, = 160 x's per SM
6500 inputs / 5 SM, = 1300 y's per SM
但我不知道该怎么办.最后,我考虑首先找到输入重量比:
6500/800 = 8.125
暗示使用32的最小网格大小X,Y必须乘以8.125*32因此,我的threadsPerBlock将是:
dim3 threadsPerBlock(32,260);
当然,每块8320个线程,远远超过每块1024个线程.
所以这是我的问题:如何保持每个块的1024个线程,同时保留我的数据的正确网格大小？
PS:我的问题不是优化代码,而是了解如何在设备上分配线程和网格数据.
分类计算问题的一种方法是讨论转换和减少.
甲减少是一个问题类别,它需要一个大的输入数据集的大小,并产生小的输出数据集大小.例如,拍摄图像并找到最大像素值将是减少.对于这个讨论,我们将忽略减少.
甲变换是计算的类别,其中所述输出数据集的大小(元素的数量)可以是"大"或"大致相同的"作为输入的数据集的大小.例如,拍摄图像并生成模糊图像将是一种转变.
对于 转换,编写cuda内核(线程代码)的常用方法("线程策略")将使一个唯一的线程负责输出数组中的每个点.因此,我必须拥有的最小线程总数等于我的输出数组的大小.线程代码只是输入数据所需的一组计算,以便产生一个输出数据点.粗略地说,你的问题和简化的内核符合这个定义; 这是一种转变.
遵循上述线程策略,我们需要网格中的线程总数等于我需要创建的输出点总数.对于2D问题,通常可以方便地考虑这些二维问题,并且为此目的,CUDA提供2D(或3D)线程块组织和2D(或3D)网格组织.
选择CUDA螺纹块尺寸通常有些随意.一般来说,我们通常希望针对每个块范围内128到512个线程的线程块(出于其他地方所涵盖的原因),并且我们希望线程块是整数倍的32(warp大小),以便在线程块获得时获得效率细分为warp,它是CUDA执行的实际单位.在当前支持的GPU上,线程块限制为每块1024个线程(总数 - 即维度的乘积).但是,对于许多问题,此范围内的线程块选择(例如,256个线程与512个线程)通常对性能的影响相对较小.为了让工作变得有效,我们不会在这一点上详述.(当你回来进行优化时,你可能会重新考虑这个选择.)
到目前为止,我们已经了解到,对于这个问题类型,我们需要一个总线程数来覆盖我们的问题空间,我们将有一个有点任意的线程块维度选择.所以让我们选择(32,16)(x,y)开始,总共512个线程.没有规则声明存在块需要"正方形",或者网格需要"正方形",或者在线程块尺寸和问题大小(或网格尺寸)之间甚至应该存在任何比例奇偶校验.
既然我们有一个(32,16)的线程块选择,我们必须问自己"我需要多少这些？".这个问题是2D,因此我们选择了一个2D线程块,以简化线程代码中的索引生成.我们也选择一个2D网格 - 它对于2D问题是有意义的,并且对于2D简单的索引生成也是如此.所以我们可以独立考虑这两个维度.
那么,在x方向上需要多少块？我至少需要(我的问题大小为x)/(我的x中的线程块大小).由于我们在这里处理所有整数,这就引出了一个问题"如果我的问题大小不能被我的线程块大小整除？" 规范的解决方案是发射足够多的线来覆盖空间,或者足够的块来覆盖空间.但是在非均匀可分的情况下,这将导致"额外的线程".我们将很快讨论和处理这些问题.因此,如果我有一个这样的dim3变量用于threadblock维度:
#define BX 32
#define BY 16
...
dim3 block(BX,BY);
然后我可以像这样构建我的dim3网格变量:
#define DX 800
#define DY 6500
...
dim3 grid((DX+block.x-1)/block.x, (DY+block.y-1)/block.y);
如果你完成这个算术,你会发现这导致我们在x和y方向上启动足够的块,这样我们将至少有足够的线程来覆盖我们的问题空间(DX,DY),每个输出一个线程点.
希望很明显Y维度是独立处理的,与x维度无关.
上述计算通常会导致在我的网格中生成"太多"线程.在我需要处理的问题空间(DX,DY)的末尾,我会有一些"额外的线程".我们希望这些线程"什么都不做".处理这个问题的规范方法是将问题空间维度传递给我的内核,在我的内核中创建一个适当的全局唯一线程索引,然后将该索引与我的问题空间中的最大索引进行比较.如果超过它,我们只需要该线程跳过所有剩余的线程代码.
以您的内核为例,它可能如下所示:
__global__ void prop_mtx( float * w, float * i, float * o, int s, const size_t d_size_x, const size_t d_size_y )
{
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
if ((x < d_size_x) && (y < d_size_y)) // thread check
o[y + x * s] = w[x] * i[y];
}
请注意,这样的线程检查将创建在后续代码中"不参与"的线程(在某些块中).这里要注意的一点是,使用__syncthreads()取决于参与的块中的所有线程.因此,我们不应该__syncthreads()直接在这种情况下使用.相反,我们必须适当地调整线程块行为:
__global__ void prop_mtx( float * w, float * i, float * o, int s, const size_t d_size_x, const size_t d_size_y )
{
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
if ((x < d_size_x) && (y < d_size_y)) // thread check
{
o[y + x * s] = w[x] * i[y];
// and other code not dependent on __syncthreads()
}
// now it is safe to use since all threads are participating
__syncthreads();
if ((x < d_size_x) && (y < d_size_y)) // thread check
{
// rest of kernel code
}
}
请注意,可以使用较少数量的线程为大量输出数据点执行必要的计算.线程和输出数据之间的1:1对应关系是思考和编写cuda内核代码的简单方法,但这不是唯一的方法.另一种可能的方法是使用某种形式的网格跨越环,以便较小的网格可以覆盖更大的问题空间.对这些策略的讨论超出了本答案的范围,在应对其他方法之前,应该理解本答案中讨论的基本方法.