Эмпирическое определение количества потоков в основе

Можно ли написать ядро CUDA, которое показывает, сколько потоков находится в warp без использования каких-либо связанных с Warp функций устройства CUDA и без использования бенчмаркинга? Если да, то как?

2 ответа

Поскольку вы указали, что решение с атомикой было бы интересно, я выдвигаю это как то, что, как я считаю, дает ответ, но я не уверен, что это обязательно ответ, который вы ищете. Я признаю, что он носит статистический характер. Я предоставляю это только потому, что нашел интересный вопрос. Я не предполагаю, что это "правильный" ответ, и я подозреваю, что кто-то умный придумает "лучший" ответ. Однако это может дать некоторые идеи.

Чтобы не использовать ничего, что явно ссылается на перекосы, я считаю, что необходимо сосредоточиться на "неявном" синхронном поведении. Сначала я пошел по пути, думая о том, как использовать конструкцию if-then-else (которая имеет некоторые синхронные последствия), но боролась с этим и вместо этого придумала такой подход:

#include <stdio.h>
#define LOOPS 100000

__device__ volatile int test2 = 0;
__device__ int test3 = 32767;

__global__ void kernel(){

 for (int i = 0; i < LOOPS; i++){
 unsigned long time = clock64();
// while (clock64() < (time + (threadIdx.x * 1000)));
 int start = test2;
 atomicAdd((int *)&test2, 1);
 int end = test2;
 int diff = end - start;
 atomicMin(&test3, diff);
 }
}

int main() {

 kernel<<<1, 1024>>>();
 int result;
 cudaMemcpyFromSymbol(&result, test3, sizeof(int));
 printf("result = %d threads\n", result);
 return 0;
}
</stdio.h>

Я компилирую с:

nvcc -O3 -arch=sm_20 -o t331 t331.cu

Я называю это "статистическим", потому что он запрашивает большое количество итераций (LOOPS) для получения правильной оценки (32). По мере уменьшения количества итераций "оценка" увеличивается.

Мы можем применить дополнительное синхронное плечо с помощью деформации, раскомментируя строку, которая закомментирована в ядре. Для моего тестового примера *, если эта строка раскоментирована, оценка правильная, даже если LOOPS= 1

* мой тестовый пример - CUDA 5, Quadro5000, RHEL 5.5


Вот несколько простых решений. Существуют и другие решения, которые используют синхронное программирование warp; однако многие решения не будут работать на всех устройствах.

РЕШЕНИЕ 1: Запустите один или несколько блоков с максимальным количеством потоков на блок, прочитайте специальные регистры % smid и % warpid, а также blockIdx и запишите значения в память. Группируйте данные по трем переменным, чтобы найти размер основы. Это еще проще, если вы ограничиваете запуск одним блоком, тогда вам нужно только% warpid.

РЕШЕНИЕ 2: Запустите один блок с максимальным количеством потоков на блок и прочитайте специальный регистр % clock. Для этого требуются следующие предположения, которые могут быть показаны на устройствах CC 1.0-3.5:

  • % clock определяется как неподписанный 32-разрядный счетчик циклов только для чтения, который обертывается молча и обновляет каждый цикл
  • все потоки в warp читают одно и то же значение для% clock
  • из-за задержек запуска warp и сборок выборки команд на одном и том же SM, но разные планировщики warp не могут выдавать первую инструкцию деформации в одном цикле

Все потоки в блоке, имеющие одинаковое время на устройствах CC1.0 - 3.5 (может измениться в будущем), будут иметь одинаковое время.

РЕШЕНИЕ 3: Используйте отладчик Nsight VSE или cuda-gdb. Представления состояния основы показывают вам достаточную информацию для определения размера основы. Также возможно сделать один шаг и увидеть изменение на адрес ПК для каждого потока.

РЕШЕНИЕ 4: Используйте Nsight VSE, Visual Profiler, nvprof и т.д. Запустите ядра из 1 блока с увеличением количества потоков на запуск. Определите, когда количество потоков, вызывающих warps_launched, будет идти от 1 до 2.

licensed under cc by-sa 3.0 with attribution.