- html - 出于某种原因,IE8 对我的 Sass 文件中继承的 html5 CSS 不友好?
- JMeter 在响应断言中使用 span 标签的问题
- html - 在 :hover and :active? 上具有不同效果的 CSS 动画
- html - 相对于居中的 html 内容固定的 CSS 重复背景?
我的 Cuda 程序获得了显着的性能提升(平均),这取决于 block 的大小和 block 的数量; “线程”的总数保持不变。 (我不确定线程是否是正确的术语......但我将在这里使用它;对于每个内核,线程总数是(# of blocks)*(block size))。我制作了一些图表来说明我的观点。
但首先请允许我先解释一下我的算法,但是我不确定它的相关性如何,因为我想这是适用于所有 GPGPU 程序的东西。但也许我我错了。
基本上,我遍历了逻辑上被视为二维数组的大型数组,其中每个线程从数组中添加一个元素,并将该值的平方添加到另一个变量,然后在最后将该值写入另一个数组,在每次读取期间,所有线程都以某种方式移动。这是我的内核代码:
__global__ void MoveoutAndStackCuda(const float* __restrict__ prestackTraces, float* __restrict__ stackTracesOut,
float* __restrict__ powerTracesOut, const int* __restrict__ sampleShift,
const unsigned int samplesPerT, const unsigned int readIns,
const unsigned int readWidth, const unsigned int defaultOffset) {
unsigned int globalId = ((blockIdx.x * blockDim.x) + threadIdx.x); // Global ID of this thread, starting from 0 to total # of threads
unsigned int jobNum = (globalId / readWidth); // Which array within the overall program this thread works on
unsigned int readIndex = (globalId % readWidth) + defaultOffset; // Which sample within the array this thread works on
globalId = (jobNum * samplesPerT) + readIndex; // Incorperate default offset (since default offset will also be the offset of
// index we will be writing to), actual globalID only needed for above two variables.
float stackF = 0.0;
float powerF = 0.0;
for (unsigned int x = 0; x < readIns; x++) {
unsigned int indexRead = x + (jobNum * readIns);
float value = prestackTraces[readIndex + (x * samplesPerT) + sampleShift[indexRead]];
stackF += value;
powerF += (value * value);
}
stackTracesOut[globalId] = stackF;
powerTracesOut[globalId] = powerF;
}
现在是这篇文章的重点,当调用这段代码时
MoveoutAndStackCuda<<<threadGroups, threadsPerGroup>>>(*prestackTracesCudaPtr,
*stackTracesOutCudaPtr, *powerTracesOutCudaPtr,
*sampleShiftCudaPtr, samplesPerT, readIns,
readWidth, defaultOffset);
我所做的只是在 <<<>>> 中改变 threadGroups 和 threadsPerGroup,其中 threadGroups.x * threadsPerGroup.x 保持不变。 (如前所述,这是一个一维问题)。
我将 block 大小增加了 64,直到达到 1024。我预计不会有任何变化,因为我认为只要 block 大小大于 32,我相信这是核心中 ALU 的数量,它就会运行得一样快尽可能。看看我制作的这张图:
对于此特定大小,线程总数为 5000 * 5120,因此例如,如果 block 大小为 64,则有 ((5000 * 5120)/64) 个 block 。 出于某种原因, block 大小为 896、768 和 512 时性能显着提升。为什么?
我知道这看起来是随机的,但这张图中的每个点都是 50 个测试的平均值!
这是另一个图表,这次线程总数为 (8000 * 8192)。这次提升是 768 和 960。
还有一个例子,这次是一个比其他两个问题小的作业(总线程数是 2000 * 2048):
事实上,这是我用这些图制作的相册,每个图代表问题的不同规模:graph album .
我正在运行这个 Quadro M5000 ,它有 2048 个 Cuda 核心。我相信每个 Cuda Core 都有 32 个 ALU,所以我假设在任何给定时间可能发生的计算总数是 (2048 * 32)?
那么如何解释这些神奇的数字呢?我认为它可能是线程总数除以 cuda 核心数,或除以 (2048 * 32),但到目前为止,我没有发现与我相册中所有图表的任何相关性。我可以做另一项测试来帮助缩小范围吗?我想找出运行此程序以获得最佳结果的 block 大小。
我也没有包括它,但我也做了一个测试,其中 block 大小从 32 减少了 1,并且速度呈指数级下降。这对我来说很有意义,因为在给定的多处理器中,我们每组的本地线程比 ALU 少。
最佳答案
基于这个声明:
I incremented the block size by 64 until I reached 1024. I expected no change, because I figured as long as block size is greater than 32, which I believe is the # of ALUs in a core, it would run as fast as possible.
我想说关于 GPU 有一个您可能不知道的重要概念:GPU 是一种“隐藏延迟”的机器。它们主要通过向它们公开大量可用(并行)工作来隐藏延迟。这可以粗略地概括为“很多线程”。对于 GPU 来说,一旦你有足够的线程来覆盖“核心”或执行单元的数量,这就足够了,这是一个完全错误的想法。 不是。
作为(初学者)GPU 程序员,您应该忽略 GPU 中的内核数量。您需要很多线程。在内核级别和每个 GPU SM。
一般来说,当您为每个 SM 提供更多线程时,GPU 在执行其他有用工作时隐藏延迟的能力就会提高。这解释了您所有图表中的总体趋势,即斜率通常从左到右向下(即平均性能通常会随着您为每个 SM 提供更多公开工作而增加)。
然而,这并没有解决高峰和低谷。 GPU 存在大量可能影响性能的架构问题。我不会在这里提供完整的治疗。但让我们来看一个案例:
Why does performance in the first graph increase up to 512 threads, then suddenly decrease at 576 threads?
这很可能是一种占用效应。 GPU 中的 SM 最多有 2048 个线程。根据前面的讨论,当我们最大化线程补充时,SM 将具有最大的隐藏延迟的能力(因此通常提供最大的平均性能),最高可达 2048。
对于 512 个线程的 block 大小,我们可以在 SM 上恰好容纳 4 个这样的线程 block ,然后它将有 2048 个线程的补充,可以从中选择工作和延迟隐藏。
但是当你把threadblock的大小改成576时,4*576 > 2048,所以我们不能再在每个SM上放4个threadblock了。这意味着,对于该内核配置,每个 SM 将使用 3 个线程 block 运行,即 2048 个线程中的 1728 个线程。从 SM 的角度来看,这实际上更糟,比之前允许 2048 个线程的情况要好,因此它可能是性能从 512 到 576 个线程下降的原因的一个指标(就像它增加一样从 448 到 512,这涉及瞬时占用率的类似变化)。
由于上述原因,当我们改变每个 block 的线程数时,经常会看到像您所显示的那样的性能图表。
其他具有粒度(量化)效果的占用限制器可能会导致性能图中出现类似的峰值行为。例如,您的问题中没有足够的信息来推测每个线程的寄存器使用情况,但占用的限制因素可能是每个线程使用的寄存器。当您改变线程补充时,您会发现每个 SM 驻留的 block 补充可能类似地发生变化,这可能会导致占用率发生变化(向上和向下),从而导致性能发生变化。
要进一步深入研究,我建议您花一些时间了解占用率、每个线程的寄存器以及各种分析器的性能分析功能。已经有很多关于这些主题的信息; google 是你的 friend ,注意 question/answers在上面的评论中链接,作为一个合理的起点。要全面研究入住率及其对性能的影响,需要比您在此处提供的信息更多的信息。它基本上需要一个 MCVE以及确切的编译命令行,以及您正在运行的平台和 CUDA 版本。编译器的每线程寄存器使用量受所有这些因素的影响,其中大部分是您没有提供的。
关于performance - GPGPU: block 大小对程序性能的影响,为什么我的程序在非常特定的大小下运行得更快?,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/41405026/
好的,所以我想从批处理文件运行我的整个工作环境... 我想要实现什么...... 打开新的 powershell,打开我的 API 文件夹并从该文件夹运行 VS Code 编辑器(cd c:\xy;
我正在查看 Cocoa Controls 上的示例并下载了一些演示。我遇到的问题是一些例子,比如 BCTabBarController ,不会在我的设备上构建或启动。当我打开项目时,它看起来很正常,没
我刚刚开始学习 C 语言(擅长 Java 和 Python)。 当编写 C 程序(例如 hello world)时,我在 ubuntu cmd 行上使用 gcc hello.c -o hello 编译
我在 php 脚本从 cron 开始运行到超时后注意到了这个问题,但是当它从命令行手动运行时这不是问题。 (对于 CLI,PHP 默认的 max_execution_time 是 0) 所以我尝试运行
我可以使用命令行运行测试 > ./node_modules/.bin/wdio wdio.conf.js 但是如果我尝试从 IntelliJ 的运行/调试配置运行它,我会遇到各种不同的错误。 Fea
Error occurred during initialization of VM. Could not reserve enough space for object heap. Error: C
将 Anaconda 安装到 C:\ 后,我无法打开 jupyter 笔记本。无论是在带有 jupyter notebook 的 Anaconda Prompt 中还是在导航器中。我就是无法让它工作。
我遇到一个问题,如果我双击我的脚本 (.py),或者使用 IDLE 打开它,它将正确编译并运行。但是,如果我尝试在 Windows 命令行中运行脚本,请使用 C:\> "C:\Software_Dev
情况 我正在使用 mysql 数据库。查询从 phpmyadmin 和 postman 运行 但是当我从 android 发送请求时(它返回零行) 我已经记录了从 android 发送的电子邮件是正确
所以这个有点奇怪 - 为什么从 Java 运行 .exe 文件会给出不同的输出而不是直接运行 .exe。 当 java 在下面的行执行时,它会调用我构建的可与 3CX 电话系统配合使用的 .exe 文
这行代码 Environment.Is64BitProcess 当我的应用单独运行时评估为真。 但是当它在我的 Visual Studio 单元测试中运行时,相同的表达式的计算结果为 false。 我
关闭。这个问题是opinion-based .它目前不接受答案。 想要改进这个问题? 更新问题,以便 editing this post 可以用事实和引用来回答它. 关闭 8 年前。 Improve
我写了一个使用 libpq 连接到 PostgreSQL 数据库的演示。 我尝试通过包含将 C 文件连接到 PostgreSQL #include 在我将路径添加到系统变量 I:\Program F
如何从 Jenkins 运行 Android 模拟器来运行我的测试?当我在 Execiute Windows bath 命令中写入时,运行模拟器的命令: emulator -avd Tester 然后
我已经配置好东西,这样我就可以使用 ssl 登录和访问在 nginx 上运行的 errbit 我的问题是我不知道如何设置我的 Rails 应用程序的 errbit.rb 以便我可以运行测试 nginx
我编写了 flutter 应用程序,我通过 xcode 打开了 ios 部分并且应用程序正在运行,但是当我通过 flutter build ios 通过 vscode 运行应用程序时,我得到了这个错误
我有一个简短的 python 脚本,它使用日志记录模块和 configparser 模块。我在Win7下使用PyCharm 2.7.1和Python 3.3。 当我使用 PyCharm 运行我的脚本时
我在这里遇到了一些难题。 我的开发箱是 64 位的,windows 7。我所有的项目都编译为“任何 CPU”。该项目引用了 64 位版本的第 3 方软件 当我运行不使用任何 Web 引用的单元测试时,
当我注意到以下问题时,我正在做一些 C++ 练习。给定的代码将不会在 Visual Studio 2013 或 Qt Creator 5.4.1 中运行/编译 报错: invalid types 'd
假设我有一个 easteregg.py 文件: from airflow import DAG from dateutil import parser from datetime import tim
我是一名优秀的程序员,十分优秀!