Избавьтесь от ожидания ожидания во время выполнения асинхронного потока CUDA

Я искал способ, как избавиться от занятого ожидания в потоке узла в следующем коде (не копируйте этот код, он только показывает идею моей проблемы, в ней много основных ошибок):

cudaStream_t steams[S_N];
for (int i = 0; i < S_N; i++) {
    cudaStreamCreate(streams[i]);
}
int sid = 0;
for (int d = 0; d < DATA_SIZE; d+=DATA_STEP) {
     while (true) {
         if (cudaStreamQuery(streams[sid])) == cudaSuccess) { //BUSY WAITING !!!!
             cudaMemcpyAssync(d_data, h_data + d, DATA_STEP, cudaMemcpyHostToDevice, streams[sid]);
             kernel<<<gridDim, blockDim, smSize streams[sid]>>>(d_data, DATA_STEP);
             break;
         }
         sid = ++sid % S_N;
     }

}

Есть ли способ простоя хост-потока и как-то дождаться завершения какого-либо потока, а затем подготовить и запустить другой поток?

РЕДАКТИРОВАТЬ: я добавил в код while(true), чтобы подчеркнуть занятость ожидания. Теперь я выполняю все потоки и проверяю, какой из них завершил запуск другого нового. cudaStreamSynchronize ожидает завершения определенного потока, но я хочу дождаться любого потока, который первым завершил работу.

EDIT2: я избавился от занятого ожидания следующим образом:

cudaStream_t steams[S_N];
for (int i = 0; i < S_N; i++) {
    cudaStreamCreate(streams[i]);
}
int sid = 0;
for (int d = 0; d < DATA_SIZE; d+=DATA_STEP) {
    cudaMemcpyAssync(d_data, h_data + d, DATA_STEP, cudaMemcpyHostToDevice, streams[sid]);
    kernel<<<gridDim, blockDim, smSize streams[sid]>>>(d_data, DATA_STEP);
    sid = ++sid % S_N;
}
for (int i = 0; i < S_N; i++) {
    cudaStreamSynchronize(streams[i]);
    cudaStreamDestroy(streams[i]);
}

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

Источник

5 ответов

Решение

Моя идея решить эту проблему - иметь один хост-поток на один поток. Этот хост-поток будет вызывать cudaStreamSynchronize для ожидания завершения потоковых команд. К сожалению, это невозможно в CUDA 3.2, так как он позволяет только одному потоку хоста иметь дело с одним контекстом CUDA, это означает, что один поток хоста на один графический процессор с поддержкой CUDA.

Надеемся, что в CUDA 4.0 это будет возможно: новости CUDA 4.0 RC

РЕДАКТИРОВАТЬ: я тестировал в CUDA 4.0 RC, используя открытый mp. Я создал один поток хоста для потока cuda. И это начало работать.

Источник

Реальный ответ - использовать cudaThreadSynchronize для ожидания завершения всех предыдущих запусков, cudaStreamSynchronize для ожидания завершения всех запусков в определенном потоке и cudaEventSynchronize для ожидания только определенного события в определенном потоке, который будет записан.

Однако вам необходимо понять, как работают потоки и синхронизация, прежде чем вы сможете использовать их в своем коде.

Что произойдет, если вы вообще не используете потоки? Рассмотрим следующий код:

kernel <<< gridDim, blockDim >>> (d_data, DATA_STEP);
host_func1();
cudaThreadSynchronize();
host_func2();

Ядро запускается, и хост переходит к одновременному выполнению host_func1 и ядра. Затем хост и устройство синхронизируются, то есть хост ожидает завершения работы ядра, прежде чем перейти к host_func2().

А что если у вас два разных ядра?

kernel1 <<<gridDim, blockDim >>> (d_data + d1, DATA_STEP);
kernel2 <<<gridDim, blockDim >>> (d_data + d2, DATA_STEP);

kernel1 запускается асинхронно! хост движется, и kernel2 запускается до завершения kernel1! однако kernel2 не будет работать до тех пор, пока не завершится kernel1, потому что они оба были запущены в потоке 0 (поток по умолчанию). Рассмотрим следующую альтернативу:

kernel1 <<<gridDim, blockDim>>> (d_data + d1, DATA_STEP);
cudaThreadSynchronize();
kernel2 <<<gridDim, blockDim>>> (d_data + d2, DATA_STEP);

Это абсолютно не нужно, потому что устройство уже синхронизирует ядра, запущенные в одном потоке.

Итак, я думаю, что функциональность, которую вы ищете, уже существует... потому что ядро всегда ожидает завершения предыдущих запусков в том же потоке, прежде чем запускать (даже если хост проходит). То есть, если вы хотите дождаться завершения любого предыдущего запуска, просто не используйте потоки. Этот код будет работать нормально:

for (int d = 0; d < DATA_SIZE; d+=DATA_STEP) {
    cudaMemcpyAsync(d_data, h_data + d, DATA_STEP, cudaMemcpyHostToDevice, 0);
    kernel<<<gridDim, blockDim, smSize, 0>>>(d_data, DATA_STEP);
 }

Теперь о потоках. Вы можете использовать потоки для управления одновременным выполнением устройства.

Думайте о потоке как о очереди. Вы можете помещать разные вызовы memcpy и запуск ядра в разные очереди. Тогда ядра в потоке 1 и запуски в потоке 2 асинхронны! Они могут быть выполнены одновременно или в любом порядке. Если вы хотите быть уверены, что на устройстве одновременно выполняется только одна memcpy/kernel, не используйте потоки. Точно так же, если вы хотите, чтобы ядра выполнялись в определенном порядке, не используйте потоки.

Тем не менее, имейте в виду, что все, что помещено в поток 1, выполняется по порядку, поэтому не беспокойтесь о синхронизации. Синхронизация предназначена для синхронизации вызовов хоста и устройства, а не двух разных вызовов устройства. Итак, если вы хотите запустить несколько ядер одновременно, поскольку они используют разную память устройства и не влияют друг на друга, используйте потоки. Что-то вроде...

cudaStream_t steams[S_N];
for (int i = 0; i < S_N; i++) {
    cudaStreamCreate(streams[i]);
}

int sid = 0;
for (int d = 0; d < DATA_SIZE; d+=DATA_STEP) {
     cudaMemcpyAsync(d_data, h_data + d, DATA_STEP, cudaMemcpyHostToDevice, streams[sid]);
     kernel<<<gridDim, blockDim, smSize streams[sid]>>>(d_data, DATA_STEP);
     sid = ++sid % S_N;
 }

Не требуется явная синхронизация устройства.

Источник

Есть: cudaEventRecord(event, stream) а также cudaEventSynchronize(event), Справочное руководство http://developer.download.nvidia.com/compute/cuda/3_2/toolkit/docs/CUDA_Toolkit_Reference_Manual.pdf содержит все подробности.

Редактировать: BTW потоки удобны для одновременного выполнения ядер и передачи памяти. Почему вы хотите сериализовать выполнение, ожидая завершения текущего потока?

Источник

Вместо cudaStreamQuery вы хотите cudaStreamSynchronize

int sid = 0;
for (int d = 0; d < DATA_SIZE; d+=DATA_STEP) {
     cudaStreamSynchronize(streams[sid]);
     cudaMemcpyAssync(d_data, h_data + d, DATA_STEP, cudaMemcpyHostToDevice, streams[sid]);
     kernel<<<gridDim, blockDim, smSize streams[sid]>>>(d_data, DATA_STEP);
     sid = ++sid % S_N;
}

(Вы также можете использовать cudaThreadSynchronize для ожидания запусков во всех потоках и событий с cudaEventSynchronize для более сложной синхронизации хоста / устройства.)

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

Источник

Вам нужно скопировать блок данных и выполнить ядро ​​на этом блоке данных в разных циклах for. Это будет более эффективным.

как это:

size = N*sizeof(float)/nStreams;

for (i=0; i<nStreams; i++){
offset = i*N/nStreams;
cudaMemcpyAsync(a_d+offset, a_h+offset, size, cudaMemcpyHostToDevice, stream[i]);
}


for (i=0; i<nStreams; i++){
offset = i*N/nStreams;
kernel<<<N(nThreads*nStreams), nThreads, 0, stream[i]>>> (a_d+offset);
}

Таким образом, копия памяти не должна ждать выполнения ядра предыдущего потока, и наоборот.

Источник
Другие вопросы по тегам