2013-01-23 66 views
2

我正在通過OpenCL在GPU上進行float[]陣列縮減(查找最小和最大值)。OpenCL中min()和max()的中性元素減少

我正在從global內存加載一些元素到local內存爲每個工作組。當全局大小不是工作組大小的倍數時,我填充全局大小,使其成爲全局大小的倍數。數組末尾的工作項將減少的中性元素放入內存中。

但是,該中性元素應爲max() - 最大功能? The OpenCL documentation給出MAXFLOAT,HUGE_VALFINFINITY作爲非常大的肯定(或無符號)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值時才返回)。

回答

2

引述的OpenCL標準:

HUGE_VALF計算爲+無窮大。

因此,有使用HUGE_VALFINFINITY之間沒有真正的區別(除了隱含的意圖);要麼減少min會正常工作。在清晰度方面,我略微偏好INFINITY,因爲HUGE_VALF在概念上用於邊緣情況下的回報,但事實並非如此。

同樣,使用-INFINITY減少max

MAX_FLOAT如果數組包含無窮大,則不會像中性元素那樣正確運行。

+1

由於該問題顯示了一些不確定性--INFINITY,我們可能想強調這是一個常規支持的事情。您可以將一元'-'運算符應用於INFINITY,並獲取-infinity的值,並且應該按照其數學意義按預期工作。 –

+1

@EricPostpischil:正確; -infinity的存在由OpenCL標準保證(特別是要求'float'是IEEE-754 binary32格式)。 –