Linux, OpenMP и GCC: многопоточная программа в 5-6 раз медленнее однопоточной?

Question

[Robotex]

Изучаю OpenMP, пытаюсь распараллелить нейронную сеть. Но почему-то постоянно получается, что расспараллеленное приложение существенно медленнее, чем обычное. В чем может быть проблема?


Однопоточное приложение:
pastebin.com/22vp3LYU

Компиляция:
g++ -lrt -O0 main.cpp -o nnlv2


Многопоточное:
pastebin.com/w3m01QTK

Компиляция:

g++ -lrt -fopenmp -O0 main_openmp.cpp -o nnlv2_openmp

P.S. Не обращайте внимания на функцию sigmoid, я в курсе, что она должна выглядеть по другому :)

Answer 1

[edeldm]

еще есть
cudaMemcpyToSymbol( «alpha», &alpha, 4);

бегло просмотрел — вроде верно все.
могу дать совет начинающего радиста — не крути две гайки одновременно.

т.е. попробуй сначала просто копировать на видюху и обратно и проверяй значения на совпадения.
затем многотредовое.

я делаю еще такое (поделюсь куском кода :) ):

__global__ void RunTest(unsigned long *heartbeat)
{
	unsigned long tid = blockIdx.x*blockDim.x + threadIdx.x;
	heartbeat[tid] = 1;
}

void Runkerneltest(int grids, int threads, unsigned long *heartbeat)
{
	RunTest <<<grids, threads>>> (heartbeat);
}




bool TestKernel(int grids, int threads)
{
	unsigned long* heartbeat =  new unsigned long [grids*threads];
	memset (heartbeat, 0, grids*threads*sizeof (unsigned long)); 
	void* heartbeat_device = 0;
	cudaError err;
	err = cudaMalloc ((void**)&heartbeat_device, grids*threads*sizeof (unsigned long));
	Check("cudaMalloc heartbeat_device",err);
	err = cudaMemset (heartbeat_device, 0, grids*threads*sizeof(unsigned long));
	Check("cudaMemset heartbeat_device",err);

	Runkerneltest(grids,threads,(unsigned long*)heartbeat_device);

	err = cudaThreadSynchronize();
	Check("cudaThreadSynchronize()", err);
	err = cudaMemcpy( heartbeat, heartbeat_device, grids*threads*sizeof(unsigned long), cudaMemcpyDeviceToHost );
	Check("cudaMemcpy( heartbeat, heartbeat_device)", err);

	bool error = false;
	for (int i=0; i<grids*threads; i++)
		if (heartbeat[i] != 1) 
		{
			LOG("tid %d test fails",i);
			error = true;
			break;
		}

	cudaFree (heartbeat_device);
	delete (heartbeat);
	return (error ? false : true); //если была ошибка, то false
}

Comments

[edeldm]

виноват, есть баг у тебя — ты не ждешь завершения работы ядер видюхи — надо делать
cudaThreadSynchronize();

— посмотри мой пример.
и на нем научись сначала.

[edeldm]

если спросишь нафиг такое надо — отвечу. что при больших размерностях гридов и тредов у тебя могут быть дырки — т.е. должны стоять 1 в ячейках матрицы а их там нет — там 0…
имхо для разных видюх — по разному — надо брать значения гридов и тредов — с инфы о видюхе автоматом.

[Robotex]

А как это автоматом делается?

Answer 2

[Sergey Lerg]

Используйте явные потоки. Например с помощью библиотеки pthread.

Comments

[Sergey Lerg]

И у Вас похоже не подключается библиотека openmp:
#include <omp.h>

[Sergey Lerg]

И распараллеливать лучше верхний цикл:
for(int i=0; i<LAYERS/2;i++)
for(int m=LAYERS/2; m<LAYERS;m++)

[Robotex]

Как распараллелить верхний, если следующий шаг зависит от предыдущего?

[Sergey Lerg]

Ой, не заметил.

Answer 3

[Robotex]

Оптимизировал код:

Однопоточная версия: pastebin.com/KAx4RmSJ
Многопоточная: pastebin.com/fbe4gZSn

Теперь многопоточная версия медленнее однопоточной всего в 2 раза (а нужно, чтобы минимум раза в 3 быстрее было). Что еще можно соптимизировать?

Comments

[Robotex]

Если вручную выставить один поток, то время выполнения одинаковое, если больше то начинает расти.

[Robotex]

У меня тоже двухядерник, но почему-то результат прямо противоположный. Сейчас попробую с O3

[Robotex]

Странно, с O3 тоже медленнее выходит.

Может в таком случае выполнить для меня одну просьбу? Можете написать время выполнения обеих версий, изменяя константу LAYERS от 10 до 100 (ставите 10, компилируете, запускаете, потом 20, компилируете, запускаете) и какая у вас конфигурация? Заранее благодарен.

[Robotex]

Спасибо.

[Robotex]

Если вручную выставить num_threads(4), то время вроде как стабильно себя ведет. Почему-то на двух потоках оно вот так вот скачет

[Robotex]

Попробуй с четырьмя.

[Robotex]

От же ж блин, что за день сегодня такой. Еще и в куде время в 35 раз большее (хотя использовал многократно оттестированный код). Может я время коряво замеряю?

[Robotex]

Можно попробовать.

[Robotex]

Со временем все впорядке :(

[Robotex]

блин, добился двойного прироста скорости в OpenMP, на следующий день опять то же самое (хотя исходник не менял)

Answer 4

[Robotex]

Здесь последние версии кода: ubuntuone.com/p/jPV/

Исходник с куда компилится так:
nvcc -lrt main_cuda.cu -o nnlv2_cuda

Все по прежнему печально. Куда мало что медленнее, так еще и считает не верно. Что я делаю не так?

Comments

[edeldm]

ты забыл сделать
cudaThreadSynchronize();!!!

Answer 5

[edeldm]

параметры видюхи «автоматом» берутся отсюда:

cudaDeviceProp prop;
if(cudaGetDeviceProperties( & prop, i) == cudaSuccess) 
{
	LOG ("Device: %s\n",prop.name);
	LOG ( "Compute capability     : %d.%d\n", prop.major, prop.minor );
	LOG ( "Name                   : %s\n", prop.name );
	LOG ( "Total Global Memory    : %ld\n", prop.totalGlobalMem );
	LOG ( "Shared memory per block: %d\n", prop.sharedMemPerBlock );
	LOG ( "Registers per block    : %d\n", prop.regsPerBlock );
	LOG ( "Warp size              : %d\n", prop.warpSize );
	LOG ( "Max Grid               : %d\n", prop.maxGridSize[0] );
	LOG ( "Max threads per block  : %d\n", prop.maxThreadsPerBlock );
	LOG ( "Total constant memory  : %d\n", prop.totalConstMem );
...
 }

Answer 6

[mikhanoid]

Так у вас же там стоит atomic перед k++. На каждой итерации нити будут синхронизовываться, чтобы друг за дружкой сделать k++. Естественно, оно будет сильно медленней.

Comments

[Robotex]

Эм, это устаревший код. В новом я убрал это, оставил только #pragma omp parallel for

Answer 7

[Robotex]

А что если ускорения на двухядерной машине и не должно быть? В одном ядре обрабатывается основной цикл, в другом параллельная зона.