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

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

3 голосов
/ 24 февраля 2011

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

}

Есть ли способ приостановить поток хоста и как-то подождать, пока какой-нибудь поток не завершится, а затем подготовить и запустить другой поток?код, чтобы подчеркнуть занятость ожидания.Теперь я выполняю все потоки и проверяю, какой из них завершил запуск другого нового.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 ]

4 голосов
/ 25 февраля 2011

Реальный ответ - использовать 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;
 }

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

3 голосов
/ 01 марта 2011

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

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

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

3 голосов
/ 24 февраля 2011

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

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

1 голос
/ 25 февраля 2011

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

вот так: 

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);
}

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

1 голос
/ 24 февраля 2011

Вместо 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 и другим. По умолчанию, вероятно, то, что вы хотите.

...