У меня есть программа, работающая на графическом процессоре, использующая CUDA с большим количеством маленьких ядер, а это означает, что для вызова ядра на моем процессоре требуется примерно столько же времени, сколько для запуска ядра на моем графическом процессоре.
Я хотел бы добавить функцию CPU в цикл моей программы, который требует примерно столько же времени, сколько одна итерация всех моих ядер.Я знаю, что после запуска ядра процессор может работать асинхронно с графическим процессором, но поскольку мой последний запуск ядра не намного опережает выполняемую работу с графическим процессором, в данном случае это не вариант.
Итак, моя идея состояла в том, чтобы использовать несколько потоков: один поток для запуска моих ядер GPU, а другой (или несколько других) для выполнения функции ЦП и параллельного запуска этих двух.
Я создал небольшой пример для проверки этой идеи:
#include <unistd.h>
#include <cuda_runtime.h>
#include <cuda_profiler_api.h>
#define THREADS_PER_BLOCK 64
__global__ void k_dummykernel1(const float* a, const float* b, float* c, const int N)
{
const int id = blockIdx.x * blockDim.x + threadIdx.x;
if(id < N)
{
float ai = a[id];
float bi = b[id];
c[id] = powf(expf(bi*sinf(ai)),1.0/bi);
}
}
__global__ void k_dummykernel2(const float* a, const float* b, float* c, const int N)
{
const int id = blockIdx.x * blockDim.x + threadIdx.x;
if(id < N)
{
float bi = b[id];
c[id] = powf(c[id],bi);
}
}
__global__ void k_dummykernel3(const float* a, const float* b, float* c, const int N)
{
const int id = blockIdx.x * blockDim.x + threadIdx.x;
if(id < N)
{
float bi = b[id];
c[id] = logf(c[id])/bi;
}
}
__global__ void k_dummykernel4(const float* a, const float* b, float* c, const int N)
{
const int id = blockIdx.x * blockDim.x + threadIdx.x;
if(id < N)
{
c[id] = asinf(c[id]);
}
}
int main()
{
int N = 10000;
int N2 = N/5;
float *a = new float[N];
float *b = new float[N];
float *c = new float[N];
float *d_a,*d_b,*d_c;
for(int i = 0; i < N; i++)
{
a[i] = (10*(1+i))/(float)N;
b[i] = (i+1)/50.0;
}
cudaMalloc((void**)&d_a,N*sizeof(float));
cudaMalloc((void**)&d_b,N*sizeof(float));
cudaMalloc((void**)&d_c,N*sizeof(float));
cudaMemcpy(d_a, a ,N*sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_b, b ,N*sizeof(float), cudaMemcpyHostToDevice);
cudaProfilerStart();
for(int k = 0; k < 100; k++)
{
k_dummykernel1<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
k_dummykernel2<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
k_dummykernel3<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
k_dummykernel4<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
k_dummykernel1<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
k_dummykernel2<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
k_dummykernel3<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
k_dummykernel4<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
for(int i = 0; i < N2; i++)
{
c[i] = pow(a[i],b[i]);
}
}
cudaDeviceSynchronize();
usleep(40000);
for(int k = 0; k <= 100; k++)
{
#pragma omp parallel sections num_threads(2)
{
#pragma omp section
{
k_dummykernel1<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
k_dummykernel2<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
k_dummykernel3<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
k_dummykernel4<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
k_dummykernel1<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
k_dummykernel2<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
k_dummykernel3<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
k_dummykernel4<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
}
#pragma omp section
{
for(int i = 0; i < N2; i++)
{
c[i] = pow(a[i],b[i]);
}
}
}
}
cudaDeviceSynchronize();
cudaProfilerStop();
delete[] a;
delete[] b;
delete[] c;
cudaFree((void*)d_a);
cudaFree((void*)d_b);
cudaFree((void*)d_c);
}
Я компилирую, используя: nvcc main.cu -O3 -Xcompiler -fopenmp
Сначала я запускаю 2x4 ядра и последовательные вычисления ЦП, а после этого я пыталсяделайте это параллельно, используя секции OpenMP.
Вот результат в профилировщике: 
Параллельная версия намного медленнее, чем последовательная ...
Если увеличить масштаб последовательной части, это будет выглядеть так: 
Можно видеть, что между каждыми 8 запусками ядра существует промежуток, гдеВычисления процессора выполнены (это я хотел бы закрыть, перекрыв его вызовами ядра).
Если я увеличу параллельную часть (тот же уровень масштабирования!), это будет выглядеть так: 
Пробелов больше нет, кроме кернеДля запуска теперь требуется около 15 микросекунд (против 5 микросекунд раньше).
Я также пробовал использовать большие размеры массивов и std::thread вместо OpenMP, но проблема всегда та же, что и раньше.
Может кто-тоскажите, если это вообще возможно, чтобы добраться до работы и если да, что я делаю не так?
Заранее спасибо
Кошка