Блог

Я изучаю возможность перехода с OpenCL на CUDA и провел несколько тестов, чтобы сравнить скорость CUDA в различных реализациях. К моему удивлению, в приведенных ниже примерах реализация PyCUDA выполняется примерно на 20% быстрее, чем в примере C CUDA.

Я читал много сообщений, в которых говорилось о «сборке релиза» кода C CUDA. Я попробовал -Xptxas -O3 использовать makefile, и это действительно ничего не изменило. Я также попытался настроить размер блока, с которым выполнялось ядро. К сожалению, это также не помогло повысить скорость.

Мои вопросы здесь:

• Какие могут быть причины, приводящие к разнице в скорости между C CUDA и PYCUDA?
• Если «продвинутая» (отсутствие лучшего слова) компиляция в PYCUDA является одной из причин, как я могу оптимизировать компиляцию моего кода C CUDA?
• Существуют ли какие-либо другие способы повысить скорость C CUDA в этом случае?

Хотя я ценю общие замечания, я ищу практические предложения, которые я могу проверить на своей машине. Спасибо!

 import pycuda.autoinit
import pycuda.driver as drv
import numpy as np

from pycuda.compiler import SourceModule
import time


mod = SourceModule(
    """
__global__ void saxpy(int n, const float a, float *x, float *y)
{
    int i = blockIdx.x * blockDim.x   threadIdx.x;
    if (i < n){
        y[i] = a * x[i]   y[i];
    }
}
"""
)

saxpy = mod.get_function("saxpy")

N = 1 << 25
time_elapse = 0.0

for i in range(100):
    # print(i)
    # print(N)

    x = np.ones(N).astype(np.float32)
    y = 2 * np.ones(N).astype(np.float32)
    start = time.time()
    saxpy(
        np.int32(N),
        np.float32(2.0),
        drv.In(x),
        drv.InOut(y),
        block=(512, 1, 1),
        grid=(int(N / 512)   1, 1),
    )
    time_elapse  = (time.time() - start)


print(time_elapse )
print(y[-100:-1])
print(y.sum())
print(N * 4.0)


 

 #include <stdio.h>
#include <time.h>
#define DIM 512



__global__ void saxpy(int n, float a, float *x, float *y)
{
    int i = blockIdx.x * blockDim.x   threadIdx.x;
    if (i < n)
        y[i] = a * x[i]   y[i];
}

int main(int num_iterations)
{
    double start;
    double cputime;
    int N = 1 << 25;
    float *x, *y, *d_x, *d_y;
    int i, j;
    for (j = 0; j < num_iterations; j  )
    {
        x = (float *)malloc(N * sizeof(float));
        y = (float *)malloc(N * sizeof(float));

        cudaMalloc(amp;d_x, N * sizeof(float));
        cudaMalloc(amp;d_y, N * sizeof(float));

        for (i = 0; i < N; i  )
        {
            x[i] = 1.0f;
            y[i] = 2.0f;
        }

        cudaMemcpy(d_x, x, N * sizeof(float), cudaMemcpyHostToDevice);
        cudaMemcpy(d_y, y, N * sizeof(float), cudaMemcpyHostToDevice);

        // Perform SAXPY on 1M elements
        start = clock();
        saxpy<<<(N   DIM) / DIM, DIM>>>(N, 2.0f, d_x, d_y);
        cputime  = ((double)(clock() - start) / CLOCKS_PER_SEC);
        cudaMemcpy(y, d_y, N * sizeof(float), cudaMemcpyDeviceToHost);

        // float maxError = 0.0f;
        // for (int i = 0; i < N; i  ){
        //     maxError = max(maxError, abs(y[i] - 4.0f));
        //     //printf("y[%d]: %fn", i,y[i]);
        // }
        // printf("Max error: %fn", maxError);

        cudaFree(d_x);
        cudaFree(d_y);
        free(x);
        free(y);
    }

 
    printf("cpu time is %fn", cputime);
    return 0;
}
 

Я сохранил приведенный выше файл как cuda_example.cu и скомпилировал его с помощью следующих команд в makefile :

 nvcc -arch=sm_61 -Xptxas -O3,-v -o main cuda_example.cu

 

Если я выполню ваш код на CUDA-C как есть, и установлю num_iterations равным 300 следующим образом:

 int num_iterations =300;
 

затем выполнение вашей программы занимает около 60 секунд на Geforce GTX 1650. Ваш код крайне неэффективен, поскольку вы копируете данные взад и вперед между графическим процессором и устройством на каждой итерации.
Итак, давайте ограничим цикл только выполнением ядра:

 #include <stdio.h>
#include <time.h>
#define DIM 512

__global__ void saxpy(int n, float a, float *x, float *y)
{
int i = blockIdx.x * blockDim.x   threadIdx.x;
if (i < n)
    y[i] = a * x[i]   y[i];
}

int main()
{
double start = clock();
int N = 1 << 25;
float *x, *y, *d_x, *d_y;
int i, j;

int num_iterations = 300;
x = (float *)malloc(N * sizeof(float));
y = (float *)malloc(N * sizeof(float));

cudaMalloc(amp;d_x, N * sizeof(float));
cudaMalloc(amp;d_y, N * sizeof(float));

for (i = 0; i < N; i  )
{
   x[i] = 1.0f;
   y[i] = 2.0f;
}
cudaMemcpy(d_x, x, N * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_y, y, N * sizeof(float), cudaMemcpyHostToDevice);

for (j = 0; j < num_iterations; j  ){
    saxpy<<<(N   DIM) / DIM, DIM>>>(N, 2.0f, d_x, d_y);
    cudaDeviceSynchronize();
}
cudaMemcpy(y, d_y, N * sizeof(float), cudaMemcpyDeviceToHost);
cudaFree(d_x);
cudaFree(d_y);
free(x);
free(y);

double cputime = ((double)(clock() - start) / CLOCKS_PER_SEC);
printf("cpu time is %fn", cputime);
return 0;
}
 

Если я это сделаю, то время выполнения составит 1,36 секунды. Выполняя sth, аналогичный коду PyCUDA, я получил около 19 секунд времени выполнения.