Я делаю следующие две операции:
a + b = AddResult
AddResult * a = MultiplyResult
В приведенной выше логике AddResult
является промежуточным результатом и используется в качестве входных данных в следующей операции умножения.
#define N 4096 // size of array
__global__ void add(const int* a, const int* b, int* c)
{
int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid < N)
{
c[tid] = a[tid] + b[tid];
}
}
__global__ void multiply(const int* a, const int* b, int* c)
{
int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid < N)
{
c[tid] = a[tid] * b[tid];
}
}
int main()
{
int T = 1024, B = 4; // threads per block and blocks per grid
int a[N], b[N], c[N], d[N], e[N];
int* dev_a, * dev_b, * dev_AddResult, * dev_Temp, * dev_MultiplyResult;
cudaMalloc((void**)&dev_a, N * sizeof(int));
cudaMalloc((void**)&dev_b, N * sizeof(int));
cudaMalloc((void**)&dev_AddResult, N * sizeof(int));
cudaMalloc((void**)&dev_Temp, N * sizeof(int));
cudaMalloc((void**)&dev_MultiplyResult, N * sizeof(int));
for (int i = 0; i < N; i++)
{
// load arrays with some numbers
a[i] = i;
b[i] = i * 1;
}
cudaMemcpy(dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(dev_b, b, N * sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(dev_AddResult, c, N * sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(dev_Temp, d, N * sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(dev_MultiplyResult, e, N * sizeof(int), cudaMemcpyHostToDevice);
//ADD
add << <B, T >> > (dev_a, dev_b, dev_AddResult);
cudaDeviceSynchronize();
//Multiply
cudaMemcpy(dev_Temp, dev_AddResult, N * sizeof(int), cudaMemcpyDeviceToDevice); //<---------DO I REALLY NEED THIS?
multiply << <B, T >> > (dev_a, dev_Temp, dev_MultiplyResult);
//multiply << <B, T >> > (dev_a, dev_AddResult, dev_MultiplyResult);
//Copy Final Results D to H
cudaMemcpy(e, dev_MultiplyResult, N * sizeof(int), cudaMemcpyDeviceToHost);
for (int i = 0; i < N; i++)
{
printf("(%d+%d)*%d=%d\n", a[i], b[i], a[i], e[i]);
}
// clean up
cudaFree(dev_a);
cudaFree(dev_b);
cudaFree(dev_AddResult);
cudaFree(dev_Temp);
cudaFree(dev_MultiplyResult);
return 0;
}
В приведенном выше примере кода я передаю результаты сложения (т. е. dev_AddResult
) в другой массив устройств (т. е. dev_Temp
) для выполнения операции умножения.
ВОПРОС: Поскольку массив результатов сложения (т. е. dev_AddResult
) уже находится на устройстве GPU, действительно ли мне нужно перенести его в другой массив? Я уже пытался запустить следующее ядро, напрямую предоставив dev_AddResult
в качестве входных данных, и это дало те же результаты. Есть ли риск, связанный с прямой передачей вывода одного ядра в качестве ввода следующего ядра? Какие-нибудь передовые практики, которым нужно следовать?
@RobertCrovella: я уже пытался, и это сработало нормально, но я просто хотел знать, действительно ли это правильный путь и есть ли какие-либо риски?
Да, для случая, который вы показали, вы можете использовать «выход» одного ядра в качестве «входа» для следующего без какого-либо копирования. Вы уже сделали это и подтвердили, что это работает, поэтому я не буду приводить примеры. Изменения в любом случае тривиальны - устраните промежуточную операцию cudaMemcpy
и используйте тот же указатель dev_AddResult
вместо указателя dev_Temp
при вызове ядра с несколькими вызовами.
Что касается «рисков», то мне ничего не известно о приведенном вами примере. Переходя от этого примера к, возможно, более общему использованию, вы хотели бы убедиться, что вычисления добавления выходных данных завершены, прежде чем использовать их где-то еще.
Ваш пример уже делает это избыточно, используя как минимум 2 механизма:
cudaDeviceSynchronize()
- это заставляет завершить ранее выпущенную работуТак что вам действительно не нужен cudaDeviceSynchronize()
в этом случае. Это ничего не «вредит» с точки зрения функциональности, но, вероятно, добавляет несколько микросекунд к общему времени выполнения.
В более общем случае, если вы выпустили ядро добавления и умножения в отдельные потоки, то CUDA не дает никаких гарантий порядка выполнения, даже если вы «выпустили» ядро умножения после ядра добавления.
В этом случае (не тот, который у вас есть здесь), если вам нужна операция умножения для использования ранее вычисленных результатов добавления, вам нужно будет каким-то образом обеспечить это (принудительно завершить добавление ядра перед ядром умножения). Вы уже показали один способ сделать это здесь, используя вызов синхронизации.
Вопрос относительно вашего последнего абзаца: в случае использования разных потоков, каковы другие способы, кроме cudaDeviceSynchronize()
, чтобы обеспечить синхронизацию, чтобы убедиться, что предыдущее ядро завершено перед запуском следующего ядра?
cudaStreamWaitEvent()
docs.nvidia.com/cuda/cuda-runtime-api/… на одном потоке и cudaEventRecord()
docs.nvidia.com/cuda/cuda-runtime-api/… на другом потоке приходят на ум для синхронизации двух потоков. Также cudaStreamSynchronize() docs.nvidia.com/cuda/cuda-runtime-api/… для синхронизации хоста и одного потока GPU.
Или — более сложно — Cuda Graphs developer.nvidia.com/blog/cuda-graphs или для конкретных случаев использования Dynamic Parallelism developer.nvidia.com/blog/…
@Sebastian: Спасибо за предложение о графиках CUDA. Я изучаю это и выглядит довольно многообещающе для моей цели. Мне просто интересно, почему в моих трех разных книгах по CUDA вообще ничего не упоминается о графиках CUDA.
Графики Cuda (задачи) — это расширенная функция, которая была представлена около 3 лет назад в CUDA 10.
да, вы можете использовать «выход» одного ядра в качестве «входа» для следующего. Почему бы просто не попробовать?