CUDA 的 Threading:Block 和 Grid 设定

本文详细解析了CUDA架构中的SM、TPC、SP及其工作原理,包括warp的概念、Block与Thread的对应关系,以及如何根据硬件特性优化Block大小以提升性能。重点讨论了在Compute Capability 1.0/1.1与1.2硬件平台下,如何合理设置Block大小以最大化资源利用率。

硬件基本架构

  实际上在 nVidia 的 GPU 里,最基本的处理单元是所谓的 SP(Streaming Processor),而一颗 nVidia 的 GPU 里,会有非常多的 SP 可以同时做计算;而数个 SP 会在附加一些其他单元,一起组成一个 SM(Streaming Multiprocessor)。几个 SM 则会在组成所谓的 TPC(Texture Processing Clusters)。

  在 G80/G92 的架构下,总共会有 128 个 SP,以 8 个 SP 为一组,组成 16 个 SM,再以两个 SM 为一个 TPC,共分成 8 个 TPC 来运作。而在新一代的 GT200 里,SP 则是增加到 240 个,还是以 8 个 SP 组成一个 SM,但是改成以 3 个 SM 组成一个 TPC,共 10 组 TPC。下面则是提供了两种不同表示方式的示意图。(可参考《NVIDIA G92终极状态!!》、《NVIDIA D10U绘图核心》)

硬件基本架构&对应到 CUDA

 

硬件基本架构&对应到 CUDA

对应到 CUDA

  而在 CUDA 中,应该是没有 TPC 的那一层架构,而是只要根据 GPU 的 SM、SP 的数量和资源来调整就可以了。

  如果把 CUDA 的 Grid - Block - Thread 架构对应到实际的硬件上的话,会类似对应成 GPU - Streaming Multiprocessor - Streaming Processor;一整个 Grid 会直接丢给 GPU 来执行,而 Block 大致就是对应到 SM,thread 则大致对应到 SP。当然,这个讲法并不是很精确,只是一个简单的比喻而已。

SM 中的 Warp 和 Block

SM 中的 Warp 和 Block

  CUDA 的 device 实际在执行的时候,会以 Block 为单位,把一个个的 block 分配给 SM 进行运算;而 block 中的 thread,又会以「warp」为单位,把 thread 来做分组计算。目前 CUDA warp 大小都是 32,也就是 32 个 thread 会被群组成一个 warp 来一起执行;同一个 warp 里的thread,会以不同的数据,执行同样的指令。此外,在 Compute Capability 1.2 的硬件中,还加入了 warp vote 的功能,可以快速的进行 warp 内的简单统计。

  基本上 warp 分组的动作是由 SM 自动进行的,会以连续的方式来做分组。比如说如果有一个 block 里有 128 个 thread 的话,就会被分成四组 warp,第 0-31 个 thread 会是 warp 132-63 是 warp 264-95 是 warp 396-127 是 warp 4

  而如果 block 里面的 thread 数量不是 32 的倍数,那他会把剩下的 thread 独立成一个 warp;比如说 thread 数目是 66 的话,就会有三个 warp0-3132-6364-65。由于最后一个 warp里只剩下两个 thread,所以其实在计算时,就相当于浪费了 30 个 thread 的计算能力;这点是在设定 block 中 thread 数量一定要注意的事!

  一个 SM 一次只会执行一个 block 里的一个 warp,但是 SM 不见得会一次就把这个 warp 的所有指令都执行完;当遇到正在执行的 warp 需要等待的时候(例如存取 global memory 就会要等好一段时间),就切换到别的 warp 来继续做运算,藉此避免为了等待而浪费时间。所以理论上效率最好的状况,就是在 SM 中有够多的 warp 可以切换,让在执行的时候,不会有「所有warp 都要等待」的情形发生;因为当所有的 warp 都要等待时,就会变成 SM 无事可做的状况了~

  下图就是一个 warp 排程的例子。一开始是先执行 thread block 1 的 warp1,而当他执行到第六行指令的时候,因为需要等待,所以就会先切到 thread block 1的 warp2 来执行;一直等到存取结束,且刚好有一个 warp 结束时,才继续执行 TB1 warp1 的第七行指令。

SM 中的 Warp 和 Block

  实际上,warp 也是 CUDA 中,每一个 SM 执行的最小单位;如果 GPU 有 16 组 SM 的话,也就代表他真正在执行的 thread 数目会是 32*16 个。不过由于 CUDA 是要透过 warp 的切换来隐藏 thread 的延迟、等待,来达到大量平行化的目的,所以会用所谓的 active thread 这个名词来代表一个 SM 里同时可以处理的 thread 数目。

  而在 block 的方面,一个 SM 可以同时处理多个 thread block,当其中有 block 的所有thread 都处理完后,他就会再去找其他还没处理的 block 来处理。假设有 16 个 SM64 block、每个 SM 可以同时处理三个 block 的话,那一开始执行时,device 就会同时处理 48 block;而剩下的 16 个 block 则会等 SM 有处理完 block 后,再进到 SM 中处理,直到所有 block都处理结束。

  为一个多处理器(SM)指定了一个或多个要执行的线程块(thread block)时,它会将其分成warp块,并由SIMT单元进行调度。将块分割为warp的方法总是相同的,每个warp都包含连续的线程,递增线程索引,第一个warp中包含全局线程过索引0-31。每发出一条指令时,SIMT单元都会选择一个已准备好执行的warp块,并将指令发送到该warp块的活动线程。Warp块每次执行一条通用指令,因此在warp块的全部32个线程执行同一条路径时,可达到最高效率。如果一个warp块的线程通过独立于数据的条件分支而分散,warp块将连续执行所使用的各分支路径,而禁用未在此路径上的线程,完成所有路径时,线程重新汇聚到同一执行路径下,其执行时间为各时间总和。分支仅在warp块内出现,不同的warp块总是独立执行的--无论它们执行的是通用的代码路径还是彼此无关的代码路径。

建议的数值?

  在 Compute Capability 1.0/1.1 中,每个 SM 最多可以同时管理 768 个 thread(768 active threads)或 个 block(8 active blocks);而每一个 warp 的大小,则是 32 个 thread,也就是一个SM 最多可以有 768 / 32 = 24 个 warp(24 active warps)。到了 Compute Capability 1.2 的话,则是 active warp 则是变为 32,所以 active thread 也增加到 1024

  在这里,先以 Compute Capability 1.0/1.1 的数字来做计算。根据上面的数据,如果一个block 里有 128 个 thread 的话,那一个 SM 可以容纳 个 block;如果一个 block 有 256 个 thread的话,那 SM 就只能容纳 个 block。不过如果一个 block 只有 64 个 thread 的话,SM 可以容纳的 block 不会是 12 个,而是他本身的数量限制的 个。

  因此在 Compute Capability 1.0/1.1 的硬件上,决定 block 大小的时候,最好让里面的thread 数目是 warp 数量(32)的倍数(可以的话,是 64 的倍数会更好);而在一个 SM 里,最好也要同时存在复数个 block。如果再希望能满足最多 24 个 warp 的情形下,block 里的 thread 数目似乎会是 96(一个 SM 中有 个 block)128(一个 SM 中有 个 block)192(一个 SM 中有 block)256(一个 SM 中有 个 block) 这些数字了~

  而官方的建议则是一个 block 里至少要有 64 个 thread192 或 256 个也是通常比较合适的数字(请参考 Programming Guide)

  但是是否这些数字就是最合适的呢?其实也不尽然。因为实际上,一个 SM 可以允许的block 数量,还要另外考虑到他所用到 SM 的资源:shared memoryregisters 等。在 G80 中,每个 SM 有 16KB 的 shared memory 和 8192 个 register。而在同一个 SM 里的 block thread,则要共享这些资源;如果资源不够多个 block 使用的话,那 CUDA 就会减少 Block 的量,来让资源够用。在这种情形下,也会因此让 SM 的 thread 数量变少,而不到最多的 768个。

  比如说如果一个 thread 要用到 16 个 register 的话(在 kernel 中宣告的变量),那一个 SM 8192 个 register 实际上只能让 512 个 thread 来使用;而如果一个 thread 要用 32 个 register,那一个 SM 就只能有 256 个 thread 了~而 shared memory 由于是 thread block 共享的,因此变成是要看一个 block 要用多少的 shread memory、一个 SM 的 16KB 能分给多少个 block 了。

  所以虽然说当一个 SM 里的 thread 越多时,越能隐藏 latency,但是也会让每个 thread 能使用的资源更少。因此,这点也就是在优化时要做取舍的了。

<think>好的,我现在需要帮助用户了解如何在jetson-inference的Python API中使用CUDA。首先,我得回顾一下用户之前的问题之前的回答,确保信息一致。用户之前询问了CUDA在jetson-inference中的应用与配置,现在他想知道具体在Python API中的使用方法。 用户提供的引用内容中有关于安装PyTorchjetson-inference的步骤,特别是环境配置编译过程。引用[1]提到了安装PyTorch的脚本,引用[2]详细说明了jetson-inference的安装步骤,包括编译设置环境变量,引用[3]则涉及Python包的安装,比如torch的whl文件。这些信息可能对回答有帮助,但用户现在关注的是Python API中CUDA的具体使用。 首先,我需要确定jetson-inference的Python API如何集成CUDA。根据之前的回答,jetson.utilsjetson.inference模块已经封装了CUDA操作,用户不需要直接写CUDA内核。但用户可能需要了解如何利用这些封装好的功能,比如数据加载到显存、使用预训练模型进行推理,以及可能的异步处理。 接下来,我需要分步骤解释如何在Python中使用这些API。例如,使用jetson.utils加载图像到CUDA显存,使用jetson.inference中的模型进行推理,处理视频流时的CUDA流控制,以及自定义CUDA扩展的方法。 此外,用户可能需要了解内存管理性能优化的技巧,比如显存的预分配混合精度支持。同时,常见问题如CUDA版本冲突或显存不足的解决方法也需要提及,引用之前的引用[4]中的内容。 需要确保回答结构清晰,分点说明,并提供代码示例。例如,展示如何加载图像、进行目标检测、处理视频流,以及如何编译调用自定义CUDA内核。还要注意环境变量的配置,如LD_LIBRARY_PATH,这在引用[2][3]中都有涉及。 最后,生成的相关问题需要围绕CUDA在jetson-inference中的深入应用,比如多GPU支持、性能监控、与cuDNN的集成等。确保问题能够引导用户进一步探索相关主题。</think>### CUDA 在 jetson-inference Python API 中的具体使用方法 jetson-inference 的 Python API 通过封装 CUDA 操作实现了高效的数据传输与计算,开发者可通过以下核心模块直接调用 CUDA 功能: --- #### 一、CUDA 显存管理与数据传输 1. **图像/视频数据加载到显存** - 使用 `jetson.utils` 模块直接将图像数据存入 CUDA 显存(避免 CPU-GPU 显存拷贝开销): ```python from jetson_utils import cudaFromNumpy, cudaToNumpy import cv2 # OpenCV 读取图像到 CPU 内存 cpu_img = cv2.imread("test.jpg") # 将 CPU 内存数据直接映射到 CUDA 显存(零拷贝) cuda_img = cudaFromNumpy(cpu_img) # 返回 CUDA 指针[^1] ``` 2. **显存与 CPU 内存互转** ```python # CUDA 显存数据转回 CPU 内存(用于后处理) numpy_array = cudaToNumpy(cuda_img) # 需显式转换[^3] ``` --- #### 二、深度学习模型推理的 CUDA 加速 1. **加载预训练模型** - 模型权重计算图自动加载至 CUDA 显存: ```python from jetson_inference import detectNet # 加载 SSD-Mobilenet 目标检测模型(自动启用 CUDA) net = detectNet("ssd-mobilenet-v2", threshold=0.5) ``` 2. **执行 CUDA 加速的推理** ```python # 使用 CUDA 显存中的图像直接推理 detections = net.Detect(cuda_img) # 输入为 CUDA 指针[^2] # 输出结果包含 CUDA 坐标(无需转回 CPU) for detection in detections: print(f"Class: {detection.ClassID}, Confidence: {detection.Confidence}") print(f"Bounding Box (CUDA坐标): {detection.Left}, {detection.Top}, {detection.Right}, {detection.Bottom}") ``` --- #### 三、视频流处理的 CUDA 优化 1. **摄像头/视频输入流的 CUDA 加速** - 使用 `videoSource` 直接捕获到显存: ```python from jetson_utils import videoSource # 创建 CSI 摄像头输入流(默认启用 CUDA 解码) stream = videoSource("csi://0", argv=["--input-width=640", "--input-height=480"]) while True: cuda_frame = stream.Capture() # 返回 CUDA 显存中的帧 if cuda_frame is None: break detections = net.Detect(cuda_frame) # 直接在显存处理 ``` 2. **异步 CUDA 流控制** - 通过多 CUDA 流提升吞吐量(需自行封装): ```python import threading # 创建多个 CUDA 流(需通过 C++ 扩展实现,Python API 未直接暴露) # 推荐使用多线程 + 预分配显存缓冲区实现类似效果 ``` --- #### 四、自定义 CUDA 操作的集成 若需扩展自定义 CUDA 内核,需通过以下步骤: 1. **编写 CUDA 内核函数** ```cpp // custom_kernel.cu __global__ void rgb2gray(float* input, float* output, int width, int height) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (x < width && y < height) { int idx = y * width * 3 + x * 3; output[y * width + x] = 0.299f * input[idx] + 0.587f * input[idx+1] + 0.114f * input[idx+2]; } } ``` 2. **编译为 Python 可调用库** ```bash nvcc -Xcompiler -fPIC -shared -o libcustom.so custom_kernel.cu ``` 3. **在 Python 中调用 CUDA 内核** ```python from ctypes import cdll, c_void_p, c_int import numpy as np lib = cdll.LoadLibrary("./libcustom.so") lib.rgb2gray.argtypes = [c_void_p, c_void_p, c_int, c_int] # 输入为 CUDA 显存指针(需通过 jetson.utils 获取) input_ptr = cuda_img.ptr # 假设 cuda_img 为 RGB 图像 output_ptr = cudaAllocMapped(width, height, format='gray8').ptr # 启动 CUDA 内核(block=(16,16), grid=(40,30) 针对 640x480) lib.rgb2gray(input_ptr, output_ptr, 640, 480) ``` --- #### 五、关键注意事项 1. **显存生命周期管理** - CUDA 显存对象在 Python 中通过引用计数自动释放,但需避免循环引用。 - 显式释放显存: ```python del cuda_img # 触发 CUDA 内存回收 ``` 2. **混合精度支持** - 启用 FP16 推理(需模型支持): ```python net = detectNet("ssd-mobilenet-v2", threshold=0.5, config="--fp16") ``` 3. **性能监控工具** - 使用 `jetson_stats` 监控 CUDA 显存与核心利用率: ```bash sudo jtop # 查看 GPU 使用率、显存占用等[^4] ``` --- ### 典型代码流程(目标检测场景) ```python from jetson_utils import videoSource, videoOutput, cudaFromNumpy from jetson_inference import detectNet # 初始化输入输出流 input = videoSource("csi://0") output = videoOutput("display://0") net = detectNet("ssd-mobilenet-v2") while True: img = input.Capture() # 获取 CUDA 显存中的帧 if img is None: break detections = net.Detect(img) # CUDA 加速推理 output.Render(img) # 渲染结果到显示屏(CUDA 加速) output.SetStatus("FPS: {:.1f}".format(net.GetNetworkFPS())) ``` --- ### 相关问题 1. 如何在 jetson-inference 中实现多模型并行推理的 CUDA 资源分配? 2. 如何通过 Python API 监控 CUDA 内核的执行时间? 3. jetson-inference 是否支持 CUDA 与 TensorRT 的混合精度动态切换? 4. 如何利用 CUDA 流实现视频采集与推理的流水线并行? 如需具体场景的代码调试或性能分析,请提供硬件配置(如 Jetson Nano 或 Xavier)任务类型(分类/检测/分割)。
评论
成就一亿技术人!
拼手气红包6.0元
还能输入1000个字符
 
红包 添加红包
表情包 插入表情
 条评论被折叠 查看
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

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

抵扣说明:

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

余额充值