stream式多处理器，块和线程（CUDA）
CUDA核心，stream式多处理器和CUDA块和线程模型之间的关系是什么？
什么被映射到什么和什么是并行的，以及如何？ 什么是更有效率，最大限度的块数或线程数？
我目前的理解是每个多处理器有8个cuda核心。 而且每个cuda核心一次可以执行一个cuda块。 并且该块中的所有线程都在该特定内核中被串行执行。
它是否正确？
CUDA编程指南中详细介绍了线程/块布局。 特别是第四章说：
CUDA架构是围绕可扩展的multithreadingstream处理器（SM）arrays构build的。 当主机CPU上的CUDA程序调用内核网格时，网格块被枚举并分配给具有可用执行能力的多处理器。 一个线程块的线程同时在一个多处理器上执行，多个线程块可以在一个多处理器上同时执行。 当线程块终止时，在空闲的多处理器上启动新的块。
每个SM包含8个CUDA内核，并且在任何时候它们都执行32个线程的单个翘曲 – 因此需要4个时钟周期来为整个翘曲发出单个指令。 您可以假定任何给定warp中的线程都是以lock-step方式执行的，但要跨越warp进行同步，则需要使用__syncthreads() 。
对于GTX 970，有13个stream式多处理器（SM），每个都有128个CUDA核心。 Cuda核心也被称为stream处理器（SP）。
您可以定义将块映射到GPU的网格。
您可以定义将线程映射到stream处理器的模块（每个SM 128个Cuda内核）。
一个warp总是由32个线程组成，一个warp的所有线程都是同时执行的。
要使用GPU的全部可能的function，每个SM需要比SM具有更multithreading的SP。 对于每个计算能力，一次可以有一定数量的线程驻留在一个SM中。 您定义的所有块都排队，等待SM获得资源（SP数量空闲），然后加载。 SM开始执行Warps。 由于一个Warp只有32个线程，一个SM有128个SP，所以SM可以在给定的时间内执行4个Warp。 事情是如果线程做内存访问线程将阻塞，直到它的内存请求得到满足。 在数字上：对SP的算术计算具有18-22个周期的等待时间，而非高速caching的全局存储器访问可能需要高达300-400个周期。 这意味着如果一个warp的线程正在等待数据，那么只有128个SP的一个子集可以工作。 因此，如果可用的话，调度器切换到执行另一个变形。 如果这个经线阻塞了，它将执行下一个等等。 这个概念被称为延迟隐藏。 经纱的数量和块的大小决定了占有率（SM可以select执行多less次经纱）。 如果占用率高，SP不太可能没有工作。
你的陈述，每个cuda核心将一次执行一个块是错误的。 如果您谈论stream式多处理器，他们可以执行驻留在SM中的所有线程的warp。 如果一个块的大小为256个线程，并且您的GPU允许每个SM有2048个线程驻留，则每个SM将有8个块，SM可以从这些块中select要执行的warp。 执行的所有线程的线程是并行执行的。
您可以在这里find不同计算能力和GPU架构的编号： https ： //en.wikipedia.org/wiki/CUDA#Limitations
您可以从Nvidia 占用率计算表（由Nvidia）下载占用率计算表。
只有当SM有足够的线程块资源（共享内存，warp，registers，barrier，…）时，计算工作分配器才会在SM上安排一个线程块（CTA）。 线程块级别的资源，如共享内存被分配。 分配为线程块中的所有线程创build足够的warp。 资源pipe理器将循环分配的经纱分配给SM子分区。 每个SM子分区都包含一个warp调度器，寄存器文件和执行单元。 一旦一个warp被分配给一个子分区，它将一直保持在subpartition上，直到它完成或被一个上下文切换（Pascal架构）预占。 在上下文切换恢复时，warp将被恢复到相同的SM相同的warp-id。
当warp中的所有线程都完成时，warp调度程序等待warp发出的所有未完成指令完成，然后资源pipe理器释放包含warp-id和register文件的warp level资源。
当一个线程块中的所有线程都变形完成时，块级资源被释放，SM通知计算工作分配器该块已经完成。
一旦将变形分配给子分区并且分配了所有资源，则变形被认为是有效的，这意味着变形调度器正在主动跟踪变形的状态。 在每个循环中，warp调度程序确定哪些活动warps被停止，哪些有资格发出指令。 warp调度程序select最高优先级的符合条件的warp，并从warp发出1-2个连续的指令。 双重问题的规则是特定于每个架构的。 如果一个warp发出一个内存的负载，它可以继续执行独立的指令，直到达到一个独立的指令。 然后经纱将报告停止，直到加载完成。 依赖math指令也是如此。 SM架构旨在通过在每个经线之间切换来隐藏ALU和内存延迟。
这个答案不使用术语CUDA核心，因为这引入了一个不正确的心理模型。 CUDA内核是stream水线单精度浮点/整数执行单元。 问题速率和依赖延迟是每个体系结构特有的。 每个SM子分区和SM具有其他执行单元，包括加载/存储单元，双精度浮点单元，半精度浮点单元，分支单元等。
为了最大限度地提高性能，开发人员必须了解块与经线与寄存器/线程的权衡。
术语占用率是SM上活动经纱与最大经纱的比率。 开普勒 – 帕斯卡尔架构（GP100除外）每个SM有4个warp调度器。 每个SM的最小的经纱数量应该至less等于经纱调度器的数量。 如果架构的执行延迟时间为6个周期（Maxwell和Pascal），那么每个调度器至less需要6个warp，每个调度器24个SM（24/64 = 37.5％占用率）来覆盖延迟。 如果线程具有指令级并行性，那么这可以被减less。 几乎所有的内核都会发出可变延迟指令，例如可能需要80-1000个周期的内存负载。 这需要每个warp调度器更多的主动warps来隐藏延迟。 对于每个内核，在经线数量和其他资源（如共享存储器或寄存器）之间存在一个折衷点，因此不build议优化100％占用率，因为可能会做出其他牺牲。 CUDA分析器可以帮助确定指令发布率，占用率和失速原因，以帮助开发人员确定平衡。
线程块的大小会影响性能。 如果内核有大块并且使用同步障碍，那么障碍失速可能会成为失速的原因。 这可以通过减less每个线程块的扭曲来缓解。
