Saya menemukan perilaku aneh ketika mengkompilasi kode cuda ke ptx. Jika fungsi global menggunakan nilai balik dari tex2DLod<uchar4> memanggil fungsi perangkat dengan pernyataan if yang kedua cabangnya berisi fungsi perangkat menggunakan uchar4 sebagai argumen, file ptx yang dihasilkan hanya memiliki kode dari cabang lain.

Contohnya ada di sini. Saya mengkompilasi kode berikut dengan cuda 10.1 update 1 dan update2. Hasilnya selalu sama. Ketika saya menghapus pernyataan if dan hanya meletakkan bagian lain di sana. Ptx yang dihasilkan tidak pernah berubah yang berarti cabang pertama telah hilang.

#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);
}

Perintah nvcc ccbin adalah clang 8.0.

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

Jika pixel bukan dari tex2DLod (misalnya dari make_uchar4), maka kedua cabang dipertahankan. Apakah ini bug yang dikenal di nvcc?

1
Wang 10 Maret 2020, 03:55

1 menjawab

Jawaban Terbaik

Ini tampaknya merupakan bug di nvcc 10.1 (satu-satunya versi yang saya uji). Tampaknya kompiler mencoba perluasan inline otomatis dari fungsi rgba2rgb dan bgra2rgb entah bagaimana rusak, sehingga hasil kompilasi ini:

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

Efektif ini:

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

Ini tidak terkait dengan tekstur saja, karena saya dapat mereproduksi masalah dengan pembacaan kode ini langsung dari memori global:

#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)

Saya berasumsi bahwa versi make_uchar4 yang Anda sebutkan berfungsi karena kompiler akan melakukan pra-perhitungan hasil karena input konstan dan menghilangkan kode fungsi konversi secara bersamaan.

Bermain-main, saya dapat memperbaikinya dengan mengubah kode seperti ini:

__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);
}

Ketika saya melakukan ini, kompilasi menyuntikkan beberapa logika swizzling ke dalam ekspansi PTX sebaris yang dihasilkannya:

    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;

Dan semuanya berfungsi dengan benar (jarak tempuh Anda mungkin berbeda):

$ 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)

Saya akan melaporkan ini sebagai bug ke NVIDIA.

2
2 revs 10 Maret 2020, 11:37