线程网格(grid)

        在并行运算中,合理处理线程网格可以为程序获得更高的性能加速比,如何合理使用线程网格,才能使并行程序效率更快?

      一个线程网格是由若干个线程块组成的,每个线程块是二维的,分为X轴与Y轴。此时,每次最多能开启Y*X*T个线程。现在,我们用一个实例进行深入理解。为简单期间,我们限制Y轴方向只有一行线程。
     假设我们现在正在看一张标准高清图片、这张图片的分辨率为1920 x 1 080。通常线程块中线程数量最好是一个线程束大小的整数倍,即32的整数倍。由于设备是以整个线程束为单位进行调度,如果我们不把线程块上的线程数目设成32的整数倍,则最后一个线程束中有一部分线程是没有用的。因此我们必须设置一个条件进行限制,防止处理的元素超出x轴方向上所规定的范围。在接下来的内容中我们会看到,如果不这样做,程序的性能将会降低。为了防止不合理的内存合并,我们尽量做到内存的分布于线程的分布达到一一映射的关系:如果我们没能做到这点,程序的性能可能会降低很多。在程序中,要尽量避免使用小的线程块,因为这样做可以充分利用硬件。在这个例子中,我们将每个线程块上开启192个线程。通常,192是我们所考虑的最少的线程数目。每个线程块192个线程,我们很容易算出处理一行图像需要10个线程块(如图)。

在这里,选择192是因为X轴方向处理的数据大小是它的整数倍,又是线程束大小的整数倍,这使我们的编程更加方便。在实际编程中,我们也要尽力做到这一点。

    在X轴方向的顶部我们可以得到线程的索引,在Y轴方向我们可以得到行号。由于每一行只处理了一行像素,每一行有10个线程块,因此我们需要1080行来处理整张图片,一共1080*10=10800个线程块。按照这种一个线程处理一个像素的方式,每个线程块开启了192个线程,一种调度了200多万个线程。我们对单个像素或数据进行单一处理,或者对同一行的数据进行处理时,这种特殊的布局方式是很有用的在当前费米架构的硬件上,一个SM可以处理8个线程块,所以上述程序从应用层的角度来说一共需要1350(10 800个线程块除以每个SM能调度的8线程块)SM来完全实现并行。但当前费米架构的硬件16SM可供使用(GT x 580 ),即每个SM将被分配675个线程块进行处理

      上述这个例子很简单,数据分布对齐因此我们很容易就找出了一个好的解决方案,但如果我们的数据不是基于行的,该如何操作呢由于数组的存在,数据往往可能不是一维的。这时,我们可以使用二维的线程块例如,在很多的图像算法中使用了8x8线程块来处理像素。

 


  • 3
    点赞
  • 2
    收藏
    觉得还不错? 一键收藏
  • 1
    评论
以下是用CUDA实现对大规模矩阵进行求和的代码,其中使用了二维的线程块和线程网格: ``` #include <stdio.h> #define N 1024 // 矩阵大小 // CUDA 核函数,对二维矩阵进行求和 __global__ void sumMatrix(float *d_matrix, float *d_sum, int width) { int tx = threadIdx.x; // 线程在块内的 x 坐标 int ty = threadIdx.y; // 线程在块内的 y 坐标 int bx = blockIdx.x; // 块在网格中的 x 坐标 int by = blockIdx.y; // 块在网格中的 y 坐标 // 计算当前线程处理的元素的全局索引 int idx = (by * blockDim.y + ty) * width + bx * blockDim.x + tx; if (idx < N * N) { atomicAdd(d_sum, d_matrix[idx]); // 对求和变量进行原子加操作 } } int main() { float *h_matrix, *d_matrix, *h_sum, *d_sum; int size = N * N * sizeof(float); int blockWidth = 32; dim3 block(blockWidth, blockWidth); // 线程块大小 dim3 grid((N + blockWidth - 1) / blockWidth, (N + blockWidth - 1) / blockWidth); // 线程网格大小 // 分配 Host 端和 Device 端内存 h_matrix = (float*)malloc(size); h_sum = (float*)malloc(sizeof(float)); cudaMalloc(&d_matrix, size); cudaMalloc(&d_sum, sizeof(float)); // 初始化矩阵和求和变量 for (int i = 0; i < N * N; i++) { h_matrix[i] = 1.0f; } *h_sum = 0.0f; // 将矩阵和求和变量从 Host 端复制到 Device 端 cudaMemcpy(d_matrix, h_matrix, size, cudaMemcpyHostToDevice); cudaMemcpy(d_sum, h_sum, sizeof(float), cudaMemcpyHostToDevice); // 调用 CUDA 核函数进行求和操作 sumMatrix<<<grid, block>>>(d_matrix, d_sum, N); // 将求和变量从 Device 端复制到 Host 端 cudaMemcpy(h_sum, d_sum, sizeof(float), cudaMemcpyDeviceToHost); // 输出结果 printf("The sum of the matrix is: %f\n", *h_sum); // 释放内存 free(h_matrix); free(h_sum); cudaFree(d_matrix); cudaFree(d_sum); return 0; } ``` 这里的核函数 `sumMatrix` 使用了二维的线程块和线程网格,每个线程计算矩阵中的一个元素,然后对求和变量进行原子加操作。在主函数中,我们将矩阵和求和变量从 Host 端复制到 Device 端,然后调用核函数进行求和操作,最后将求和变量从 Device 端复制到 Host 端并输出结果。
评论 1
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值