Как изменить код CUDA, чтобы получить 100% загрузку графического процессора
Как я могу изменить этот код, чтобы получить 100% загрузку моего графического процессора?
#include <iostream>
using namespace std;
__global__ void saxpy_parallel(int n, float a, float *x, float *y)
{
// Get the unique ID of this kernel instance
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n)
{
y[i] = a*x[i] + y[i];
}
}
int main(int argc, char const *argv[])
{
// Tensors length
int const n = 100;
// Define tensors
float x[n], y[n];
for (int i = 0; i < n; ++i)
{
x[i] = 1.0f*i;
y[i] = 1.0f*i;
}
// Device pointers
float *d_x, *d_y;
cudaMalloc(&d_x, n*sizeof(float));
cudaMalloc(&d_y, n*sizeof(float));
if (cudaMemcpy(d_x, &x, n*sizeof(float), cudaMemcpyHostToDevice) != cudaSuccess)
{
printf("Memory Error!\n");
return 0;
}
if (cudaMemcpy(d_y, &y, n*sizeof(float), cudaMemcpyHostToDevice) != cudaSuccess)
{
printf("Memory Error!\n");
return 0;
}
// Run the kernel
saxpy_parallel<<<4096, 512>>>(n, 2.0, d_x, d_y);
// Retrieve results from the device memory
cudaMemcpy(&y, d_y, n*sizeof(float), cudaMemcpyDeviceToHost);
cudaFree(d_y);
cudaFree(d_x);
printf("%s\n",y[0]);
system("PAUSE");
return 0;
}
Решение
Хорошо, давайте проигнорируем цель загрузки 100% графического процессора, поскольку она нереальна и не поддается измерению. Итак, давайте предположим, что вы хотите оптимизировать этот код, чтобы он работал быстрее. Какие рычаги нацелены? Ваш алгоритм очень прост, поэтому он не предоставляет много возможностей. Тем не менее, я вижу следующие цели
1) Размер блока
saxpy_parallel<<<4096, 512>>>
Если 512 — это лучшее число, я бы начал с 32 или 64 и удвоил бы размер при настройке запуска ядра, чтобы найти лучшее значение этого параметра.
2) Удалить ненужный код
if( i < n )
Оператор if может быть отброшен, если n всегда меньше, чем i. Это можно контролировать внешне для ядра. Может потребоваться заполнить массив нечетного размера, чтобы он был кратен размеру блока, чтобы заставить это работать.
3) Изучите использование векторных типов
CUDA имеет типы float2 и float4. Так что переработайте код, чтобы использовать любой из них, с НАДЕЖДОЙ, что будет более быстрый доступ к памяти благодаря меньшему количеству выборок и сохранений и арифметических операций, происходящих параллельно.
4) Разомкнуть петлю
Каждый поток в настоящее время выбирает один x, a и y. Попробуйте получить 2, 4 или 8 значений.
...
y[i] = a*x[i] + y[i];
y[i+1] = a*x[i+1] + y[i+1];
y[i+2] = a*x[i+2] + y[i+2];
y[i+3] = a*x[i+3] + y[i+3];
Это требует меньше потоков, но каждый поток делает больше работы. Попробуйте снять глушение с 2,4,6 или 8
ценности.
5) Сохраните результат в другой переменной
Передайте дополнительный параметр для результата. Затем перекодировать
__global__ void saxpy_parallel(int n, float a, float *x, float *y, float * b)
...
b[i] = a*x[i] + y[i];
Это тратит больше памяти за то, что не читает и не записывает в одно и то же место.
Не пробуя каждый из подходов по отдельности и не измеряя эффект до и после, вы не сможете определить, что работает. Тогда некоторые комбинации могут работать лучше
или хуже.
Попробуйте и получайте удовольствие и дайте нам знать!
Другие решения
Других решений пока нет …