почему nv cc удаляет мою ветку if во время компиляции?

Question

Я обнаружил странное поведение при компиляции кода cuda в ptx. Если глобальная функция, использующая возвращаемое значение из tex2DLod<uchar4>, вызывает функцию устройства с оператором if, чьи обе ветви содержат функцию устройства, использующую в качестве аргумента uchar4, результирующий файл ptx имеет только код из ветви else.

Пример здесь. Я скомпилировал следующий код, используя cuda 10.1 update 1 и update2. Результат всегда один и тот же. Когда я удаляю оператор if и помещаю туда только часть else. Результирующий ptx никогда не изменяется, что означает, что первая ветка потеряла.

#include <stdint.h>
#include <cuda.h>
__device__ float3 rgba2rgb(uchar4 p)
{
    return make_float3(p.x/255.0f, p.y/255.0f, p.z/255.0f);
}
__device__ float3 bgra2rgb(uchar4 p)
{
    return make_float3(p.z/255.0f, p.y/255.0f, p.x/255.0f);
}
__device__ float3 pixel2rgb(uchar4 p, bool flag)
{
    if(flag)
    {
        return bgra2rgb(p);
    }
    else
    {
        return rgba2rgb(p);
    }
}

extern "C" __global__ void func2(
    CUtexObject rgb_mip_texture,
    size_t width, size_t height,
    bool flag
)
{
    size_t x_p = blockIdx.x * blockDim.x + threadIdx.x;
    size_t y_p = blockIdx.y * blockDim.y + threadIdx.y;


    if (x_p >= width || y_p >= height)
        return;
    uchar4 pixel = tex2DLod<uchar4>(rgb_mip_texture, x_p, y_p, (float)0);
    //uchar4 pixel = make_uchar4(1, 2, 3, 4);
    float3 rgb = pixel2rgb(pixel, flag);
    printf("rgb=(%f,%f,%f)", rgb.x, rgb.y, rgb.z);
}

ccbin команды nv cc: clang 8.0.

/usr/bin/nvcc -ptx \
    -v --ptxas-options=-v \
    --compiler-options "-v" \
    -ccbin "${ccbin}" \
    "${input_file}" \
    -o "${ptx_file}"

Если pixel не из tex2DLod (например, из make_uchar4), то обе ветви сохраняются. Это известная ошибка в nv cc?

Answer 1

Это может показаться ошибкой в ​​nv cc 10.1 (единственная версия, которую я тестировал). Похоже, что компилятор пытается автоматическое c встроенное расширение функций rgba2rgb и bgra2rgb как-то ломается, так что результат компиляции этого:

__device__ float3 pixel2rgb(uchar4 p, bool flag)
{
    if(flag)
    {
        return bgra2rgb(p);
    }
    else
    {
        return rgba2rgb(p);
    }
}

фактически таков:

__device__ float3 pixel2rgb(uchar4 p, bool flag)
{
    return rgba2rgb(p);
}

Это не относится к текстурам как таковым, потому что я могу воспроизвести проблему с этим чтением кода непосредственно из глобальной памяти:

#include <stdint.h>
#include <cuda.h>
#include <cstdio>

__device__ float3 rgba2rgb(uchar4 p)
{
    return make_float3(p.x/255.0f, p.y/255.0f, p.z/255.0f);
}
__device__ float3 bgra2rgb(uchar4 p)
{
    return make_float3(p.z/255.0f, p.y/255.0f, p.x/255.0f);
}
__device__ float3 pixel2rgb(uchar4 p, bool flag)
{
    if(flag)
    {
        return bgra2rgb(p);
    }
    else
    {
        return rgba2rgb(p);
    }
}

__global__ void func2(
    uchar4* pixels,
    size_t width, size_t height,
    bool flag
)
{
    size_t x_p = blockIdx.x * blockDim.x + threadIdx.x;
    size_t y_p = blockIdx.y * blockDim.y + threadIdx.y;

    if ((x_p < width) && (y_p < height)) {

    size_t idx = x_p * width + y_p;
    uchar4 pixel = pixels[idx];
    float3 rgb = pixel2rgb(pixel, flag);

    printf("flag=%d idx=%ld rgb=(%f,%f,%f)\n", flag, idx, rgb.x, rgb.y, rgb.z);
    }
}

int main()
{
    int width = 2, height = 2;
    uchar4* data;
    cudaMallocManaged(&data, width * height * sizeof(uchar4));

    data[0] = make_uchar4(1, 2, 3, 4);
    data[1] = make_uchar4(2, 3, 4, 5);
    data[2] = make_uchar4(3, 4, 5, 6);
    data[3] = make_uchar4(4, 5, 6, 7);

    dim3 bdim(2,2);
    func2<<<1, bdim>>>(data, width, height, true);
    cudaDeviceSynchronize();

    func2<<<1, bdim>>>(data, width, height, false);
    cudaDeviceSynchronize();

    cudaDeviceReset();

    return 0;
}

$ nvcc  -arch=sm_52 -o wangwang wangwang.cu 
$ ./wangwang 
flag=1 idx=0 rgb=(0.003922,0.007843,0.011765)
flag=1 idx=2 rgb=(0.011765,0.015686,0.019608)
flag=1 idx=1 rgb=(0.007843,0.011765,0.015686)
flag=1 idx=3 rgb=(0.015686,0.019608,0.023529)
flag=0 idx=0 rgb=(0.003922,0.007843,0.011765)
flag=0 idx=2 rgb=(0.011765,0.015686,0.019608)
flag=0 idx=1 rgb=(0.007843,0.011765,0.015686)
flag=0 idx=3 rgb=(0.015686,0.019608,0.023529)

Я предполагаю, что версия make_uchar4 Упоминание работает, потому что компилятор выполняет предварительное вычисление результатов из-за постоянных входных данных и полностью исключает код функции преобразования.

Играя, я смог исправить это, изменив код следующим образом:

__device__ __inline__ float3 rgba2rgb(uchar4 p)
{
    return make_float3(p.x/255.0f, p.y/255.0f, p.z/255.0f);
}
__device__ __inline__ float3 bgra2rgb(uchar4 p)
{
    return make_float3(p.z/255.0f, p.y/255.0f, p.x/255.0f);
}

Когда я делаю это, компиляция внедряет несколько мерцающих логик c во встроенное расширение PTX, которое генерирует:

    ld.global.v4.u8         {%rs2, %rs3, %rs4, %rs5}, [%rd10];
    and.b16         %rs8, %rs1, 255;   <---- %rs1 is the input bool
    setp.eq.s16     %p4, %rs8, 0;
    selp.b16        %rs9, %rs2, %rs4, %p4;
    and.b16         %rs10, %rs9, 255;
    selp.b16        %rs11, %rs4, %rs2, %p4;
    and.b16         %rs12, %rs11, 255;

, и все работает правильно (ваш пробег может отличаться):

$ nvcc  -arch=sm_52 -o wangwang wangwang.cu 
$ ./wangwang 
flag=1 idx=0 rgb=(0.011765,0.007843,0.003922)
flag=1 idx=2 rgb=(0.019608,0.015686,0.011765)
flag=1 idx=1 rgb=(0.015686,0.011765,0.007843)
flag=1 idx=3 rgb=(0.023529,0.019608,0.015686)
flag=0 idx=0 rgb=(0.003922,0.007843,0.011765)
flag=0 idx=2 rgb=(0.011765,0.015686,0.019608)
flag=0 idx=1 rgb=(0.007843,0.011765,0.015686)
flag=0 idx=3 rgb=(0.015686,0.019608,0.023529)

Я бы сообщил об этом как об ошибке в NVIDIA.