0

我有以下（减少的测试用例！）CUDA 内核

__global__
void test(int n, const double* __restrict__ in, double* __restrict__ out)
{
    int idx = blockIdx.x * blockDim.x * threadIdx.x;

    if (idx < n)
    {
        out[idx] = 0.0*in[idx] + 1.0;
    }
}

我希望生成与out[idx] = 1.0. 0.0*in[idx]（当使用模板引擎自动生成内核时出现无操作表达式，其中0.0开始的生命为${template_parameter}.）但是，nvcc -arch sm_20 -ptx ...生成：

//
// Generated by NVIDIA NVVM Compiler
// Compiler built on Sat Sep 22 01:35:14 2012 (1348274114)
// Cuda compilation tools, release 5.0, V0.2.1221
//

.version 3.1
.target sm_20
.address_size 64

[...]

    mul.wide.s32    %rd5, %r1, 8;
    add.s64     %rd6, %rd2, %rd5;
    ld.global.f64   %fd1, [%rd6];
    fma.rn.f64  %fd2, %fd1, 0d0000000000000000, 0d3FF0000000000000;
    add.s64     %rd7, %rd1, %rd5;
    st.global.f64   [%rd7], %fd2;

有明确的全局负载和 FMA。然而，当-arch sm_10指定为 nvcc 时，它会生成out[idx] = 1.0. 是否有任何编译器选项/标志可以诱使它执行上述优化？

4

1 回答 1

2

由于 CUDA 通常遵循 IEEE-754 语义，因此不会优化浮点乘以零。特别是，IEEE-754 指定 +-0 * +-infinity = NaN、+-0 * NaN = NaN 和 +0 * -0 = -0。请参阅 C99 标准的“F.8.2 表达式转换”部分，了解浮点表达式的这种转换和其他转换。

于 2012-12-29T19:11:11.990 回答