Мой вопрос является расширением Float4 не быстрее, чем float в cuda . На самом деле, я использую слегка модифицированную версию ядра, приведенную в связанном вопросе, чтобы добавить элементы двух массивов. Я не использую флаг -G, как предложено в ответе на связанный вопрос. Ядро float4 превосходит ядро float для массивов размером менее 2e + 6 (2M). Однако для больших массивов ядро float4 дает более низкую производительность, чем обычное ядро float.
Я пробовал блоки разных размеров и пробовал циклы с сеткой в сочетании с векторизованными загрузками / хранилищами float4. Ни один из этих методов не может превзойти простое ядро с плавающей точкой для больших массивов.
Мой вопрос состоит из двух частей:
- Почему ядро, использующее векторизованную (float4) загрузку / хранение медленнеечем обычное ядро, использующее скалярные (плавающие) загрузки / хранилища для больших массивов?
- Что можно сделать, чтобы повысить производительность по сравнению с простым плавающим ядром в контексте добавления элементов двух больших массивов?
Пример программы, представленной ниже, имеет четыре ядра:
- Простое ядро добавления с использованием загрузок / хранилищ float
- Простое ядро добавления с использованием загрузок / хранилищ float4
- Ядро сложенного добавления с использованием float загружает / хранит
- Ядро сложного добавления использует float4 загружает / сохраняет
Я тестировал следующий код на двух разных графических процессорах под CUDA 9.0.
- Nvidia GeForce 940M (ширина шины 64-битной памяти)
- Nvidia GeForce GTX 960 (ширина шины памяти 128)
Пример программы:
Редактировать В свете комментария @talonmies код был изменен таким образом, что вместо функции clock () в time.h теперь он использует события CUDA для целей синхронизации. Тем не менее, временные тренды остаются неизменными.
Редактировать Дальнейшие изменения кода, так что теперь он показывает время, затраченное каждым ядром, которое усредняется по 10 измерениям. Теперь простое ядро начинает превосходить другие ядра даже с меньшим размером элементов 2M.
#include <stdio.h>
#include <stdlib.h>
#include <time.h>
///////////////// Helper functions ////////////////////
void print_float_vec(const char *name, float *v, int n)
{
int i;
printf("%s\n", name);
for(i = n-5; i < n; i++)
printf("%.5f, ", v[i]);
printf("\n");
}
void init_float_vec(float *v, int n)
{
int i;
for(i = 0; i < n; i++)
v[i] = rand(); //(float)sqrt(float(i)); //or something else
}
//////////////////////////////////////////////////////////////////
////////////////// Kernels ///////////////////////////////////////
__global__ void add_float_simple(float *c, const float *a, const float *b, int size)
{
int i = (blockIdx.x * blockDim.x) + threadIdx.x;
if(i < size)
c[i] = a[i] + b[i];
}
__global__ void add_float4_simple(float4 *c, const float4 *a, const float4 *b, int size)
{
int i = blockDim.x * blockIdx.x + threadIdx.x;
if(i < size/4)
{
const float4 a1 = a[i];
const float4 b1 = b[i];
float4 c1;
c1.x = a1.x + b1.x;
c1.y = a1.y + b1.y;
c1.z = a1.z + b1.z;
c1.w = a1.w + b1.w;
c[i] = c1;
}
}
__global__ void add_float_strided(float *c, const float *a, const float *b, int size)
{
int i = (blockIdx.x * blockDim.x) + threadIdx.x;
for(; i < size; i += gridDim.x*blockDim.x)
{
c[i] = a[i] + b[i];
}
}
__global__ void add_float4_strided(float4 *c, const float4 *a, const float4 *b, int size)
{
int i = (blockIdx.x * blockDim.x) + threadIdx.x;
for(; i < size/4; i += gridDim.x*blockDim.x)
{
float4 a1 = a[i];
float4 b1 = b[i];
float4 c1;
c1.x = a1.x + b1.x;
c1.y = a1.y + b1.y;
c1.z = a1.z + b1.z;
c1.w = a1.w + b1.w;
c[i] = c1;
}
}
////////////////////////////////////////////////////////////
int main()
{
int size, ctr;
int BS = 128, NB1, NB4;
float gpu_time, total_time, avg_time;
cudaEvent_t t1, t2;
cudaEventCreate(&t1);
cudaEventCreate(&t2);
float *h_a, *h_b, *h_c;
float *d_a, *d_b, *d_c;
printf("Enter length of arrays ");
scanf("%d", &size);
NB1 = (int)ceil((float) size / BS);
NB4 = (int)ceil((float) (size / 4) / BS);
h_a = (float*) malloc (size * sizeof(float));
h_b = (float*) malloc (size * sizeof(float));
h_c = (float*) malloc (size * sizeof(float));
init_float_vec(h_a, size);
init_float_vec(h_b, size);
cudaMalloc( &d_a, size * sizeof(float));
cudaMalloc( &d_b, size * sizeof(float));
cudaMalloc( &d_c, size * sizeof(float));
cudaMemcpy( d_a, h_a, size * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy( d_b, h_b, size * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy( d_c, h_c, size * sizeof(float), cudaMemcpyHostToDevice);
/////////////////////////////////////////////////////////////////////////
total_time = 0.0;
for(ctr = 0; ctr < 10; ctr++)
{
cudaEventRecord(t1, 0);
add_float_simple << <NB1, BS>> >(d_c, d_a, d_b, size);
cudaEventRecord(t2, 0);
cudaEventSynchronize(t2);
cudaEventElapsedTime(&gpu_time, t1, t2);
total_time += gpu_time;
}
avg_time = total_time / 10.0;
printf("Simple float kernel average time = %f ms\n", avg_time);
//cudaMemcpy( h_c, d_c, size * sizeof(float), cudaMemcpyDeviceToHost);
//print_float_vec("c1. ", h_c, size);
//printf("\n\n");
/////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////
memset( h_c, 0, size * sizeof(float));
cudaMemset( d_c, 0, size * sizeof(float));
total_time = 0.0;
for(ctr = 0; ctr < 10; ctr++)
{
cudaEventRecord(t1, 0);
add_float4_simple << <NB4, BS>> >((float4*)d_c, (float4*)d_a, (float4*)d_b, size);
cudaEventRecord(t2, 0);
cudaEventSynchronize(t2);
cudaEventElapsedTime(&gpu_time, t1, t2);
total_time += gpu_time;
}
avg_time = total_time / 10.0;
printf("Simple float4 kernel average time = %f ms\n", avg_time);
//cudaMemcpy( h_c, d_c, size * sizeof(float), cudaMemcpyDeviceToHost);
//print_float_vec("c1. ", h_c, size);
//printf("\n\n");
/////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////
memset( h_c, 0, size * sizeof(float));
cudaMemset( d_c, 0, size * sizeof(float));
total_time = 0.0;
for(ctr = 0; ctr < 10; ctr++)
{
cudaEventRecord(t1, 0);
add_float_strided <<<128, 128>>> (d_c, d_a, d_b, size);
cudaEventRecord(t2, 0);
cudaEventSynchronize(t2);
cudaEventElapsedTime(&gpu_time, t1, t2);
total_time += gpu_time;
}
avg_time = total_time / 10.0;
printf("Strided float kernel average time = %f ms\n", avg_time);
//cudaMemcpy( h_c, d_c, size * sizeof(float), cudaMemcpyDeviceToHost);
//print_float_vec("c1. ", h_c, size);
//printf("\n\n");
/////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////
memset( h_c, 0, size * sizeof(float));
cudaMemset( d_c, 0, size * sizeof(float));
total_time = 0.0;
for(ctr = 0; ctr < 10; ctr++)
{
cudaEventRecord(t1, 0);
add_float4_strided <<<32, 128>>> ((float4*)d_c, (float4*)d_a, (float4*)d_b, size);
cudaEventRecord(t2, 0);
cudaEventSynchronize(t2);
cudaEventElapsedTime(&gpu_time, t1, t2);
total_time += gpu_time;
}
avg_time = total_time / 10.0;
printf("Strided float4 kernel average time = %f ms\n", avg_time);
//cudaMemcpy( h_c, d_c, size * sizeof(float), cudaMemcpyDeviceToHost);
//print_float_vec("c1. ", h_c, size);
//printf("\n\n");
/////////////////////////////////////////////////////////////////////////
}
Время, затраченное различными ядрами для массивов размером 500K, 1M, 2M, 4M, 5M, 6M и 8M, указанониже. Эти измерения относятся к устройству GeForce 940M, имеющему ширину шины 64 бита. Однако измерения, проведенные для GeForce GTX 960 с шириной шины 128 бит, также показали аналогичные тенденции. Я пробовал приведенный выше код с (64-битным) типом float2, но не смог получить другие результаты. Можно видеть, что для массивов, имеющих более 2M элементов, простое ядро с плавающей запятой начинает бить все остальные ядра.
Редактировать Измерения времени, полученные с использованием событий CUDA.
Редактировать Измерения времени усредняются за 10 прогонов каждого ядра
$ ./add_kernels_comp_avgd_time
Enter length of arrays 500000
Simple float kernel average time = 0.453997 ms
Simple float4 kernel average time = 0.446851 ms
Strided float kernel average time = 0.450666 ms
Strided float4 kernel average time = 0.476714 ms
$ ./add_kernels_comp_avgd_time
Enter length of arrays 1000000
Simple float kernel average time = 0.892458 ms
Simple float4 kernel average time = 0.888886 ms
Strided float kernel average time = 0.895814 ms
Strided float4 kernel average time = 0.902890 ms
Простое ядро с плавающей запятой начинает превосходить другие ядра для больших массивов, как показано ниже
$ ./add_kernels_comp_avgd_time
Enter length of arrays 2000000
Simple float kernel average time = 1.781978 ms
Simple float4 kernel average time = 1.803517 ms
Strided float kernel average time = 2.262067 ms
Strided float4 kernel average time = 1.865056 ms
$ ./add_kernels_comp_avgd_time
Enter length of arrays 4000000
Simple float kernel average time = 3.549984 ms
Simple float4 kernel average time = 4.322560 ms
Strided float kernel average time = 4.297559 ms
Strided float4 kernel average time = 3.912022 ms
$ ./add_kernels_comp_avgd_time
Enter length of arrays 6000000
Simple float kernel average time = 5.307305 ms
Simple float4 kernel average time = 6.263428 ms
Strided float kernel average time = 6.255444 ms
Strided float4 kernel average time = 5.995670 ms
$ ./add_kernels_comp_avgd_time
Enter length of arrays 8000000
Simple float kernel average time = 7.059725 ms
Simple float4 kernel average time = 8.264868 ms
Strided float kernel average time = 7.799684 ms
Strided float4 kernel average time = 7.764359 ms
Я будуБуду очень признателен, если кто-нибудь может помочь мне определить проблему, связанную с недостаточной производительностью ядер, использующих загрузки float4 и циклы с пошаговыми операциями для добавления больших массивов. Любое руководство по созданию ядра, способного превзойти простое ядро с плавающей запятой для больших массивов, будет наиболее ценно.
PS Я не хочу использовать CUBLAS / Thrust, потому что операция в моем реальном ядре не является дополнением. Кроме того, я хочу понять причину такой потери производительности, которую несут векторизованные нагрузки / хранилища при добавлении двух больших массивов.