我正在通過OpenCL在GPU上進行float[]
陣列縮減(查找最小和最大值)。OpenCL中min()和max()的中性元素減少
我正在從global
內存加載一些元素到local
內存爲每個工作組。當全局大小不是工作組大小的倍數時,我填充全局大小,使其成爲全局大小的倍數。數組末尾的工作項將減少的中性元素放入內存中。
但是,該中性元素應爲max()
- 最大功能? The OpenCL documentation給出MAXFLOAT
,HUGE_VALF
和INFINITY
作爲非常大的肯定(或無符號)float
值。 以中性元素爲例如-INFINITY
是否合理?
現在我正在使用HUGE_VALF
作爲min()
的中性元素,但文檔中還指出HUGE_VALF
被用作錯誤值,所以這可能不是一個好主意。
減少內核(代碼):
#define NEUTRAL_ELEMENT HUGE_VALF
#define REDUCTION_OP min
__kernel void reduce(__global float* weights,
__local float* weights_cached
)
{
unsigned int id = get_global_id(0);
// Load data
if (id < {{ point_count }}) {
weights_cached[get_local_id(0)] = weights[id];
} else {
weights_cached[get_local_id(0)] = NEUTRAL_ELEMENT;
}
barrier(CLK_LOCAL_MEM_FENCE);
// Reduce
for(unsigned int stride = get_local_size(0)/2; stride >= 1; stride /= 2) {
if (get_local_id(0) < stride) {
weights_cached[get_local_id(0)] = REDUCTION_OP(weights_cached[get_local_id(0)], weights_cached[get_local_id(0) + stride]);
barrier(CLK_LOCAL_MEM_FENCE);
}
// Save
weights[get_group_id(0)] = weights_cached[0];
}
編輯: 其實,我結束了使用fmin()
和fmax()
與NAN
作爲中性元素在一起 - 這基本上是有保證根據OpenCL documentation作爲工作數值將始終返回(NAN
僅在給出兩個NAN
值時才返回)。
由於該問題顯示了一些不確定性--INFINITY,我們可能想強調這是一個常規支持的事情。您可以將一元'-'運算符應用於INFINITY,並獲取-infinity的值,並且應該按照其數學意義按預期工作。 –
@EricPostpischil:正確; -infinity的存在由OpenCL標準保證(特別是要求'float'是IEEE-754 binary32格式)。 –