OpenCL логическое выражение нежелательной ленивой оценки

Question

Из спецификации OpenCL 2.0, глава "6.3 Операторы", стр. 29:

г. Логические операторы и (&&), или (||) работают со всеми скалярными и векторными встроенными типами. Только для скалярных встроенных типов, и (&&) оценивает правый операнд только в том случае, если левый операнд сравнивается с 0. Только для скалярных встроенных типов, или (||) оценивает правый операнд, только если левый операнд сравнивается равным 0. Для встроенных векторных типов оба операнда оцениваются, а операторы применяются компонентно. Если один операнд является скаляром, а другой - вектором, скаляр может подвергаться обычному арифметическому преобразованию в тип элемента, используемый векторным операндом. Затем скалярный тип расширяется до вектора, который имеет то же количество компонентов, что и векторный операнд. Операция выполняется компонентно, в результате чего получается вектор того же размера.

Это означает, что использование выражений с логическими операторами приведет к ветвлению и расхождению потоков, что, в свою очередь, приведет к потере производительности на некоторых параллельных платформах. Пример:

int min_nonzero(int a, int b)
{
    return (a < b && a != 0)? a : b; // branch
}

Это может быть частично исправлено, как:

int min_nonzero(int a, int b)
{
    return select(b, a, a < b && a != 0); // branch because of &&
}

Где встроенная функция select вероятно, будет реализован с использованием арифметики, чтобы избежать ветвления (например, в виде линейной интерполяции). Но есть еще ветка в &&, Возможно лучший способ:

int min_nonzero(int a, int b)
{
    return select(b, a, (int)(a < b) & (int)(a != 0)); // branch free
}

Но это быстро становится нечитаемым.

Поэтому мой вопрос: есть ли лучший способ убедить компилятор OpenCL отказаться от ленивой оценки булевых выражений (не глобально, а в отдельных случаях) ?

---

Ниже мои практические эксперименты в этом вопросе, больше не вопрос. Ленивая оценка все еще требуется в некоторых случаях, таких как:

if(i < N && array[i] == x) // will go OOB without lazy evaluation

Поэтому маловероятно, что оптимизатор отключит его полностью или, по крайней мере, во всех применимых случаях.

Я смотрю на некоторые PTX, сгенерированные драйвером NVIDIA 320.49, и он будет оптимизировать только случаи без доступа к массиву справа:

if(p[i] == n_end && i)
    return;

компилируется в одну ветку:

    setp.ne.s32     %p2, %r17, %r5; // p[i] != n_end
    setp.eq.s32     %p3, %r28, 0; // !i
    or.pred     %p4, %p2, %p3; // (p[i] != n_end || !i) = !(p[i] == n_end && i)
    @%p4 bra    BB2_3; // branch
    ret;

BB2_3:

И все же это:

int n_increment = 1;
for(++ i; i < n_cols_B && p[i + 1] == n_end; ++ i)
    ++ n_increment;

компилируется в:

    mov.u32     %r29, 1;

BB2_4:
    mov.u32     %r6, %r28;
    add.s32     %r8, %r6, 1;
    ld.param.u32    %r24, [Fill_ColsTailFlags_v0_const_param_0];
    setp.ge.u32     %p5, %r8, %r24;
    @%p5 bra    BB2_6; // branch if i < n_cols_B

    shl.b32     %r19, %r6, 2;
    ld.param.u32    %r25, [Fill_ColsTailFlags_v0_const_param_1];
    add.s32     %r20, %r19, %r25;
    ld.const.u32    %r21, [%r20+8];
    setp.eq.s32     %p6, %r21, %r5;
    @%p6 bra    BB2_7; // branch if p[i + 1] == n_end

BB2_6:
    shl.b32     %r22, %r5, 2;
    ld.param.u32    %r27, [Fill_ColsTailFlags_v0_const_param_2];
    add.s32     %r23, %r27, %r22;
    st.global.u32   [%r23], %r29;
    ret;

BB2_7:
    add.s32     %r29, %r29, 1;
    mov.u32     %r28, %r8;
    bra.uni     BB2_4;

Похоже, что он стесняется доступа к массиву, поскольку он не знает, как анализировать правильность доступа к массиву по отношению к условию слева. Порядок переключения условий на p[i + 1] == n_end && i < n_cols_B избавляется от ветки в этом случае. Изменение индекса на постоянную i < n_cols_B && B_p[j] == n_end (где j = get_global_id(0) инициализируется в начале) не избавляется от ветки.