为什么OpenCL内核调用sin而不是cos的浮点异常？

Question

我的 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 的浮点数，因为行：

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

运行良好，但

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

抛出异常。这似乎表明正在调用 sin_half 而不是 sin...我想知道这是否是优化所做的替换。

如果我在上面一行之前添加一个 printf 语句

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

然后错误从“浮点异常”变为有关“divisionErrorHandler”的内容（尽管由于输出文本混乱而难以辨认）。

Answer 1

由于 sin(x) 似乎不适用于 x 的小值，我怀疑 Nvidia 驱动程序正在用 native_sin(x) 替换它，它评估为直接在硬件中实现的函数，但可能不是那么准确或支持所有数字。我建议在对 clBuildProgram 的调用中添加构建选项“-cl-opt-disable”，因为这应该会禁用所有优化，告诉我们编译器优化是否有问题。