Программирование на графическом процессоре (5): эффективное использование общей памяти

искусственный интеллект

содержание

  • предисловие
  • Транспонирование матрицы ЦП
  • Реализация графического процессора
  • простое портирование
  • единый блок
  • tile
  • Расчет использования
  • shared memory
  • Наконец

предисловие

Я сравнивал ЦП и ГП в главе 3, и разница очень велика.На этот раз давайте взглянем на оптимизацию самого ГП, в основном на использование разделяемой памяти.

Транспонирование матрицы ЦП

В транспонировании матриц нет ничего сложного, реализовать его на процессоре просто:

#include <stdio.h>
#include <stdlib.h>
#include <sys/time.h>

#define LOG_
#define N 1024

/* 转置 */
void transposeCPU( float in[], float out[] )
{
	for ( int j = 0; j < N; j++ )
	{
		for ( int i = 0; i < N; i++ )
		{
			out[j * N + i] = in[i * N + j];
		}
	}
}


/* 打印矩阵 */
void logM( float m[] )
{
	for ( int i = 0; i < N; i++ )
	{
		for ( int j = 0; j < N; j++ )
		{
			printf( "%.1f ", m[i * N + j] );
		}
		printf( "\n" );
	}
}


int main()
{
	int	size	= N * N * sizeof(float);
	float	*in	= (float *) malloc( size );
	float	*out	= (float *) malloc( size );

	/* 矩阵赋值 */
	for ( int i = 0; i < N; ++i )
	{
		for ( int j = 0; j < N; ++j )
		{
			in[i * N + j] = i * N + j;
		}
	}

	struct timeval	start, end;
	double		timeuse;
	int		sum = 0;
	gettimeofday( &start, NULL );

	transposeCPU( in, out );

	gettimeofday( &end, NULL );
	timeuse = end.tv_sec - start.tv_sec + (end.tv_usec - start.tv_usec) / 1000000.0;
	printf( "Use Time: %fs\n", timeuse );

#ifdef LOG
	logM( in );
	printf( "\n" );
	logM( out );
#endif

	free( in );
	free( out );
	return(0);
}

Реализация графического процессора

простое портирование

Если ничего не рассматривается, просто портируйте код на GPU:

#include <stdio.h>
#include <stdlib.h>
#include <sys/time.h>

#define N 1024
#define LOG_

/* 转置 */
__global__ void transposeSerial( float in[], float out[] )
{
	for ( int j = 0; j < N; j++ )
		for ( int i = 0; i < N; i++ )
			out[j * N + i] = in[i * N + j];
}

/* 打印矩阵 */
void logM( float m[] ){...}

int main()
{
	int size = N * N * sizeof(float);

	float *in, *out;

	cudaMallocManaged( &in, size );
	cudaMallocManaged( &out, size );

	for ( int i = 0; i < N; ++i )
		for ( int j = 0; j < N; ++j )
			in[i * N + j] = i * N + j;

	struct timeval	start, end;
	double		timeuse;
	gettimeofday( &start, NULL );

	transposeSerial << < 1, 1 >> > (in, out);

	cudaDeviceSynchronize();

	gettimeofday( &end, NULL );
	timeuse = end.tv_sec - start.tv_sec + (end.tv_usec - start.tv_usec) / 1000000.0;
	printf( "Use Time: %fs\n", timeuse );


#ifdef LOG
	logM( in );
	printf( "\n" );
	logM( out );
#endif

	cudaFree( in );
	cudaFree( out );
}

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

耗时

единый блок

Один блок может открывать до 1024 потоков, здесь 1024 потока открываются для выполнения.

/* 转置 */
__global__ void transposeParallelPerRow( float in[], float out[] )
{
	int i = threadIdx.x;
	for ( int j = 0; j < N; j++ )
		out[j * N + i] = in[i * N + j];
}

int main()
{
    ...
    transposeParallelPerRow << < 1, N >> > (in, out);
    ...
}

Эффективность повышается немедленно, а затраты времени значительно сокращаются.

耗时

tile

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

/* 转置 */
__global__ void transposeParallelPerElement( float in[], float out[] )
{
	int i = blockIdx.x * K + threadIdx.x;
	/* column */
	int j = blockIdx.y * K + threadIdx.y;
	/* row */
	out[j * N + i] = in[i * N + j];
}

int main()
{
	...
	dim3 blocks( N / K, N / K );
	dim3 threads( K, K );

	...
	
	transposeParallelPerElement << < blocks, threads >> > (in, out);
	...
}

Это обычные операции графического процессора, но на самом деле их использование все еще ограничено.

耗时

Расчет использования

Утилизацию можно примерно посчитать, например, здесьMemory Clock rateиMemory Bus WidthЭто 900 МГц и 128 бит, поэтому пиковая скорость составляет 14,4 ГБ/с.

GPU参数

Предыдущее кратчайшее время составляет 0,001681 с. Объем данных составляет 1024 * 1024 * 4 (байт) * 2 (чтение и запись). Таким образом, это 4,65 ГБ / с. Коэффициент использования составляет 32%. Ошибка.

shared memory

Как это улучшить?Проблема в том что при чтении данных читается непрерывно.Один варп считывает 32 данных,которые могут работать синхронно,но при записи пишется в разбросанном виде,с большим размером шага.Это Это приводит к снижению эффективности, поэтому необходимо использовать общую память для транспонирования данных, чтобы запись была также непрерывной и эффективной.

/* 转置 */
__global__ void transposeParallelPerElementTiled( float in[], float out[] )
{
	int	in_corner_i	= blockIdx.x * K, in_corner_j = blockIdx.y * K;
	int	out_corner_i	= blockIdx.y * K, out_corner_j = blockIdx.x * K;

	int x = threadIdx.x, y = threadIdx.y;

	__shared__ float tile[K][K];

	tile[y][x] = in[(in_corner_i + x) + (in_corner_j + y) * N];
	__syncthreads();
	out[(out_corner_i + x) + (out_corner_j + y) * N] = tile[x][y];
}

int main()
{

	...
	dim3 blocks( N / K, N / K );
	dim3 threads( K, K );

	struct timeval	start, end;
	double		timeuse;
	gettimeofday( &start, NULL );

	transposeParallelPerElementTiled << < blocks, threads >> > (in, out);
	...

}

Это доводит коэффициент использования до 44%, что является пропуском.

耗时

Итак, чтобы разработать алгоритм в соответствии с архитектурой, просмотрите диаграмму архитектуры:

GPU存储架构

Наконец

Но 44% достигли проходной черты, то есть предстоит еще более глубокая работа по оптимизации.Это содержание также размещено в последующих статьях, и вы можете увидеть их в области комментариев, если у вас есть какие-либо комментарии. или предложения~

Категории