Одновременное выполнение двух процессов, совместно использующих Tesla K20

Question

У меня возникло странное поведение, когда я запускаю 2 экземпляра ядра для одновременного запуска при совместном использовании ресурсов графического процессора.

Я разработал ядро ​​CUDA, которое предназначено для работы в одном SM (многопроцессор), где потоки выполняют операцию несколько раз (с циклом).

Ядро готово создать только блок, поэтому использовать только один SM.

simple.cu

#include <cuda_runtime.h> 
#include <stdlib.h> 
#include <stdio.h> 
#include <helper_cuda.h> 
using namespace std; 

__global__ void increment(float *in, float *out) 
{ 
    int it=0, i = blockIdx.x * blockDim.x + threadIdx.x; 
    float a=0.8525852f; 

    for(it=0; it<99999999; it++) 
      out[i] += (in[i]+a)*a-(in[i]+a); 
} 

int main(int argc, char* argv[]) 
{ 
    int i; 
    int nBlocks = 1; 
    int threadsPerBlock = 1024; 
    float *A, *d_A, *d_B, *B; 
    size_t size=1024*13; 

    A = (float *) malloc(size * sizeof(float)); 
    B = (float *) malloc(size * sizeof(float)); 

    for(i=0;i<size;i++){ 
      A[i]=0.74; 
      B[i]=0.36; 
    } 

    cudaMalloc((void **) &d_A, size * sizeof(float)); 
    cudaMalloc((void **) &d_B, size * sizeof(float)); 

    cudaMemcpy(d_A, A, size, cudaMemcpyHostToDevice); 
    cudaMemcpy(d_B, B, size, cudaMemcpyHostToDevice); 

    increment<<<nBlocks,threadsPerBlock>>>(d_A, d_B); 

    cudaDeviceSynchronize(); 

    cudaMemcpy(B, d_B, size, cudaMemcpyDeviceToHost); 

    free(A); 
    free(B); 

    cudaFree(d_A); 
    cudaFree(d_B); 

    cudaDeviceReset(); 

    return (0); 
} 

Так что, если я исполню ядро:

time ./simple

я

real 0m36.659s user 0m4.033s sys 0m1.124s

В противном случае, если я выполнить два экземпляра:

time ./simple & time ./simple

я для каждого процесса:

real 1m12.417s user 0m29.494s sys 0m42.721s

real 1m12.440s user 0m36.387s sys 0m8.820s

Насколько мне известно, расстрелы должны работать одновременно длительностью один (около 36 секунд). Однако они продолжаются в два раза больше базового времени. Мы знаем, что GPU имеет 13 SM, каждый должен выполнить один блок, поэтому ядра создают только 1 блок.

Выполняются ли они в том же SM?

Не должны ли они работать одновременно в разных SM?

EDITED

Для того, чтобы сделать мне яснее я приложу профили одновременного выполнения, полученного от nvprof:

Profile, первый экземпляр

профиля, второй экземпляр

Теперь я хотел бы показать вам поведение одного и того же сценария, но одновременно выполнять два экземпляры matrixMul образца:

Profile, первый экземпляр

профиля, второй экземпляр

Как вы можете видеть, в первом сценарии, ядро ​​ждет другой, чтобы закончить. Хотя во втором сценарии (matrixMul) ядра из обоих контекстов запускаются одновременно.

спасибо.

Answer 1

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

Имеются технологии (MPS, Hyper-Q), которые могут делать то, что вы хотите, но способ, которым вы пытаетесь это сделать, не будет работать.

---

Редактировать, чтобы ответить на обновление в вашем вопросе:

Добавленный используя образец MatrixMul не показывает, что вы думаете, он делает пример. Это приложение запускает 300 коротких ядер и вычисляет показатель производительности в среднем за 300 запусков. Экран профилирования был настроен на очень грубое временное разрешение, так что он выглядит так, как будто у него есть один длинный запуск ядра, когда на самом деле это серия очень коротких ядер времени.

Чтобы проиллюстрировать это, рассмотрим следующее:

Это нормальное профилирование прогон для одного процесса MatrixMul работает на устройстве Kepler. Обратите внимание, что существует много отдельных ядер, работающих непосредственно друг за другом.

Это профилирование следы двух одновременных процессов MatrixMul, работающих на одной и той же Kepler устройства:

Следует отметить, что существуют пробелы в профиле следов каждого процесса, это где переключение контекста между происходят два процесса. Поведение идентично вашему первоначальному примеру, только при гораздо более тонкой детализации. Как неоднократно повторялось несколько разных людей в ходе этого обсуждения - CUDA не поддерживает множественные контексты на образце устройства одновременно с использованием стандартного API времени выполнения. Сервер MPS делает, разрешив это добавлением демона, который перенаправляет API с помощью большого общего внутреннего конвейера Hyper-Q, но вы не используете его и не имеете никакого отношения к результатам, которые вы указали в этом вопросе.