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_10nvcc に を指定すると、期待される のコードが生成されます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 に答える