3

OpenCL カーネルが浮動小数点例外をスローしています。問題を引き起こしていると思われる行だけに減らしました。

ライン交換したら

acc.x += sin(distSqr);

と

acc.x += cos(distSqr);

また

acc.x += sqrt(distSqr);

あるいは単に

acc.x += (distSqr);

カーネルは正常に動作します。なんで？注: 私のグローバル ワーク サイズは、ローカル ワーク サイズで割り切れます。

ありがとう。

カーネルは次のとおりです。

__kernel void compute_forces(
                      __global float3 *x3a,
                      __global float3 *p3a,
                      __global float3 *x3b,
                      __global float3 *p3b,
                      __global float3 *f3a,
                      float dt,
                      float qQa,
                      float qQb,
                      float qma,
                      float qmb,
                      int n0a,
                      int n1a,
                      int n0b,
                      int n1b,
                      float xmin,
                      float ymin,
                      float epsSqr,
                      float force_fac,
                        __local float3 *localx
                      )
{


//we are going to compute the force between parts (n0a-n1a) and (n0b-n1b)
//Each particle loads the particle in the current block into local memory, so
unsigned int tid = get_local_id(0);
unsigned int gid = get_global_id(0);

unsigned int ninter=0;

// position of this work-item
float3 myPos = x3a[gid];
float3 acc = (float3)(0.0f, 0.0f, 0.0f);


    // Synchronize to make sure data is available for processing
    barrier(CLK_LOCAL_MEM_FENCE);


    for(int j = 0; j < 2; ++j)
    {

        float3 r=-myPos;
        float distSqr = r.x * r.x;


        // accumulate effect of all particles
        acc.x += sin(distSqr);

        ninter++;

    }//j

    // Synchronize so that next tile can be loaded
    barrier(CLK_LOCAL_MEM_FENCE);



f3a[gid]+=acc;
f3a[gid].x=(float)ninter;

}

カーネルを次のように呼び出します。

err=clSetKernelArg(k_compute_forces, 0, sizeof(_x3), &_x3);
err=clSetKernelArg(k_compute_forces, 1, sizeof(_p3), &_p3);
err=clSetKernelArg(k_compute_forces, 2, sizeof(_x3), &_x3);
err=clSetKernelArg(k_compute_forces, 3, sizeof(_p3), &_p3);
err=clSetKernelArg(k_compute_forces, 4, sizeof(_f3), &_f3);
err=clSetKernelArg(k_compute_forces, 5, sizeof(dt_float), &dt_float);
err=clSetKernelArg(k_compute_forces, 6, sizeof(qQa), &qQa);
err=clSetKernelArg(k_compute_forces, 7, sizeof(qQb), &qQb);
err=clSetKernelArg(k_compute_forces, 8, sizeof(qma), &qma);
err=clSetKernelArg(k_compute_forces, 9, sizeof(qmb), &qmb);
    err=clSetKernelArg(k_compute_forces,10, sizeof(n0a), &n0a);
    err=clSetKernelArg(k_compute_forces,11, sizeof(n1a), &n1a);
    err=clSetKernelArg(k_compute_forces,12, sizeof(n0b), &n0b);
    err=clSetKernelArg(k_compute_forces,13, sizeof(n1b), &n1b);
err=clSetKernelArg(k_compute_forces,14, sizeof(xmin_float), &xmin_float);
err=clSetKernelArg(k_compute_forces,15, sizeof(ymin_float), &ymin_float);
err=clSetKernelArg(k_compute_forces,16, sizeof(epsSqr), &epsSqr);
err=clSetKernelArg(k_compute_forces,17, sizeof(force_fac), &force_fac);
err=clSetKernelArg(k_compute_forces,18, parts_per_block*sizeof(cl_float3),NULL);

    err=clEnqueueNDRangeKernel(queue, k_compute_forces, work_dim, NULL, global_work_size, local_work_size, 0, NULL, &k_compute_forces_completion);

編集: sin 関数は約 1.0e-12 より小さい float を処理できないと思います。

acc.x += sin(1.0e-12);

正常に動作しますが、

acc.x += sin(1.0e-13);

例外をスローします。これは、sin の代わりに sin_half が呼び出されていることを示唆しているように思われます...これは、最適化によって行われた置き換えでしょうか。

上記の行の直前にprintfステートメントを追加すると

printf("distSqr=%g\n",distSqr);

その後、エラーは「浮動小数点例外」から「divisionErrorHandler」に関するものに変わります（出力テキストがごちゃごちゃしているため、わかりにくいですが）。

4

1 に答える 1

0

sin(x) は x の小さな値に対しては機能していないように見えるので、Nvidia ドライバーがそれを、ハードウェアに直接実装された関数に評価される native_sin(x) に置き換えているのではないかと疑っていますが、そうではない可能性があります。正確であるか、または数値の全範囲をサポートしています。ビルド オプション「-cl-opt-disable」を clBuildProgram の呼び出しに追加することをお勧めします。これにより、すべての最適化が無効になり、コンパイラの最適化に問題があるかどうかがわかります。

于 2013-04-01T17:05:40.650 に答える