РЕДАКТИРОВАТЬ: новый минимальный рабочий пример, чтобы проиллюстрировать вопрос и лучшее объяснение результатов nvvp (следующие предложения, приведенные в комментариях).
Итак, я создал «минимальный» рабочий пример, который следует:
#include <cuComplex.h>
#include <iostream>
int const n = 512 * 100;
typedef float real;
template < class T >
struct my_complex {
T x;
T y;
};
__global__ void set( my_complex< real > * a )
{
my_complex< real > & d = a[ blockIdx.x * 1024 + threadIdx.x ];
d = { 1.0f, 0.0f };
}
__global__ void duplicate_whole( my_complex< real > * a )
{
my_complex< real > & d = a[ blockIdx.x * 1024 + threadIdx.x ];
d = { 2.0f * d.x, 2.0f * d.y };
}
__global__ void duplicate_half( real * a )
{
real & d = a[ blockIdx.x * 1024 + threadIdx.x ];
d *= 2.0f;
}
int main()
{
my_complex< real > * a;
cudaMalloc( ( void * * ) & a, sizeof( my_complex< real > ) * n * 1024 );
set<<< n, 1024 >>>( a );
cudaDeviceSynchronize();
duplicate_whole<<< n, 1024 >>>( a );
cudaDeviceSynchronize();
duplicate_half<<< 2 * n, 1024 >>>( reinterpret_cast< real * >( a ) );
cudaDeviceSynchronize();
my_complex< real > * a_h = new my_complex< real >[ n * 1024 ];
cudaMemcpy( a_h, a, sizeof( my_complex< real > ) * n * 1024, cudaMemcpyDeviceToHost );
std::cout << "( " << a_h[ 0 ].x << ", " << a_h[ 0 ].y << " )" << '\t' << "( " << a_h[ n * 1024 - 1 ].x << ", " << a_h[ n * 1024 - 1 ].y << " )" << std::endl;
return 0;
}
Когда я компилирую и запускаю приведенный выше код, ядрам duplicate_whole и duplicate_half требуется примерно одинаковое время для запуска.
Однако, когда я анализирую ядра с помощью nvvp Iполучить разные отчеты для каждого из ядер в следующем смысле.Для ядра duplicate_whole nvvp предупреждает меня, что в строке 23 (d = { 2.0f * d.x, 2.0f * d.y };) ядро выполняет
Global Load L2 Transaction/Access = 8, Ideal Transaction/Access = 4
Я согласен, что загружаю 8-байтовые слова.Я не понимаю, почему 4 байта - это идеальный размер слова.В частности, нет никакой разницы в производительности между ядрами.
Я полагаю, что должны быть обстоятельства, когда этот шаблон доступа к глобальному хранилищу может привести к снижению производительности.Что это такое?
И почему я не получаю удар по производительности?
Я надеюсь, что это изменение прояснило некоторые неясные моменты.
+++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++
Я начну с некоторого кода ядра, чтобы проиллюстрировать мой вопрос, который будет следовать ниже
template < class data_t >
__global__ void chirp_factors_multiply( std::complex< data_t > const * chirp_factors,
std::complex< data_t > * data,
int M,
int row_length,
int b,
int i_0
)
{
#ifndef CUGALE_MUL_SHUFFLE
// Output array length:
int plane_area = row_length * M;
// Process element:
int i = blockIdx.x * row_length + threadIdx.x + i_0;
my_complex< data_t > const chirp_factor = ref_complex( chirp_factors[ i ] );
my_complex< data_t > datum;
my_complex< data_t > datum_new;
for ( int i_b = 0; i_b < b; ++ i_b )
{
my_complex< data_t > & ref_datum = ref_complex( data[ i_b * plane_area + i ] );
datum = ref_datum;
datum_new.x = datum.x * chirp_factor.x - datum.y * chirp_factor.y;
datum_new.y = datum.x * chirp_factor.y + datum.y * chirp_factor.x;
ref_datum = datum_new;
}
#else
// Output array length:
int plane_area = row_length * M;
// Element to process:
int i = blockIdx.x * row_length + ( threadIdx.x + i_0 ) / 2;
my_complex< data_t > const chirp_factor = ref_complex( chirp_factors[ i ] );
// Real and imaginary part of datum (not respectively for odd threads):
data_t datum_a;
data_t datum_b;
// Even TIDs will read data in regular order, odd TIDs will read data in inverted order:
int parity = ( threadIdx.x % 2 );
int shuffle_dir = 1 - 2 * parity;
int inwarp_tid = threadIdx.x % warpSize;
for ( int i_b = 0; i_b < b; ++ i_b )
{
int data_idx = i_b * plane_area + i;
datum_a = reinterpret_cast< data_t * >( data + data_idx )[ parity ];
datum_b = __shfl_sync( 0xFFFFFFFF, datum_a, inwarp_tid + shuffle_dir, warpSize );
// Even TIDs compute real part, odd TIDs compute imaginary part:
reinterpret_cast< data_t * >( data + data_idx )[ parity ] = datum_a * chirp_factor.x - shuffle_dir * datum_b * chirp_factor.y;
}
#endif // #ifndef CUGALE_MUL_SHUFFLE
}
Давайте рассмотрим случай, когда data_t является float, который ограничен по пропускной способности памяти.Как видно выше, существует две версии ядра: одна читает / записывает 8 байтов (целое комплексное число) на поток, а другая - читает / записывает 4 байта на поток, а затем перетасовывает результаты, так что сложный продукт получаетсявычислено правильно.
Причина, по которой я написал версию с использованием shuffle, заключается в том, что nvvp настаивал на том, что чтение 8 байт на поток не было лучшей идеей, поскольку этот шаблон доступа к памяти был бы неэффективным.Это так, хотя в обеих протестированных системах (GTX 1050 и GTX Titan Xp) пропускная способность памяти была очень близка к теоретическому максимуму.
Конечно, я знал, что улучшения не произойдет, и это действительноcase: для запуска обоих ядер требуется примерно одинаковое время.Итак, мой вопрос заключается в следующем:
Почему nvvp сообщает, что чтение 8 байтов будет менее эффективным, чем чтение 4 байтов на поток?При каких обстоятельствах это имело бы место?
В качестве примечания, одинарная точность важнее для меня, но двойная полезна и в некоторых случаях.Интересно, что в случае, когда data_t является двойным, между двумя версиями ядра также нет разницы во времени выполнения, даже если в этом случае ядро привязано к вычислениям, а версия с произвольным воспроизведением выполняет несколько больше флопов, чем исходная версия.
Примечание: ядра применяются к набору данных row_length * M * b (b изображения with row_length столбцы и M строки), а массив chirp_factor равен row_length * M.Оба ядра работают отлично (я могу отредактировать вопрос, чтобы показать вам вызовы обеих версий, если у вас есть сомнения по этому поводу).