CUDA——线程束分化

本文是阅读《CUDA C 编程权威指南》所做的笔记

1. 线程束分化

线程束是SM中基本的执行单元。当一个线程块的网格被启动后，网格中的线程块分布在SM中。一旦线程块被调度到一个SM上，线程块中的线程会被进一步划分为线程束。一个线程束由32个连续的线程组成，在一个线程束中，所有的线程按照单指令多线程（SIMT）方式执行；也就是说，所有线程都执行相同的指令，每个线程在私有数据上进行操作。下图展示了线程块的逻辑视图和硬件视图之间的关系。

GPU是相对简单的设备，它没有复杂的分支预测机制。一个线程束中的所有线程在同一周期中必须执行相同的指令，如果一个线程执行一条指令，那么线程束中的所有线程都必须执行该指令。如果在同一线程束中的线程使用不同的路径通过同一个应用程序，这可能会产生问题。思考下面语句：

if(cond)
{...
}
else
{...
}

一半的线程束需要执行if语句块中的指令，而另一半需要执行else语句块中的指令。在同一线程束中的线程执行不同的指令，被称为线程束分化。
如果一个线程束中的线程产生分化，线程束将连续执行每一个分支路径，而禁用不执行这一路径的线程。线程束分化会导致性能明显地下降。在前面的例子中可以看到，线程束中并行线程的数量减少了一半：只有16个线程同时活跃地执行，而其他16个被禁用了。条件分支越多，并行性削弱越严重。
注意，线程束分化只发生在同一个线程束中。在不同的线程束中，不同的条件值不会引起线程束分化。

2. 简单的线程束分化代码：

#include <stdio.h>
#include <cuda_runtime.h>
#include "device_launch_parameters.h"
#include <stdlib.h>
#include <time.h>// kernel1
__global__ void mathKernel1(float *c)
{int tid = blockIdx.x * blockDim.x + threadIdx.x;float ia, ib;ia = ib = 0.0f;if (tid % 2 == 0) // 偶数线程{ia = 100.0f;}else // 奇数线程{ib = 200.0f;}c[tid] = ia + ib;
}// kernel2
__global__ void mathKernel2(float *c)
{int tid = blockIdx.x * blockDim.x + threadIdx.x;float ia, ib;ia = ib = 0.0f;// 分支粒度是线程束(warpsize = 32)大小的倍数if (tid / 32 % 2 == 0) // 偶数{ia = 100.0f;}else // 奇数{ib = 200.0f;}c[tid] = ia + ib;
}// Kernel3
__global__ void mathKernel3(float *c)
{int tid = blockIdx.x * blockDim.x + threadIdx.x;float ia, ib;ia = ib = 0.0f;bool ipred = (tid % 2 == 0);if (ipred){ia = 100.0f;}if (!ipred){ib = 200.0f;}c[tid] = ia + ib;
}// Kernel4
__global__ void mathKernel4(float *c)
{int tid = blockIdx.x * blockDim.x + threadIdx.x;float ia, ib;ia = ib = 0.0f;int itid = tid >> 5;if (itid & 0x01 == 0) // 偶数{ia = 100.0f;}else // 奇数{ib = 200.0f;}c[tid] = ia + ib;
}// Kernelwarm
__global__ void warmingup(float *c)
{int tid = blockIdx.x * blockDim.x + threadIdx.x;float ia, ib;ia = ib = 0.0f;if ((tid / warpSize) % 2 == 0){ia = 100.0f;}else{ib = 200.0f;}c[tid] = ia + ib;
}// 主函数
int main(int argc, char **argv)
{// set up deviceint dev = 0;cudaDeviceProp deviceProp;cudaGetDeviceProperties(&deviceProp, dev);printf("%s using Device %d: %s\n", argv[0], dev, deviceProp.name);int a, bc, d;// set up data sizeint size = 1 << 25;int blocksize = 256;printf("Data size %d ", size);// set up execution configurationdim3 block(blocksize, 1);dim3 grid((size + block.x - 1) / block.x, 1);printf("Execution Configure (block %d grid %d)\n", block.x, grid.x);// allocate gpu memoryfloat *d_C;size_t nBytes = size * sizeof(float);cudaMalloc((float**)&d_C, nBytes);clock_t iStart, iElaps;// run a warmup kernel to remove overheadcudaDeviceSynchronize();iStart = clock();warmingup << <grid, block >> > (d_C);cudaDeviceSynchronize();iElaps = clock();float time = (float)(iElaps - iStart) / CLOCKS_PER_SEC;printf("warmup      <<< %4d %4d >>> elapsed %f sec \n", grid.x, block.x, time);cudaGetLastError();// run kernel 1iStart = clock();mathKernel1 << <grid, block >> > (d_C);cudaDeviceSynchronize();iElaps = clock();float time1 = (float)(iElaps - iStart) / CLOCKS_PER_SEC;printf("mathKernel1 <<< %4d %4d >>> elapsed %f sec \n", grid.x, block.x, time1);cudaGetLastError();// run kernel 2iStart = clock();mathKernel2 << <grid, block >> > (d_C);cudaDeviceSynchronize();iElaps = clock();float time2 = (float)(iElaps - iStart) / CLOCKS_PER_SEC;printf("mathKernel2 <<< %4d %4d >>> elapsed %f sec \n", grid.x, block.x, time2);cudaGetLastError();// run kernel 3iStart = clock();mathKernel3 << <grid, block >> > (d_C);cudaDeviceSynchronize();iElaps = clock();float time3 = (float)(iElaps - iStart) / CLOCKS_PER_SEC;printf("mathKernel3 <<< %4d %4d >>> elapsed %f sec \n", grid.x, block.x, time3);cudaGetLastError();// run kernel 4iStart = clock();mathKernel4 << <grid, block >> > (d_C);cudaDeviceSynchronize();iElaps = clock();float time4 = (float)(iElaps - iStart) / CLOCKS_PER_SEC;printf("mathKernel4 <<< %4d %4d >>> elapsed %f sec \n", grid.x, block.x, time4);cudaGetLastError();// free gpu memory and reset divececudaFree(d_C);cudaDeviceReset();return EXIT_SUCCESS;
}

3. 运行结果

CUDA——线程束分化相关推荐

cuda线程束原语 __shfl_xor、__shfl、__shfl_up()、__shfl_down()
在CC3.0以上,支持了shuffle指令,允许thread直接读其他thread的寄存器值,只要两个thread在同一个warp中,这种比通过shared Memory进行thread间的通讯效果 ...
【并行计算-CUDA开发】CUDA线程、线程块、线程束、流多处理器、流处理器、网格概念的深入理解
GPU的硬件结构,也不是具体的硬件结构,就是与CUDA相关的几个概念:thread,block,grid,warp,sp,sm. sp: 最基本的处理单元,streaming processor 最 ...
深入理解CUDA线程层次以及关于设置线程数的思考
深入理解CUDA线程层次以及关于设置线程数的思考 2015-09-16 08:45 215人阅读评论(0) 收藏举报分类: cuda(24) GPU线程以网格(grid)的方式组织,而每个网格中 ...
GPU（CUDA）学习日记（十一）------ 深入理解CUDA线程层次以及关于设置线程数的思考
GPU(CUDA)学习日记(十一)------ 深入理解CUDA线程层次以及关于设置线程数的思考标签: cuda存储线程结构网格 2012-12-07 16:30 6298人阅读评论(4)收藏举 ...
最优的cuda线程配置
最优的cuda线程配置 1 每个SM上面失少要有192个激活线程,寄存器写后读的数据依赖才能被掩盖 2 将寄存器的bank冲突降到最低,应尽量使每个block含有的线程数是64的倍数 3 ...
(CUDA 编程1).CUDA 线程执行模型分析（一）招兵 ------ GPU的革命
(CUDA 编程1).CUDA 线程执行模型分析(一)招兵 ------ GPU的革命作者:赵开勇来源:http://www.hpctech.com/2009/0818/198.html 序:或许 ...
CUDA线程、线程块、线程束、流多处理器、流处理器、网格概念的深入理解
一.与CUDA相关的几个概念:thread,block,grid,warp,sp,sm. sp: 最基本的处理单元,streaming processor 最后具体的指令和任务都是在sp上处理的.G ...
CUDA——线程配置
前言线程的组织形式对程序的性能影响是至关重要的,本篇博文主要以下面一种情况来介绍线程组织形式: 2D grid 2D block 线程索引矩阵在memory中是row-major线性存储的: 在k ...
CUDA 线程索引公式
一.线程(thread).线程块(block).线程格(grid) 二.threadIdx.blockIdx.blockDim 和 gridDim 索引分析 1.先找到当前线程位于线程格中的哪一个线程 ...

CUDA——线程束分化

CUDA——线程束分化相关推荐

最新文章

热门文章