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

最新推荐文章于 2024-06-15 12:06:58 发布
最新推荐文章于 2024-06-15 12:06:58 发布
阅读量423 收藏 2
点赞数
分类专栏： cuda 文章标签： cuda
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.youkuaiyun.com/claroja/article/details/103817180
版权
本文深入探讨GPU并行计算的基本概念，包括streaming multiprocessor（SM）的作用，block和warp的组织方式，以及线程束分化的影响。通过具体示例解释了如何优化线程配置，避免资源浪费，提高计算效率。

摘要生成于 C知道 ，由 DeepSeek-R1 满血版支持， 前往体验 >

基础

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

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

  1. block被分配到一个SM时,会被划分为多个warp

  2. 一个warp由32个连续的core组成

  3. 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

  4. 若果block的大小不是warp大小的整数倍,就会造成资源浪费,如启动80个thread,会分配3个warp,共96个core,多余的16个core仍消耗sm资源,如寄存器

线程束分化

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

在这里插入图片描述

__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大小准则

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

同步

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

避免分支分化

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

相邻配对

在这里插入图片描述

// 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 check

    unsigned int tid = threadIdx.x;
    int *idata = g_idata + blockIdx.x * blockDim.x;//将全局内存转换为block内存
    // in-place reduction in global memory
    for (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时,分化就会出现.

交错配对

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

在这里插入图片描述

__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 ID
    unsigned 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 memory
    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];
}

在这里插入图片描述

展开线程归约

  1. __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];
}

完全展开归约

模板函数归约

动态并行

嵌套执行

  1. 内核执行分为父母和孩子,只有在所有的子网格都完成后,父母才会完成
#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 size
    int 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执行模型
YYRAN_ZZU的博客
10-29 1026
CUDA执行模型概述说明如何在制定的计算机构上指令。本节会重点介绍指令吞吐量。GPU架构是围绕着一个流式多处理器(SM)的扩展阵列搭建的。GPU中每一个SM都能支持说白格线程并发执行，每个GPU通常有多个SM。CUDA采用单指令多线程(SIMT)架构来管理和执行线程，每32个线程为一组，被称为线程束(warp)。线程束中所有线程同时执行相同的指令。每个线程都有自己的指令地址计数器和状态寄存器，利用自身的数据执行当前的指令。SIMT确保可以编写独立的线程级并行代码、标量线程以及用于协调线程数据并行代码。
CUDA编程模型介绍
最新发布
YYRAN_ZZU的博客
10-26 1692
CUDA是一种通用的并行计算平台和编程模型，在C语言基础扩展，借助于cuda可以像编写C语言一样的实现并行算法，本文主要介绍基础的编程概念和程序入门的第一个cuda程序。cuda编程模型使用C语言扩展成的代码在异构计算系统中执行应用程序，爱一个异构环境中包含多个CPU和GPU，下面介绍两个常用的专用概念词主机：CPU及其内存设备：GPU及其内存设备(kernel)-指代码在GPU上运行的代码如下图所示一般一个典型的CUDA程序实现的流程如下数据从CPU内存拷贝到GPU内存。
cuda c权威编程指南_CUDA C++程序设计模型
weixin_39802814的博客
11-15 486
CUDA C++程序设计模型 本章介绍了CUDA编程模型背后的主要概念，概述了它们在C++中的暴露方式。在编程接口中给出了CUDA C++的广泛描述。 使用的矢量加法示例的完整代码可以在矢量加法CUDA示例中找到。 一. 内核 CUDA C++通过允许程序员定义C++函数，称为内核，扩展了C++，当调用时，用n个不同的CUDA线程并行执行n次，而不是像常规C++函数那样只执行一次。 一个内核使用...
CUDA C编程权威指南Chapter1&2
m0_73535031的博客
03-13 878
基于这些坐标，你可以将部分数据分配给 不同的线程。全局类似于CPU的系统内存，而共享内存类似于CPU的缓存。然而GPU的共享内存可以由CUDA C的内核直接控制。构成，一个线程块包含一组线程，同一线程块内 的线程协作可以通过同步和共享内存来实现。不同块内的线程不能协作。同一网格中的所有线程共享相同 的全局内存空间。（线程格的维度，用每个线程格中的线程数来表示）（线程块的维度，用每个线程块中的线程数来表示）形式，一个线程块会被组织成线程的。一个线程格会被组织成线程块的。（线程块在线程格内的索引）
《CUDA C编程权威指南》——1.1节并行计算
weixin_34417635的博客
05-02 281
本节书摘来自华章社区《CUDA C编程权威指南》一书中的第1章，第1.1节并行计算，作者[美]　马克斯·格罗斯曼（Max Grossman）　，更多章节内容可以访问云栖社区“华章社区”公众号查看 1.1　并行计算在过去的几十年间，人们对并行计算产生了越来越多的兴趣。并行计算的主要目标是提高运算速度。从纯粹的计算视角来看，并行计算可以被定义为计算的一种形式...
《CUDA C编程权威指南》——1.2节异构计算
weixin_33851429的博客
05-02 230
本节书摘来自华章社区《CUDA C编程权威指南》一书中的第1章，第1.2节异构计算，作者[美]　马克斯·格罗斯曼（Max Grossman）　，更多章节内容可以访问云栖社区“华章社区”公众号查看 1.2　异构计算最初，计算机只包含用来运行编程任务的中央处理器（CPU）。近年来，高性能计算领域中的主流计算机不断添加了其他处理元素，其中最主要的就是GPU。G...
CUDA C编程权威指南.pdf
07-08
#### 第3章 CUDA执行模型 **3.1 CUDA执行模型概述** - **3.1.1 GPU架构概述** - GPU架构的发展历程及其特点。 - Fermi架构和Kepler架构的特点。 - **3.1.2 Fermi架构** - Fermi架构的主要改进点和技术特性。 ...
CUDA专家手册 GPU编程权威指南
02-24
《CUDA专家手册：GPU编程权威指南》详细讨论CUDA的硬件和软件，包括CUDA 5.0和开普勒架构的最新特性。每个CUDA开发人员，不论新手还是高手，都可以在这里找到感兴趣的内容并即时上手。新晋的CUDA开发者将理解硬件...
CUDA_C_Programming_Guide.pdf
12-23
针对编程接口的详细阐述，本指南还提供了对CUDA编程模型、内存管理、异步并发执行、多设备系统等关键概念的深入讲解和代码示例。这些内容不仅帮助开发者掌握CUDA的基础知识，更是深入理解如何在实际应用中充分利用...
cudac编程权威指南pdf
08-01
《CUDA编程权威指南》是一本全面介绍了CUDA并行计算框架的经典著作。该书作者是David B. Kirk和Wen-mei Hwu，他们是在NVIDIA公司工作的专家。 该书分为四个部分，共25章，内容详尽且深入。第一部分介绍了CUDA的背景...
CUDA C编程权威指南 第二章 CUDA编程模型
Claroja
01-01 513
CUDA6.0开始 有"统一寻址"(Unified Memory)编程模型,可以用单个指针访问CPU和GPU内存,无须手动拷贝 主机启动内核后,管理权立刻返回给主机(类似启动线程后,不join) C函数 CUDA C函数 malloc cudaMalloc memcpy cudaMemcpy memset cudaMemset free cudaFree cuda...
CUDA C编程权威指南 第3章 CUDA执行模型
weixin_44500921的博客
06-15 826
CUDA编程模型中两个主要的抽象概念：内存层次结构和线程层次结构。它们能够控制大规模并行GPU。GPU架构是围绕一个流式多处理器（SM）的可扩展阵列搭建的。GPU中每一个SM都能支持数百个线程并发执行，每个GPU通常有多个SM，所以在一个GPU上并发执行数千个线程是可能的。当启动一个内核网格时，它的线程块被分布在了可用的SM上来执行。多个线程块可能会被分配到同一个SM上，而且是根据SM资源的可用性进行调度的。
《CUDA C编程权威指南》—读书笔记—3.1CUDA执行模型与硬件
RiddleC的博客
04-18 431
软硬总归是不能分家。上一章简单介绍了CUDA的编程模型，本章学习主要关于CUDA的相关硬件架构、线程执行的本质，以及CUDA在线程级并行的一些优化方法，比如规约问题。话不多数，开始学习。 GPU实际上是由多个流式多处理器构建的处理器阵列（可以这样理解，当然所有SM会由GPU管理，并非完全独立）。先上一张经典Fermi架构的GPU的SM示意图。需要注意的时，以下结构实际上只是 GPU的compute...
cuda c语言编程指南,CUDA C编程权威指南
weixin_39948439的博客
05-18 520
译者序推荐序自序作者简介技术审校者简介前言致谢章基于CUDA的异构并行计算11.1并行计算11.1.1串行编程和并行编程21.1.2并行性31.1.3计算机架构41.2异构计算61.2.1异构架构71.2.2异构计算范例91.2.3CUDA：一种异构计算平台101.3用GPU输出HelloWorld121.4使用CUDAC编程难吗151.5总结161.6习题16第2章CUDA编程模型182.1CU...
getDimension方法的一些坑
zhuhai__yizhi的专栏
02-16 5321
（转载）http://luckyandyzhang.github.io/2015/09/25/getDimension%E6%96%B9%E6%B3%95%E7%9A%84%E4%B8%80%E4%BA%9B%E5%9D%91/ 问题背景 在Android开发中，我们经常会在dimens.xml中定义一些值，比如边距，字体大小等 1 2 3 r
《CUDA C编程权威指南》——2.5节总结
weixin_33907511的博客
05-02 185
本节书摘来自华章社区《CUDA C编程权威指南》一书中的第2章，第2.5节总结，作者[美]　马克斯·格罗斯曼（Max Grossman）　，更多章节内容可以访问云栖社区“华章社区”公众号查看 2.5　总结与C语言中的并行编程相比，CUDA程序中的线程层次结构是其独有的结构。通过一个抽象的两级线程层次结构，CUDA能够控制一个大规模并行环境。通过本章的例子...
CUDA 执行模型
A-Egoist的博客
12-28 479
CUDA 执行模型 CUDA 执行模型能够提供有助于在指令吞吐量和内存访问方面编写高效代码的见解。 1 GPU 架构概述 GPU 架构是围绕一个流式多处理器(SM)的拓展阵列搭建的，通过复制这种结构来实现 GPU 的硬件并行。 如图所示，其中包含了 Fermi SM 关键组件： CUDA Cores Shared Memory / L1 Cache Register File 加载/存储单元(Load / Store Units) 特殊功能单元(Special Function Units) 线程束调度器
CUDA C 编程权威指南 学习笔记：第一章 基于CUDA的异构并行计算
JackZhangNJU的专栏
03-01 504
并行性 这个主要分为任务并行和数据并行，CPU适应任务并行，GPU更假适合数据并行，虽然CPU和GPU结构类似，但是其实很不一样， CPU的核心比较重，用来处理非常复杂的控制逻辑，已优化串行程序执行 GPU的核心比较轻，用于优化具有简单控制逻辑的数据并行任务，注重并行程序的吞吐量 CUDA CUDA提供两层API来管理GPU设备和组织线程 CUDA驱动API CUDA运行时AP...
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

抵扣说明：

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

余额充值