CUDA核心,流式多处理器和块和线程的CUDA模型之间有什么关系？
什么被映射到什么和什么是并行化以及如何？什么是更有效,最大化块数或线程数？
我目前的理解是每个多处理器有8个cuda核心.并且每个cuda核心都能够一次执行一个cuda块.并且该块中的所有线程在该特定核心中串行执行.
它是否正确？
Edr*_*ric 63
线程/块布局在CUDA编程指南中有详细描述.特别是,第4章指出:
CUDA架构围绕可扩展的多线程流式多处理器(SM)阵列构建.当主机CPU上的CUDA程序调用内核网格时,将枚举网格块并将其分配给具有可用执行容量的多处理器.线程块的线程在一个多处理器上并发执行,并且多个线程块可以在一个多处理器上并发执行.当线程块终止时,在空出的多处理器上启动新块.
每个SM包含8个CUDA内核,并且在任何时候它们都执行32个线程的单个warp - 因此需要4个时钟周期才能为整个warp发出单个指令.您可以假设任何给定warp中的线程都以锁定步骤执行,但要在warp之间进行同步,则需要使用__syncthreads().
Joe*_*Fox 32
对于GTX 970,有13个流式多处理器(SM),每个处理器有128个Cuda核心.Cuda Cores也称为流处理器(SP).
您可以定义将块映射到GPU的网格.
您可以定义将线程映射到流处理器的块(每个SM的128个Cuda内核).
一个warp总是由32个线程组成,并且warp的所有线程同时执行.
要使用GPU的全部功能,每SM需要比SM具有SP更多的线程.对于每个计算能力,有一定数量的线程可以一次驻留在一个SM中.您定义的所有块都排队等待SM拥有资源(SP的数量空闲),然后加载它.SM开始执行Warps.由于一个Warp只有32个线程而SM有例如128个SP,因此SM可以在给定时间执行4个Warp.问题是如果线程进行内存访问,线程将阻塞,直到满足内存请求.在数字上:SP上的算术计算具有18-22个周期的延迟,而非缓存的全局内存访问可能需要300-400个周期.这意味着如果一个warp的线程正在等待数据,则只有128个SP的子集可以工作.因此,调度程序切换到执行另一个warp(如果可用).如果这个warp阻塞它执行下一个,依此类推.这个概念称为延迟隐藏.warp的数量和块大小决定了占用率(从SM可以选择执行的warp数量).如果占用率很高,则SP不太可能没有工作.
您声明每个cuda核心一次执行一个块是错误的.如果您谈论流式多处理器,他们可以从驻留在SM中的所有线程执行warp.如果一个块的大小为256个线程,并且您的GPU允许每个SM驻留2048个线程,则每个SM将有8个块,SM可以从中选择warp来执行.执行的warp的所有线程并行执行.
您可以在此处找到不同的计算功能和GPU架构的数字:https: //en.wikipedia.org/wiki/CUDA#Limitations
您可以从Nvidia 占用计算表(由Nvidia)下载占用计算表.
仅当SM具有足够的线程块资源（共享内存，扭曲，寄存器，屏障等）时，Compute Work Distributor才会在SM上调度线程块（CTA）。线程块级资源（如共享内存）已分配。分配为线程块中的所有线程创建足够的扭曲。资源管理器使用循环机制向SM子分区分配扭曲。每个SM子分区都包含一个扭曲调度程序，寄存器文件和执行单元。将扭曲分配给子分区后，它将保留在子分区上，直到完成或被上下文切换（Pascal架构）抢占为止。在上下文切换恢复时，扭曲将恢复为相同的SM相同的扭曲ID。
当warp中的所有线程都已完成时，warp调度程序将等待warp发出的所有未完成指令完成，然后资源管理器将释放warp级别的资源，其中包括warp-id和注册文件。
当线程块中的所有扭曲均完成时，则释放块级资源，并且SM通知Compute Work Distributor该块已完成。
一旦将warp分配给一个子分区并分配了所有资源，就将warp视为活动状态，这意味着warp调度程序正在主动跟踪warp的状态。在每个周期上，翘曲调度程序都会确定哪些活动的翘曲已停止，哪些合格的翘曲有资格发出指令。Warp调度程序选择优先级最高的合格Warp，并从Warp发出1-2个连续的指令。双重问题的规则特定于每种体系结构。如果warp发出内存负载，它可以继续执行独立的指令，直到到达相关指令为止。然后，扭曲将报告已停止，直到加载完成。依存数学指令也是如此。SM体系结构旨在通过在扭曲之间每个周期进行切换来隐藏ALU和内存延迟。
此答案未使用术语CUDA核心，因为这会引入错误的心理模型。CUDA内核是流水线式单精度浮点/整数执行单元。发行率和依赖性等待时间特定于每种体系结构。每个SM子分区和SM都有其他执行单元，包括加载/存储单元，双精度浮点单元，半精度浮点单元，分支单元等。
为了使性能最大化，开发人员必须了解块与扭曲与寄存器/线程之间的权衡。
占用率是指SM上活动翘曲与最大翘曲的比率。开普勒-Pascal架构（GP100除外）每个SM具有4个warp调度程序。每个SM的最小经纱数量应至少等于经纱调度程序的数量。如果架构具有6个周期（麦克斯韦和帕斯卡）的相关执行延迟，则每个调度程序至少需要6个扭曲，每个SM为24（24/64 = 37.5％占用），以覆盖延迟。如果线程具有指令级并行性，则可以减少并行度。几乎所有内核都发出可变的延迟指令，例如可能需要80-1000个周期的内存负载。这需要每个warp调度程序更多活动的warp以隐藏延迟。对于每个内核，在扭曲数量和其他资源（例如共享内存或寄存器）之间存在一个折衷点，因此不建议对100％的占用进行优化，因为可能会做出其他牺牲。CUDA探查器可以帮助确定指令的发出率，占用率和停顿原因，以帮助开发人员确定平衡。
线程块的大小会影响性能。如果内核具有大块并使用同步屏障，则屏障停顿可能是导致停顿的原因。这可以通过减少每个线程块的扭曲来缓解。