基础

每个GPU有多个SM（streaming multiprocessor）
当启动一个grid时,它的block会被分配给多个SM上执行,一个block一旦被调度到一个SM上,则这个block只会在那个SM上执行
多个block可以被分配到一个SM上执行
没32个线程未一组,被称为线程束(warp)
block里的thread逻辑上可以并行运行,单并不是所有的thread可以同时在物理层面执行,既block中,不同的thread可能会以不同的速度前进.

warp(线程束)和block(线程块)

block被分配到一个SM时,会被划分为多个warp
一个warp由32个连续的core组成
block可以被配置为一维,二维或三维的,但是物理上都被组织成了一维
1. 对于一维block,唯一threadid被存在threadIdx.x中,并且,threadIdx.x拥有连续值的线程被分组到同一个warp中
  假设有128个thread的block,是被分配给了4个warp
2. 二维的block,每个thread的id可以通过threadIdx和blockDim来计算:
  threadIdx.y * blockDim.x + threadIdx.x
3. 三维的block
  thradIdx.z * blockDim.y * blockDim.x + threadIdx.y * blockDim.x + threadIdx.x
若果block的大小不是warp大小的整数倍,就会造成资源浪费,如启动80个thread,会分配3个warp,共96个core,多余的16个core仍消耗sm资源,如寄存器

线程束分化

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

假设32个thread,16个执行true,16个执行false,这样在同一warp中执行不同的指令就称为warp分化
warp并行线程数量减少了一半,16个线程同时活跃执行,其余16个被禁用了(16个执行if时,另外16个等待,16个执行else时,另外16个等待)
应避免同一warp中有不同的执行路径
不同的if-then-else分支会连续执行
调整分支力度以适应线程束大小的倍数

__global__ void mathKernel1(float *c){int tid = blockIdx.x * blockDim.x + threadIdx.x;float ia, ib;ia = ib = 0.0f;if (tid % 2 == 0){//奇数,偶数thread分开ia = 100.0f;}else{ib = 200.0f;}c[tid] = ia + ib;
}__global__ void mathKernel2(float *c){int tid = blockIdx.x * blockDim.x + threadIdx.x;float ia, ib;ia = ib = 0.0f;if ((tid / warpSize) % 2 == 0){//奇数,偶数warp分开ia = 100.0f;}else{ib = 200.0f;}c[tid] = ia + ib;
}__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){//使用两个if语句,而不是if-else来减少if-else的分化,只有在if的条件为0时,编译器才会下达指令ia = 100.0f;}if (!ipred){ib = 200.0f;}c[tid] = ia + ib;
}__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;
}

占用率

#include <stdio.h>
#include <cuda_runtime.h>int main(int argc, char *argv[])
{int iDev = 0;cudaDeviceProp iProp;cudaGetDeviceProperties(&iProp, iDev);printf("Device %d: %s\n", iDev, iProp.name);printf("  Number of multiprocessors:                     %d\n",iProp.multiProcessorCount);printf("  Total amount of constant memory:               %4.2f KB\n",iProp.totalConstMem / 1024.0);printf("  Total amount of shared memory per block:       %4.2f KB\n",iProp.sharedMemPerBlock / 1024.0);printf("  Total number of registers available per block: %d\n",iProp.regsPerBlock);printf("  Warp size:                                     %d\n",iProp.warpSize);printf("  Maximum number of threads per block:           %d\n",iProp.maxThreadsPerBlock);printf("  Maximum number of threads per multiprocessor:  %d\n",iProp.maxThreadsPerMultiProcessor);printf("  Maximum number of warps per multiprocessor:    %d\n",iProp.maxThreadsPerMultiProcessor / 32);return EXIT_SUCCESS;
}

grid和block大小准则

每个block的thread数时warp(32)的倍数
每个block至少有128或256个thread(避免block太小)
根据内核资源调整block大小
block数量要多与SM的数量

同步

系统级:等待host和device完成
lock级:每个block所有的thread到大同一点
cudaError_t cudaDeviceSynchronize(void)这色host,等待device返回
__device__ void __syncthreads(void);同一block每个thread不许等待,知道所有thread都到大同一点

避免分支分化

执行满足交换律和结合律的运算,被称为规约问题,并行归约是并行执行

相邻配对

// Neighbored Pair Implementation with divergence
__global__ void reduceNeighbored (int *g_idata, int *g_odata, unsigned int n){unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;if (idx >= n) return;// boundary checkunsigned int tid = threadIdx.x;int *idata = g_idata + blockIdx.x * blockDim.x;//将全局内存转换为block内存// in-place reduction in global memoryfor (int stride = 1; stride < blockDim.x; stride *= 2){//每次归约步长变为2倍if ((tid % (2 * stride)) == 0){//相邻两个数相加(这里会导致warp分化)idata[tid] += idata[tid + stride];}__syncthreads();// 等待block中的所有线程结束}if (tid == 0) g_odata[blockIdx.x] = idata[0];// 将第0个线程的block内存写入全局内存中
}

优化(没看懂)

__global__ void reduceNeighboredLess (int *g_idata, int *g_odata, unsigned int n){unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;if(idx >= n) return;unsigned int tid = threadIdx.x;int *idata = g_idata + blockIdx.x * blockDim.x;for (int stride = 1; stride < blockDim.x; stride *= 2){int index = 2 * stride * tid;if (index < blockDim.x){idata[index] += idata[index + stride];}__syncthreads();}if (tid == 0) g_odata[blockIdx.x] = idata[0];
}

int index = 2 * stride * tid;``index < blockDim.x 对512个线程block来说,前8个warp(32*8=256)完成第一轮归约,剩下8个线程什么也不做.第二轮里,前4个warp完成归约,剩下的12个线程束什么也不做.当thread总数,小于warp时,分化就会出现.

交错配对

strike跨度是block大小的一半,每次迭代归约减少一半
与相邻归约相比,交错归约的工作线程没有变化,但是线程在全局内存中的加载位置是不同的

__global__ void reduceInterleaved (int *g_idata, int *g_odata, unsigned int n)
{unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;if(idx >= n) return;unsigned int tid = threadIdx.x;int *idata = g_idata + blockIdx.x * blockDim.x;for (int stride = blockDim.x / 2; stride > 0; stride >>= 1){if (tid < stride){idata[tid] += idata[tid + stride];}__syncthreads();}if (tid == 0) g_odata[blockIdx.x] = idata[0];
}

循环展开

for (int i = 0; i< 100;i++){//循环检查100次a[i] = b[i] + c[i];
}for (int i = 0; i< 100;i+=2){//循环检查50次a[i] = b[i] + c[i];a[i+1] = b[i+1] + c[i+1];
}

__global__ void reduceUnrolling2 (int *g_idata, int *g_odata, unsigned int n)
{// set thread IDunsigned int tid = threadIdx.x;unsigned int idx = blockIdx.x * blockDim.x * 2 + threadIdx.x;int *idata = g_idata + blockIdx.x * blockDim.x * 2;if (idx + blockDim.x < n) g_idata[idx] += g_idata[idx + blockDim.x];__syncthreads();// in-place reduction in global memoryfor (int stride = blockDim.x / 2; stride > 0; stride >>= 1){if (tid < stride){idata[tid] += idata[tid + stride];}__syncthreads();}if (tid == 0) g_odata[blockIdx.x] = idata[0];
}

展开线程归约

__syncthreads用于块内同步,确保thread进入下一轮之前,每一轮的所有thread已经将局部结果写入全局内存中

__global__ void reduceUnrolling8 (int *g_idata, int *g_odata, unsigned int n){unsigned int tid = threadIdx.x;unsigned int idx = blockIdx.x * blockDim.x * 8 + threadIdx.x;int *idata = g_idata + blockIdx.x * blockDim.x * 8;if (idx + 7 * blockDim.x < n){int a1 = g_idata[idx];int a2 = g_idata[idx + blockDim.x];int a3 = g_idata[idx + 2 * blockDim.x];int a4 = g_idata[idx + 3 * blockDim.x];int b1 = g_idata[idx + 4 * blockDim.x];int b2 = g_idata[idx + 5 * blockDim.x];int b3 = g_idata[idx + 6 * blockDim.x];int b4 = g_idata[idx + 7 * blockDim.x];g_idata[idx] = a1 + a2 + a3 + a4 + b1 + b2 + b3 + b4;}__syncthreads();for (int stride = blockDim.x / 2; stride > 0; stride >>= 1){if (tid < stride){idata[tid] += idata[tid + stride];}__syncthreads();}if (tid == 0) g_odata[blockIdx.x] = idata[0];
}

完全展开归约

模板函数归约

动态并行

嵌套执行

内核执行分为父母和孩子,只有在所有的子网格都完成后,父母才会完成

#include <stdio.h>
#include <cuda_runtime.h>__global__ void nestedHelloWorld(int const iSize, int iDepth)
{int tid = threadIdx.x;printf("Recursion=%d: Hello World from thread %d block %d\n", iDepth, tid, blockIdx.x);if (iSize == 1) return;int nthreads = iSize >> 1; //减少一半if(tid == 0 && nthreads > 0){nestedHelloWorld<<<1, nthreads>>>(nthreads, ++iDepth);printf("-------> nested execution depth: %d\n", iDepth);}
}int main(int argc, char **argv)
{int size = 8;int blocksize = 8;   // initial block sizeint igrid = 1;if(argc > 1){igrid = atoi(argv[1]);size = igrid * blocksize;}dim3 block (blocksize, 1);dim3 grid  ((size + block.x - 1) / block.x, 1);printf("%s Execution Configuration: grid %d block %d\n", argv[0], grid.x,block.x);nestedHelloWorld<<<grid, block>>>(block.x, 0);cudaDeviceReset();return 0;
}

嵌套归约

CUDA C编程权威指南第三章 CUDA执行模型相关推荐

CUDA C编程权威指南第六章流和并发
流 cuda流流分为两种类型:1)隐式声明流(空流),2)显式声明流(非空流) cudaError_t cudaMemcpyAsync(void* dst,const void* src, size ...
CUDA C编程权威指南第五章共享内存和常量内存
共享内存是较小的片上内存,具有较低的延迟(相比全局,低20~30倍),提供更高的带宽(相比全局,10倍) block通信用于全局内存数据的缓存 __shared__来申请共享变量,如果共享内存大小在 ...
CUDA C编程权威指南第四章全局内存
基础一般内存的设计:寄存器->缓存->主存->磁盘存储器 GPU内存设计修饰符变量名存储器作用域生命周期 float var 寄存器线程线程 float var[10 ...
CUDA C编程权威指南第七章调整指令级原语
将程序分为两类:IO密集型和计算密集型 double value = in1 * in2 + in3 乘法后紧跟加法的模式被称为乘加法,或者MAD 简单的编译器会将一个MAD指令转换为:一个乘法指令和 ...
《CUDA C编程权威指南》——1.5节总结
本节书摘来自华章社区<CUDA C编程权威指南>一书中的第1章,第1.5节总结,作者[美] 马克斯·格罗斯曼(Max Grossman) ,更多章节内容可以访问云栖社区"华章社区 ...
《CUDA C编程权威指南》——3.4　避免分支分化
本节书摘来自华章计算机<CUDA C编程权威指南>一书中的第3章,第3.4节,作者 [美] 马克斯·格罗斯曼(Max Grossman),译颜成钢殷建李亮,更多章节内容可以访问云栖社 ...
c cuda 指定gpu_《CUDA C编程权威指南》——1.3　用GPU输出Hello World-阿里云开发者社区...
本节书摘来自华章计算机<CUDA C编程权威指南>一书中的第1章,第1.3节,作者 [美] 马克斯·格罗斯曼(Max Grossman),译颜成钢殷建李亮,更多章节内容可以访问云栖社 ...
《CUDA C编程权威指南》——2.2　给核函数计时
本节书摘来自华章计算机<CUDA C编程权威指南>一书中的第2章,第2.2节,作者 [美] 马克斯·格罗斯曼(Max Grossman),译颜成钢殷建李亮,更多章节内容可以访问云栖社 ...
《CUDA C编程权威指南》——2.4节设备管理
本节书摘来自华章社区<CUDA C编程权威指南>一书中的第2章,第2.4节设备管理,作者[美] 马克斯·格罗斯曼(Max Grossman) ,更多章节内容可以访问云栖社区"华章 ...

CUDA C编程权威指南第三章 CUDA执行模型

基础

warp(线程束)和block(线程块)

线程束分化

占用率

grid和block大小准则

同步

避免分支分化

相邻配对

优化(没看懂)

交错配对

循环展开

展开线程归约

完全展开归约

模板函数归约

动态并行

嵌套执行

嵌套归约

CUDA C编程权威指南第三章 CUDA执行模型相关推荐

最新文章

热门文章

CUDA C编程权威指南 第三章 CUDA执行模型

基础

warp(线程束)和block(线程块)

线程束分化

占用率

grid和block大小准则

同步

避免分支分化

相邻配对

优化(没看懂)

交错配对

循环展开

展开线程归约

完全展开归约

模板函数归约

动态并行

嵌套执行

嵌套归约

CUDA C编程权威指南 第三章 CUDA执行模型相关推荐

最新文章

热门文章

CUDA C编程权威指南第三章 CUDA执行模型

CUDA C编程权威指南第三章 CUDA执行模型相关推荐