我可以强制进行某些计算，尽管它们的结果没有在内核中使用吗？

Question

我正在尝试分析我的 CUDA 内核执行的计算的某一部分的“成本”。当然，还有分析的使用。但我想要确定的是，通过投资加速或替换该特定部分，我可以获得多少收益。

我想通过确定当我删除该部分时内核的行为方式来进行估计 - 不用更有效的东西替换它，只是跳过它。所以，我的内核可能看起来像（半伪代码，尽管技术上有效的 CUDA C++）：

// -- snip --
auto data = read_data(input);
auto massaged_data = massage(data);
auto result = heavy_lifting(massaged_data); // (*)
update(output, result);

我将

(*)

行替换为

auto result = dummy_result();

默认初始化，或设置为 0 等。

不幸的是，这种方法不起作用，因为 nvcc 前端或 ptxas 都会优化掉与未使用的数据相关的大部分代码：大部分

read_data()

和大部分

massage()

消失，甚至可能全部消失。

我还尝试将一些变量和字段标记为易失性 - 要么是结果，要么是计算过程中的中间值，或者两者都是；但这没有帮助。

我的问题：我可以强制计算

massaged_data

，尽管它以后从未使用过，并且不会禁用任何编译器优化吗？

注意：如果您相信这可以可以用

volatile

来完成，但也许我没有正确地挥发我的程序，这将是一个有效的答案，但尝试给出一个具体的例子来说明如何实际做到这一点是的，如果您知道的话，请提及潜在的“陷阱”。

Answer 1

我可以强制计算 Massaged_data，尽管它以后从未使用过，并且不会禁用任何编译器优化吗？

是的。 Paulius Micikevicius 在 CUDA 性能分析的早期提出了一种设计模式，如下所示：

__global__ void mykernel(....., T *dummy, int benchmark=0)
{

    .....

    // Some code

    .....

    // Code you want to conditionally disable for benchmarking
    T result;
    if (benchmark > 0) {
         result = dobigcalculation();
    }

    // More code

    if (benchmark > 0) *dummy = result;

}

因为

dobigcalculation

的结果连接到全局内存写入，并且基准测试在编译时不会解析为立即常量，所以优化不会将任何内容识别为死代码并将其删除。

完美吗？没有。

您可能会更改内核的寄存器配置文件，并且您至少添加一个额外的参数，这可能会改变编译器构建代码的方式。在某些情况下，代码重新排序启发式方法还可能使复制未修改和修改代码的真实流程变得困难，这意味着您确实需要检查编译器发出的 PTX。虽然对于

dobigcalculation

返回简单 POD 类型的情况来说，设计模式很简单，但如果函数在类或类上运行，则设计赋值操作来真正确保部分代码不会被默默地删除可能会更加困难。一个引用，而不是一个返回值。

但是，这是您在完全优化的代码中衡量给定代码节的影响所能做的最好的事情。