Действительно ли sort_by_key вызывает блокирующий вызов?

Question

я неоднократно епдиеие последовательность ядер:

for 1..100: 
    for 1..10000: 
     // Enqueue GPU kernels 
     Kernel 1 - update each element of array 
     Kernel 2 - sort array 
     Kernel 3 - operate on array 
    end 
    // run some CPU code 
    output "Waiting for GPU to finish" 
    // copy from device to host 
    cudaMemcpy ... D2H(array) 
end 

Ядро 3 имеет порядок O (N^2), так что на сегодняшний день самым медленным из всех. Для ядра 2 я использую упорную :: sort_by_key непосредственно на устройстве:

thrust::device_ptr<unsigned int> key(dKey); 
thrust::device_ptr<unsigned int> value(dValue); 
thrust::sort_by_key(key,key+N,value); 

кажется, что этот призыв к тяги блокировки, так как код CPU только получает выполняется один раз внутренний цикл завершен. Я вижу это, потому что, если я удалю вызов до sort_by_key, код хоста (правильно) выводит строку «Ожидание» до завершения внутреннего цикла, в то время как он не работает, если я запускаю сортировку.

Есть ли способ позвонить thrust::sort_by_key асинхронно?

Answer 1

1. Прежде всего, подумайте, что есть очередь запуска ядра, в которой может храниться столько ожидающих запусков. После того, как очередь запуска заполнена, запускаются дополнительные запуска ядра, любого типа блокируются. Хост-поток не будет выполняться (помимо тех запросов на запуск), пока не станут доступными слоты очереди. Я уверен, что 10000 итераций из 3 запусков ядра заполнит эту очередь, пока не достигнет 10000 итераций. Таким образом, будет какая-то латентность (я думаю) с любыми нетривиальными запусками ядра, если вы запускаете 30000 из них в последовательности. (в конце концов, когда все ядра добавляются в очередь, потому что некоторые из них уже завершены, тогда вы увидите сообщение «Ожидание ...», прежде чем все ядра завершатся, если не было другого поведения блокировки.)

2. thrust::sort_by_keyrequires temporary storage (размер примерно соответствует вашему набору данных размер). Это временное хранилище выделяется каждый раз, когда вы его используете, посредством операции cudaMalloc под капотом. Это cudaMalloc операция is блокировка. Когда cudaMalloc запускается из хост-потока, он ждет пробела в активности ядра, прежде чем он сможет продолжить.

Для работы вокруг пункта 2, кажется, что может быть по крайней мере 2 возможных подходов:

1. обеспечивают thrust custom allocator. В зависимости от характеристик этого распределителя вы можете устранить блокировку поведения cudaMalloc. (но см. обсуждение ниже)

2. Использование cub SortPairs. Преимущество здесь (как я вижу - ваш пример неполный) заключается в том, что вы можете сделать выделение один раз (при условии, что вы знаете наихудший временный размер хранилища во время итераций цикла) и исключить необходимость временного распределения памяти в вашем петля.

Метод тяги (1, выше), насколько я знаю, все равно будет эффективно сделать какое-то временное выделение/свободного шаг на каждую итерации, даже если поставить пользовательский распределитель. Если у вас хорошо спроектированный пользовательский распределитель, возможно, это почти «нет-op». Метод cub имеет недостаток в необходимости знать максимальный размер (чтобы полностью исключить необходимость выделения/свободного шага), но я утверждаю, что такое же требование будет иметь место для настраиваемого распределителя тяги. В противном случае, если вам нужно было выделить больше памяти в какой-то момент, пользовательскому распределителю фактически придется делать что-то вроде cudaMalloc, которое будет бросать ключ в работу.