CUDA C pow() 函数导致 modf() 和类似功能的舍入错误

Question

我通过将线程添加到某个正在排序的值来跨线程跟踪索引：value + threadIdx.x*.001。这（理论上）有效，因为我排序的值是整数，但它们可以是负数。

当我去拉小数部分时，我运行出现舍入错误。具体来说，小数部分在某些情况下计算为 1.0（例如 -3.0）。

重现问题的代码：

code ="""

/* CUDA's random number stuff */
#include <stdio.h>
#include <stdlib.h>
#include <curand.h>
#include <curand_kernel.h>
#include <math.h>

#define NUM_THREADS 64
#define INDEX_BASE_10 .001

extern "C" {

__device__ double decimal_part_to_index(double value) {
  
  value = fabs(value);
  value -= (long)value;
  value *= 1000;
  
  //if (value > NUM_THREADS) {
  //  return 0;
  //}
  
  return value;

}/*DECIMAL PART TO INDEX*/


/****GLOBAL****GLOBAL****GLOBAL****/
/*********GLOBAL****GLOBAL*********/
__global__ void foo(double *input_array,
                    double *result_array) {

  bool bool_array[2] = {0,1};

  int some_int = 3;
  
  input_array[threadIdx.x] = pow(some_int,1) + threadIdx.x*.INDEX_BASE_10;

  result_array[threadIdx.x] = decimal_part_to_index(input_array[threadIdx.x]);
  


}
}
"""
mod = SourceModule(code, no_extern_c=True)
foo = mod.get_function("foo")


GPU_result_1 = gpuarray.zeros(64, dtype=numpy.double)

GPU_result_2 = gpuarray.zeros(64, dtype=numpy.double)

for i in range(1):
    foo(GPU_result_1,
    GPU_result_2,
    block=(64,1,1), grid=(1,1,1))

正如评论中所建议的那样，滥用 pow() 和 int,int 参数可能是罪魁祸首。

原始 POST 下面：

__device__ double decimal_part_to_index(double value) {
  
  value = abs(value);
  value -= (long)value;
  value *= 1000;
  
  //This catches rounding error for NUM_THREADS = 256
  //When value = -3.0, value -= (long)value is 1.0???
  if (value > NUM_THREADS) {
    return 0;
  }
  
  return value;

}/*DECIMAL PART TO INDEX*/

我也尝试了 modf()，得到了类似的结果，但我必须捕捉到更多的舍入错误实例。我的问题是上述方法是否适用于更大的 NUM_THREADS 值，或者我是否会遇到更多必须捕获的舍入实例？例如，上面的方法是否有其他值会给出 1.0 的小数部分，这会破坏我的方案？在我的例子中，整数部分通常很小但可以变大。有什么解决方案可以安全地取小数部分吗？似乎无论我尝试什么，我都会遇到类似的舍入问题。

编辑：除以和乘以 64 的输出示例：NUM_THREADS = 64

在第一个示例中，第二个位置显示 -3.0 评估为 64，但应该评估为 0。所以我们必须捕捉并设置为 0。

array([ 2.359375, -3.      , -3.015625, -3.03125 , -3.046875, -3.0625  ,
       -3.078125, -3.09375 , -3.109375, -3.125   , -3.140625, -3.15625 ,
       -3.171875, -3.1875  , -3.203125, -3.21875 , -3.234375, -3.25    ,
       -3.265625, -3.28125 , -3.296875, -3.3125  , -3.328125, -3.34375 ,
       -3.375   , -3.390625, -3.40625 , -3.421875, -3.4375  , -3.453125,
       -3.46875 , -3.484375, -3.5     , -3.515625, -3.53125 , -3.546875,
       -3.5625  , -3.578125, -3.59375 , -3.609375, -3.625   , -3.640625,
       -3.65625 , -3.671875, -3.6875  , -3.703125, -3.71875 , -3.734375,
       -3.75    , -3.765625, -3.78125 , -3.796875, -3.8125  , -3.828125,
       -3.84375 , -3.859375, -3.875   , -3.890625, -3.90625 , -3.921875,
       -3.9375  , -3.953125, -3.96875 , -3.984375])

array([23., 64.,  1.,  2.,  3.,  4.,  5.,  6.,  7.,  8.,  9., 10., 11.,
       12., 13., 14., 15., 16., 17., 18., 19., 20., 21., 22., 24., 25.,
       26., 27., 28., 29., 30., 31., 32., 33., 34., 35., 36., 37., 38.,
       39., 40., 41., 42., 43., 44., 45., 46., 47., 48., 49., 50., 51.,
       52., 53., 54., 55., 56., 57., 58., 59., 60., 61., 62., 63.])

在第二个示例中，第一个位置显示 -2.0 像预期的那样计算为 0。

array([-2.      , -2.140625, -2.203125, -2.4375  , -2.46875 , -2.5     ,
       -2.6875  , -3.015625, -3.03125 , -3.046875, -3.0625  , -3.078125,
       -3.09375 , -3.109375, -3.125   , -3.15625 , -3.171875, -3.1875  ,
       -3.21875 , -3.234375, -3.25    , -3.265625, -3.28125 , -3.296875,
       -3.3125  , -3.328125, -3.34375 , -3.359375, -3.375   , -3.390625,
       -3.40625 , -3.421875, -3.453125, -3.484375, -3.515625, -3.53125 ,
       -3.546875, -3.5625  , -3.578125, -3.59375 , -3.609375, -3.625   ,
       -3.640625, -3.65625 , -3.671875, -3.703125, -3.71875 , -3.734375,
       -3.75    , -3.765625, -3.78125 , -3.796875, -3.8125  , -3.828125,
       -3.84375 , -3.859375, -3.875   , -3.890625, -3.90625 , -3.921875,
       -3.9375  , -3.953125, -3.96875 , -3.984375])

array([ 0.,  9., 13., 28., 30., 32., 44.,  1.,  2.,  3.,  4.,  5.,  6.,
        7.,  8., 10., 11., 12., 14., 15., 16., 17., 18., 19., 20., 21.,
       22., 23., 24., 25., 26., 27., 29., 31., 33., 34., 35., 36., 37.,
       38., 39., 40., 41., 42., 43., 45., 46., 47., 48., 49., 50., 51.,
       52., 53., 54., 55., 56., 57., 58., 59., 60., 61., 62., 63.])

Answer 1

上述CUDA内核代码更改：

int some_int = 3;

至：

double some_int = 3;

将解决 pow() 乱拉小数部分的问题。

在我的例子中，我试图使用 pow() 函数来执行 if/else 逻辑来确定符号。使用 if/else 语句而不是 pow() 解包逻辑也解决了问题，而无需更改类型或重铸任何整数。

CUDA C pow() 函数导致 modf() 和类似功能的舍入错误

CUDA C pow() function causing rounding errors with modf() and similar functionality

cuda