Как изменить код CUDA, чтобы получить 100% нагрузку на графический процессор

Question

Как я могу изменить этот код, чтобы получить 100% нагрузку на мой GPU?

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

Answer 1

ОК, давайте проигнорируем цель загрузки 100% GPU как ее нереалистичную и нелегко измеримую. Поэтому позвольте предположить, что вы хотите оптимизировать этот код, чтобы работать быстрее. Каковы рычаги для прицеливания? Ваш алгоритм очень прост, поэтому он не предоставляет много возможностей. Тем не менее, я могу видеть последующие цели

1) Размер блока

saxpy_parallel<<<4096, 512>>> 

Является 512 лучшим числом, я хотел бы начать 32 или 64, и в два раза больше, как вы настроить запуск ядра, чтобы найти наилучшее значение этого параметра.

2) Удалите ненужный код

if(i < n) 

если заявление может быть отброшен, если п всегда меньше, чем я. Это может управляться извне к ядру. Может потребоваться, чтобы массив с нечетным размером был кратным размеру блока, чтобы заставить это работать.

3) изучить использование векторных типов

CUDA имеет float2 и Float4 типов. Поэтому переработайте код, чтобы использовать любой из них, с HOPE, что будет быстрее доступ к памяти через меньшее количество выборок и хранилищ и арифметические операции, происходящие параллельно.

4) помогает предотвратить заедание Loop

Каждый поток в настоящее время, забирающий один х, а и у. Попробуйте извлечь 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) Сохраните результат в другой переменной

Pass дополнительный параметр в течение результата. Затем переконфигурация

__global__ void saxpy_parallel(int n, float a, float *x, float *y, float * b) 

... 

    b[i] = a*x[i] + y[i]; 

Этот объект больше хранит для хранения и записи в том же месте.

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

Попробуйте и получайте удовольствие и дайте нам знать!