OpenCL：輸出可變長度數組

我們正在爲GPGPU課程進行分配。我們選擇了一種算法，在CPU上實現，現在將其轉換爲OpenCL。OpenCL：輸出可變長度數組

我們選擇的算法將模型加載爲一組三角形，並將它們柵格化爲體素。體素被定義爲點數據的VBO。然後我們使用幾何着色器將這些點轉換爲三維像素。

所以我們的OpenCL程序需要一個三角形列表並輸出一個可變的點列表。

並輸出一個可變長度的數組似乎是一個問題。

我們找到的解決方案是以原子方式遞增計數器並將該計數器用作輸出數組的索引和數組的最終大小。除了......我們的GPU都不支持原子操作的擴展。

這是我們到目前爲止有：

#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable 
#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable 
#pragma OPENCL EXTENSION cl_khr_int64_extended_atomics : enable 

#define POS1  i0 * 3 + 0 
#define POS2  i0 * 3 + 1 
#define POS3  i0 * 3 + 2 

void WritePosition(__global float* OutBuffer, uint inIndex, __global float* inPosition) 
{ 
    OutBuffer[ inIndex * 3 ] = inPosition[0]; 
    OutBuffer[ inIndex * 3 + 1] = inPosition[1]; 
    OutBuffer[ inIndex * 3 + 2] = inPosition[2]; 
} 

__kernel void Voxelize( 
    __global float* outPointcloudBuffer, 
    __global float* inTriangleBuffer, 
    __global uint* inoutIndex 
) 
{ 
    size_t i0 = get_global_id(0); 
    size_t i1 = get_local_id(0); 

    WritePosition(outPointcloudBuffer, inIndex[0], &inTriangleBuffer[ i0 ]); 

    //atomic_inc(inoutIndex[0]); 
    inoutIndex[0] = max(inoutIndex[0], i0); 
}

和這個輸出是非常奇怪的。我們正在測試一個非常小的模型（12個三角形，36個位置，108個浮標），我們得到的結果是31,63或95.總是16的倍數減去1.

如何獲得長度我們的可變長度輸出數組？

在此先感謝。

來源

2011-12-20 knight666

您的結果是16N-1，因爲你在整個warp上運行內核。要修復，請將三角形的總數傳遞給內核。如果'global_id'大於三角形數量，則返回。這樣，您只能運行與三角形一樣多的內核。 – 2011-12-20 18:33:58

我猜想，這通常解決如下：

第一遍：計算上使用scan（並行前綴總和）原語的GPU陣列的所需尺寸。以上鍊接包含Apple的示例實施。
使用掃描算法的結果在主機端分配所需的資源。請注意，掃描算法的結果通常可用作各個工作項結果的索引提示。
第二遍（可選）：將數組壓縮到需要在第三遍中考慮的那些元素。
第三遍：重新運行傳遞目標索引和分配數組的算法。

您可能想看看NVIDIA的OpenCL行軍立方體implementation，上面提到的所有三個通道都已實現。

Best，Christoph

來源

2011-12-20 09:24:36 cheind

OpenCL：輸出可變長度數組

回答

相關問題