CUDA как правильно оптимизировать?

Question

[Типа Программист]

Есть кусок кода, как по мне он достаточно простой, но выполняется целых 5 миллисекунд, как-то это много, как для параллельных вычислений.

__global__ void kernel_compute_global_lighting(float* device_lenght_buff, CudaRenderWindow render_window, CudaRenderCamera camera, CudaRenderMap map, CudaRenderTextures textures, CudaRenderLight lights) {
	int pixel_coordinate_y = blockIdx.x * blockDim.x + threadIdx.x;
	int pixel_coordinate_x = blockIdx.y * blockDim.y + threadIdx.y;
	if (pixel_coordinate_y >= render_window.render_height || pixel_coordinate_x >= render_window.render_width)
		return;
	render_window.device_rendered_window[pixel_coordinate_y*render_window.render_width + pixel_coordinate_x].r = render_window.device_render_window[pixel_coordinate_y*render_window.render_width + pixel_coordinate_x].r * lights.device_light_pointers[0]->r;
	render_window.device_rendered_window[pixel_coordinate_y*render_window.render_width + pixel_coordinate_x].g = render_window.device_render_window[pixel_coordinate_y*render_window.render_width + pixel_coordinate_x].g * lights.device_light_pointers[0]->g;
	render_window.device_rendered_window[pixel_coordinate_y*render_window.render_width + pixel_coordinate_x].b = render_window.device_render_window[pixel_coordinate_y*render_window.render_width + pixel_coordinate_x].b * lights.device_light_pointers[0]->b;
}

Для того чтобы вызвать этот Kernel, использую такие параметры:
Сетка
X : (900 + 31) / 32
Y : (1400 + 31) / 32);
Потоки
X: 32
Y: 32

Пытался сжимать косвенную адресацию, делал промежуточный обьект для проведениче вычислений, чтоб сократить количество обращений к памяти. Но все что как по мне должно ускорять выполнение, наоборот делает задержку только больше.
Как мне тогда быть с изображениями выше 1024х800

Просто подскажите что именно тут больше всего влияет, устал рыть без результата...

Comments

[freeExec]

Сделайте пустое тело метода, за сколько выполнится?

[Типа Программист]

freeExec, Пустое выполняется за 0 миллисекунд.

Измеряю время выполнения так:

cudaDeviceSynchronize(); // Ожидание выполнения предыдущего кернела
	end = clock(); // остановка таймера от предыдущего кернала итд.
	cout << "Draw floor and roof-- " << ((double)end - start) / ((double)CLOCKS_PER_SEC) * 1000 << " , milliseconds." << endl;
	start = clock(); // запускаю таймер для кернала о котором идёт речь, ну и дальше понятно.
	kernel_compute_global_lighting << <grid, threads >> > (render_window->device_lenght_buff, *render_window, *render_window->camera, *render_window->map, *render_window->textures, *render_window->light);
	cudaDeviceSynchronize();
	end = clock();
	cout << "Global light--- " << ((double)end - start) / ((double)CLOCKS_PER_SEC) * 1000 << " , milliseconds." << endl;

[freeExec]

MegaCraZy6, А если оставить расчёт только одного канала. Ещё у nvidia была тулза для показа затыков при копировании и доступа к памяти.

[Типа Программист]

freeExec, Ля.
Забавно, переделал код, под одномерный массив, и полетела ласточка
0-1 миллисекунд.
:|

__global__ void kernel_compute_global_lighting(float* device_lenght_buff, CudaRenderWindow render_window, CudaRenderCamera camera, CudaRenderMap map, CudaRenderTextures textures, CudaRenderLight lights) {
	int addr = threadIdx.x + blockIdx.x * blockDim.x;
	render_window.device_rendered_window[addr].r = render_window.device_render_window[addr].r * lights.device_light_pointers[0]->r;
	render_window.device_rendered_window[addr].g = render_window.device_render_window[addr].g * lights.device_light_pointers[0]->g;
	render_window.device_rendered_window[addr].b = render_window.device_render_window[addr].b * lights.device_light_pointers[0]->b;
}

[freeExec]

MegaCraZy6, Так то в этом главная проблема, для производительности ему требуется, чтобы одному варпу нужны были данные в пределах 256 байт. Тогда он их запросит за одни раз. А когда у тебя двумерный массив, то данные разбросаны и основное время простоя это просто ожидание данных.

[Типа Программист]

freeExec, На самом деле большое спасибо, у меня с этим есть проблемы во всем коде, я читал куча статей, и там столько этих правил оптимизации, немного голова кругом идёт. Спасибо, пожалуй пока-что везде буду одномерный использовать).

[freeExec]

MegaCraZy6, Можно и двумерный, если он достаточно большой и запускать треды построчно.

Answer 1

[Сергей Тихонов]

1. Во-первых, стоит удостовериться, что Вы измерили именно скорость выполнения кода, а не загрузки данных, выполнения и выгрузки результата.
   
2. Во-вторых, попробуйте избавиться от if совсем (можно добавить неиспользуемые поля до размера блока).
   
3. В-третьих, используйте векторное умножение вместо отдельных операций в три строки.
   
4. В-четвертых, посмотрите размеры блока для вашей карты, может не влезаете.
   

Comments

[freeExec]

В данном случае if не проблема, потому, что нет второй ветви кода.