CUDA C pow() 函数导致 modf() 和类似功能的舍入错误
CUDA C pow() function causing rounding errors with modf() and similar functionality
我通过将线程添加到某个正在排序的值来跨线程跟踪索引: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.])
上述CUDA内核代码更改:
int some_int = 3;
至:
double some_int = 3;
将解决 pow()
乱拉小数部分的问题。
在我的例子中,我试图使用 pow()
函数来执行 if/else
逻辑来确定符号。使用 if/else
语句而不是 pow()
解包逻辑也解决了问题,而无需更改类型或重铸任何整数。
我通过将线程添加到某个正在排序的值来跨线程跟踪索引: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.])
上述CUDA内核代码更改:
int some_int = 3;
至:
double some_int = 3;
将解决 pow()
乱拉小数部分的问题。
在我的例子中,我试图使用 pow()
函数来执行 if/else
逻辑来确定符号。使用 if/else
语句而不是 pow()
解包逻辑也解决了问题,而无需更改类型或重铸任何整数。