Что означает "sizeof type T на моем устройстве (ах) CUDA"?

Question

Размеры типов на устройствах CUDA теоретически могут отличаться от их размеров на платформе хоста. Итак, каков идиоматический способ выражения "sizeof(T) на моем устройстве CUDA" в коде, кроме как сворачивание вашей собственной карты идентификаторов типов в значения, которые вы знаете?

Answer 1

На любой поддерживаемой в настоящее время поддерживаемой платформе CUDA ничего не требуется. Одна из причин, по которой набор инструментов CUDA так тесно интегрирован с компилятором хоста и библиотекой времени выполнения C++ хоста, заключается в том, что гарантируется, что размер базовых типов на хосте и устройстве всегда совпадает. Никакого идиоматического перевода размеров не требуется. Результат sizeof всегда будет одинаковым как для хоста, так и для устройства. Обратите внимание, что размер основных типов может различаться в зависимости от платформы (Windows - это платформа LLP64/IL32P64, linux и OS X - это платформы LP64/I32LP64), но это не влияет на GPU.

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

Например, рассмотрим следующий тривиальный пример кода:

#include <cstdio>

__device__ __host__ __noinline__ void printsizes(const char* title)
{
    printf("%s\n", title);
    printf("sizeof(void*) = %ld\n", (unsigned long)sizeof(void*));
    printf("sizeof(char) = %ld\n", (unsigned long)sizeof(char));
    printf("sizeof(bool) = %ld\n", (unsigned long)sizeof(bool));
    printf("sizeof(short) = %ld\n", (unsigned long)sizeof(short));
    printf("sizeof(int) = %ld\n", (unsigned long)sizeof(int));
    printf("sizeof(long) = %ld\n", (unsigned long)sizeof(long));
    printf("sizeof(long long) = %ld\n", (unsigned long)sizeof(long long));
}

__global__ void printkernel()
{
    printsizes("On the device:");
}

int main()
{
    printsizes("On the host:");

    printkernel<<<1,1>>>();
    cudaDeviceSynchronize();
    cudaDeviceReset();

    return 0;
}

Скомпилированный и запущенный на платформе Linux 64 дает следующее:

$ nvcc -arch=sm_52 -m64 -o sizeof64 sizeof.cu
$ ./sizeof64
On the host:
sizeof(void*) = 8
sizeof(char) = 1
sizeof(bool) = 1
sizeof(short) = 2
sizeof(int) = 4
sizeof(long) = 8
sizeof(long long) = 8
On the device:
sizeof(void*) = 8
sizeof(char) = 1
sizeof(bool) = 1
sizeof(short) = 2
sizeof(int) = 4
sizeof(long) = 8
sizeof(long long) = 8

Построенный на 64-битной платформе Windows, он дает следующее:

>nvcc -arch=sm_21 -m64 sizes.cu
sizes.cu
   Creating library a.lib and object a.exp
>a.exe
On the host:
sizeof(void*) = 8
sizeof(char) = 1
sizeof(bool) = 1
sizeof(short) = 2
sizeof(int) = 4
sizeof(long) = 4
sizeof(long long) = 8
On the device:
sizeof(void*) = 8
sizeof(char) = 1
sizeof(bool) = 1
sizeof(short) = 2
sizeof(int) = 4
sizeof(long) = 4
sizeof(long long) = 8

Построенный на 32-битной платформе Windows, он дает следующее:

>nvcc -arch=sm_21 -m32 sizes.cu
sizes.cu
   Creating library a.lib and object a.exp

C:\Users\david\Documents>a.exe
On the host:
sizeof(void*) = 4
sizeof(char) = 1
sizeof(bool) = 1
sizeof(short) = 2
sizeof(int) = 4
sizeof(long) = 4
sizeof(long long) = 8
On the device:
sizeof(void*) = 4
sizeof(char) = 1
sizeof(bool) = 1
sizeof(short) = 2
sizeof(int) = 4
sizeof(long) = 4
sizeof(long long) = 8

Обратите внимание, что размеры void * а также longможет варьироваться между платформами. Но в любом случае размеры графического процессора соответствуют размерам хоста. Это базовый принцип разработки драйвера CUDA и среды выполнения GPU.