Обработка граничных условий в OpenCL/CUDA

Question

Учитывая трехмерную однородную сетку, я хотел бы установить значения пограничных ячеек относительно значений их ближайшего соседа внутри сетки. Например, для сетки 10x10x10 для вокселя в координатах (0, 8, 8) я хотел бы установить значение следующим образом: val (0, 8, 8) = a * val (1,8,8).

Поскольку может быть любого вещественного числа, я не думаю, что текстура + пробоотборники могут быть использован в данном случае. Кроме того, метод должен работать и с обычными буферами.

Также, поскольку координата вокселя границы может быть либо частью угла сетки, края или грани, 26 (= 8 + 12 + 6), существуют разные варианты поиска ближайшего соседа (например, если координата была в (0,0,0), его ближайший сосед сказал, что сетка будет (1, 1, 1)). Таким образом, существует много потенциальных ветвлений.

Есть ли «элегантный» способ сделать это в OpenCL/CUDA? Кроме того, целесообразно ли обрабатывать границу с помощью отдельного ядра?

Answer 1

Самый обычный способ обработки границ в CUDA, чтобы проверить наличие всех возможных условий границы и действовать соответствующим образом, то есть:

• Если «этот элемент» выходит за пределы, то return (это очень полезно в CUDA, где вы, вероятно, запускаете больше потоков, чем это строго необходимо, поэтому дополнительные потоки должны выйти рано, чтобы избежать записи в памяти за пределами границ).
• Если «этот элемент» находится на/около левой границы (минимум x), то выполняйте специальные операции для левой границы.
• То же, что и справа, вверх, вниз (и спереди и сзади, в 3D).

К счастью, в большинстве случаев вы можете использовать max/min для упрощения этих операций, поэтому вы избегаете слишком много ifs. Я хотел бы использовать выражение этой формы:

source_pixel_x = max(0, min(thread_2D_pos.x + j, MAX_X)); 
source_pixel_y = ... // you get the idea 

В результате этих выражений всегда связана между 0 и некоторым MAX, тем самым зажимом out_of_bounds источником пикселей пограничных пикселями.

EDIT: Как прокомментировал DarkZeros, проще использовать (и меньше подвергать ошибкам) ​​использовать функцию clamp(). Мало того, что он проверяет как min, так и max, он также допускает типы векторов, такие как float3, и фиксирует каждый размер отдельно. См: clamp

Вот пример, который я сделал в качестве упражнения, 2D-Gaussian Blur:

__global__ 
void gaussian_blur(const unsigned char* const inputChannel, 
        unsigned char* const outputChannel, 
        int numRows, int numCols, 
        const float* const filter, const int filterWidth) 
{ 
    const int2 thread_2D_pos = make_int2(blockIdx.x * blockDim.x + threadIdx.x, 
             blockIdx.y * blockDim.y + threadIdx.y); 
    const int thread_1D_pos = thread_2D_pos.y * numCols + thread_2D_pos.x; 

    if (thread_2D_pos.x >= numCols || thread_2D_pos.y >= numRows) 
    { 
     return; // "this output pixel" is out-of-bounds. Do not compute 
    } 

    int j, k, jn, kn, filterIndex = 0; 
    float value = 0.0; 
    int2 pixel_2D_pos; 
    int pixel_1D_pos; 

    // Now we'll process input pixels. 
    // Note the use of max(0, min(thread_2D_pos.x + j, numCols-1)), 
    // which is a way to clamp the coordinates to the borders. 
    for(k = -filterWidth/2; k <= filterWidth/2; ++k) 
    { 
     pixel_2D_pos.y = max(0, min(thread_2D_pos.y + k, numRows-1)); 
     for(j = -filterWidth/2; j <= filterWidth/2; ++j,++filterIndex) 
     { 
      pixel_2D_pos.x = max(0, min(thread_2D_pos.x + j, numCols-1)); 
      pixel_1D_pos = pixel_2D_pos.y * numCols + pixel_2D_pos.x; 

      value += ((float)(inputChannel[pixel_1D_pos])) * filter[filterIndex]; 
     } 
    } 

    outputChannel[thread_1D_pos] = (unsigned char)value; 
} 

Answer 2

В OpenCL вы можете использовать Image3d для обработки своей 3d-сетки. Граничное обращение может быть achived с пробоотборником и определенным режимом адреса:

• CLK_ADDRESS_REPEAT - из-за границами диапазона координат изображения завернуты в допустимый диапазон. Этот режим адреса можно использовать только с нормализованными координатами. Если нормализованные координаты не используются, этот режим адресации может генерировать координаты изображения, которые не определены.
• CLK_ADDRESS_CLAMP_TO_EDGE - координаты изображения за пределами диапазона зажимаются.
• CLK_ADDRESS_CLAMP32 - координаты изображения вне диапазона возвращают цвет границы. Цвет рамки равен (0.0f, 0.0f, 0.0f, 0.0f), если порядок канала изображения CL_A, CL_INTENSITY, CL_RA, CL_ARGB, CL_BGRA или CL_RGBA и равен (0.0f, 0.0f, 0.0f, 1.0f), если изображение канал - CL_R, CL_RG, CL_RGB или CL_LUMINANCE.
• CLK_ADDRESS_NONE - для этого режима адресата программатор гарантирует, что координаты изображения, используемые для выборки элементов изображения, относятся к местоположению внутри изображения; в противном случае результаты не определены.

Кроме того, вы можете определить режим фильтрации для интерполяции (ближайший сосед или линейной).

Это соответствует вашим потребностям? В противном случае, пожалуйста, дайте нам более подробную информацию о ваших данных и их граничных требованиях.