Является ли 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), поэтому является самым медленным из всех. Для Kernel 2 я использую thrust::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);

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

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

Answer 1

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

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

Чтобы обойти пункт 2, кажется, может быть как минимум 2 возможных подхода:

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

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

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