OpenCL/C pow(x,0.5) != sqrt(x)

OpenCL/C pow(x,0.5) != sqrt(x)

我想我遇到了一些非常奇怪的边界情况,可能有双精度问题,我想知道发生了什么。

在我使用的 OpenCL 内核中:

#pragma OPENCL EXTENSION cl_khr_fp64 : enable

__private int k = 2; // I need k to be an int, because I want to use as a counter
__private double s = 18;
__private double a = 1;

a = a/(double)k; // just to show, that I make in-place typecasting of k
a = k+1;
k = (int)a; //to show that I store k in a double buffer in an intermediate-step
if ((k-1)==2)
{
//    k = 3;
    s = pow(s/(double)(k-1),0.5);
}

这让我想到 s = 2.999[...]6 但是,如果我取消注释 k=3 行,我会得到(在我看来)正确的结果 s = 3。这是为什么?

作为附带信息:当我这样做时不会发生相同的行为

s = sqrt(s/(double)(k-1))

下面是 pyopencl 的完整、最小内核和主机代码

内核(Minima.cl):

#pragma OPENCL EXTENSION cl_khr_fp64 : enable

__kernel void init_z(__global double * buffer)
{
    __private int x = get_global_id(0);
    __private int y = get_global_id(1);
    //w,h
    __private int w_y = get_global_size(1);
    __private int address = x*w_y+y;
    //h,w
    __private double init = 3.0;
    buffer[address]=init;
}

__kernel void root(__global double * buffer)
{
    __private int x = get_global_id(0);
    __private int y = get_global_id(1);
    //w,h
    __private int w_y = get_global_size(1);
    __private int address = x*w_y+y;
    //h,w
    __private double value = 18;
    __private int k;
    __private double out;
    k = (int) buffer[address];
  //k = 3;  If this line is uncommented, the result will be exact.
    out = pow(value/(double)(k-1), 0.5);
    buffer[address] = out;
}

主持人:

import pyopencl as cl
import numpy as np


platform = cl.get_platforms()[0]
devs = platform.get_devices()
device1 = devs[1]
h_buffer = np.empty((10,10)).astype(np.float64)
mf = cl.mem_flags
ctx = cl.Context([device1])
Queue1 = cl.CommandQueue(ctx,properties=cl.command_queue_properties.PROFILING_ENABLE)
Queue2 = cl.CommandQueue(ctx,properties=cl.command_queue_properties.PROFILING_ENABLE)
mf = cl.mem_flags
m_dic = {0:mf.READ_ONLY,1:mf.WRITE_ONLY,2:mf.READ_WRITE}


fi = open('Minimal.cl', 'r')
fstr = "".join(fi.readlines())
prg = cl.Program(ctx, fstr).build()
knl = prg.init_z
knl.set_scalar_arg_dtypes([None,])
knl_root = prg.root
knl_root.set_scalar_arg_dtypes([None,])



def f():
    d_buffer =  cl.Buffer(ctx,m_dic[2], int(10 * 10  * 8))
    knl.set_args(d_buffer)
    knl_root.set_args(d_buffer)
    a = cl.enqueue_nd_range_kernel(Queue2,knl,(10,10),None)
    b = cl.enqueue_nd_range_kernel(Queue2,knl_root,(10,10),None, wait_for = [a,])
    cl.enqueue_copy(Queue1,h_buffer,d_buffer,wait_for=[b,])
    return h_buffer
a = f()
a[0,0] # Getting the result on the host.

编辑: 由于一些不明确的地方,我再次更新了这个问题。我明白,对于相同的输入,powsqrt 的值不必相同。我的问题是,为什么 pow 显示相同输入的不同输出,这取决于我从哪里得到它。

二进制文件在 pastebin 上: k_explicit and k_read

printf("a%\n", out) 使用 k=3 行导致 0x1.8p+1,当它被注释掉时导致 0x1.7ffffffffffffp+1

浮点计算不准确。因此,不能指望使用一种算法 (sqrt) 和具有相同输入的不同算法 (pow()) 会给出按位相同的结果。如果两个结果都在数学真值的 ±epsilon 范围内,那么两种实现都是可以接受的。

通常,pow() 是根据 ln()exp()(中间有乘法)实现的,而 sqrt() 可以使用更快的实现(这第一步可能涉及将尾数减半)。

因此,经过大量讨论,我的代码似乎没有问题。该行为似乎与硬件有关,目前只能在 Nvidia 硬件上重现。我仍然对发生的 "why" 和 "how" 感兴趣,但在这种情况下问题得到了回答。