NVCC - разные размеры блоков в зависимости от арки во время компиляции

Question

У меня есть ядро, которое показывает максимальную производительность для другой Размеры блоков при запуске на оборудовании Kepler и Fermi. Я хотел бы, во время компиляции проверить текущую цель архитектуры и определить THREADS_PER_BLOCK макрос к i) запустить ядро; ii) определить количество необходимых блоков; iii) Статически установить общий размер памяти в ядре.

Ниже демонстрирует, что я пытаюсь сделать. Предположим, я нацелен на оборудование GK104 и, следовательно, использую nvcc -arch=sm_30. Анкет Это все равно приведет к THREADS_PER_BLOCK = 256 поскольку __CUDA_ARCH__ не определено для компиляции кода хоста. (Я понимаю, из EG этот ответ, почему это не может работать таким образом.)

#if __CUDA_ARCH__ >= 300
#define THREADS_PER_BLOCK 512
#else
#define THREADS_PER_BLOCK 256
#endif

__global__ void some_kernel(int* a, int* b) {
    __shared__ sm_data[THREADS_PER_BLOCK];
    // Do something.
}

int main(void) {
    // Initialize data.
    // Calculate blocks based on THREADS_PER_BLOCK, problem size and some max.
    some_kernel<<blocks, THREADS_PER_BLOCK>>>(d_a, d_b)
    return 0;
}

Я мог бы проверить свойства устройства во время выполнения и использовать динамическую общую память, но хотел бы знать, можно ли это жестко кодировать во время компиляции без необходимости вручную добавлять -dfermi или -dkepler и настройки THREADS_PER_BLOCK на основе этого. NB: Любые пользователи этого кода будут компилироваться сами, почти Конечно, для одной архитектуры, так что это не необоснованный вариант. Это просто кажется излишним в свете прохождения -arch= флаг.

Answer 1

nvcc Компилятор не обнаруживает локально доступных графических процессоров, он всегда нацелен на SM 1.0 по умолчанию. В противном случае это может ввести довольно запутанное поведение при создании различных систем.

Чтобы скомпилировать доступное устройство, вам нужно либо попросить пользователя указать версию SM или запустить некоторый код обнаружения во время сборки. Я убежден, что легче поместить код обнаружения оборудования в среду выполнения, а затем настроить запуск вашего ядра по желанию.