Я пытаюсь использовать тягу, чтобы уменьшить массив из 1 млн элементов до одного значения. Мой код выглядит следующим образом:
#include<chrono>
#include<iostream>
#include<thrust/host_vector.h>
#include<thrust/device_vector.h>
#include<thrust/reduce.h>
int main()
{
int N,M;
N = 1000;
M = 1000;
thrust::device_vector<float> D(N*M,5.0);
int sum;
auto start = std::chrono::high_resolution_clock::now();
sum = thrust::reduce(D.begin(),D.end(),(float)0,thrust::plus<float>());
auto end = std::chrono::high_resolution_clock::now();
auto duration = std::chrono::duration_cast<std::chrono::microseconds>(end-start);
std::cout<<duration.count()<<" ";
std::cout<<sum;
}
Проблема в том, что только тяга :: уменьшить запускается около 4 мс на графическом процессоре моего ноутбука RTX 3070. Это значительно медленнее, чем код, который я могу написать на основе сокращения № 4 в эта ссылка на CUDA Марка Харриса, что занимает около 150 микросекунд. Я делаю что-то не так здесь?
РЕДАКТИРОВАТЬ 1: Изменено high_resolution_clock на Steady_Clock. Теперь для запуска Thrust::reduce требуется 2 мс. Обновленный код выглядит следующим образом:
#include<chrono>
#include<iostream>
#include<thrust/host_vector.h>
#include<thrust/device_vector.h>
#include<thrust/reduce.h>
int main()
{
int N,M;
N = 1000;
M = 1000;
thrust::device_vector<float> D(N*M,5.0);
int sum;
auto start = std::chrono::steady_clock::now();
sum = thrust::reduce(D.begin(),D.end(),(float)0,thrust::plus<float>());
auto end = std::chrono::steady_clock::now();
auto duration = std::chrono::duration<double,std::ratio<1,1000>>(end-start);
std::cout<<duration.count()<<" ";
std::cout<<sum;
}
Дополнительная информация:
Я запускаю CUDA C++ на Ubuntu в WSL2
Версия CUDA - 11.4
Я использую компилятор nvcc для компиляции:
nvcc -o reduction reduction.cu
Бежать:
./reduction
На самом деле вот что говорит один из людей, наиболее ответственных за создание high_resolution_clock
об использовании.
Хотя тяга проста, она, к сожалению, часто не очень эффективна. Вместо этого попробуйте CUB, который должен насытить (память) GPU. Обратите внимание, что CUB теперь зависит от тяги и является последней версией AFAIK.
Возможно, вам следует также включить код измерения для эталонного решения.
Я получаю ~ 500 us, запуская ваш код. Так что да, с вашим тестовым случаем что-то не так. Возможно, вы создаете отладочный проект для случая тяги. Когда я создаю отладочную версию вашего кода, я получаю отчет о 5 мс.
@paleonix Нет отдельного кода, который я использую для измерения, я просто использовал хроно, чтобы найти и отобразить продолжительность в самом коде.
@user4581301 user4581301 Я удалил high_resolution_clock и заменил его на Steady_Clock. Теперь он работает за 2 мс.
@RobertCrovella Я не понимаю, что ты имеешь в виду. Я просто использую компилятор nvcc для компиляции и запуска программы без включения режима отладки. Не могли бы вы уточнить?
Я имею в виду, что если вы компилируете с переключателем отладки устройства (-G
), это может повлиять на производительность кода CUDA. Для вопросов, связанных с производительностью, я обычно предлагаю, чтобы люди предоставили, в дополнение к тому, что вы предоставили до сих пор, используемую операционную среду (например, Linux или Windows, а также версии CUDA) и точную используемую команду компиляции.
Когда вы редактируете вопрос, чтобы включить дополнительные детали и команду сборки, не стесняйтесь обновлять код до текущей версии.
Am I doing something wrong here?
Я бы не сказал, что вы делаете что-то не так здесь. Однако это может быть вопросом мнения. Немного распакуем его, используя профайлер. Я не использую ту же настройку, что и вы (я использую другой графический процессор — Tesla V100, в Linux, CUDA 11.4). В моем случае измерение, выдаваемое кодом, составляет ~ 0,5 мс, а не 2 мс.
thrust::reduce
выполняется скрытно с помощью вызова cub::DeviceReduceKernel
, за которым следует cub::DeviceReduceSingleTileKernel
. Этот подход с двумя ядрами должен быть вам понятен, если вы изучили материалы по сокращению Марка Харриса. Профилировщик говорит мне, что вместе эти два вызова составляют ~40 мкс из ~500 мкс общего времени. Это время будет наиболее сопоставимо с измерением, которое вы сделали для своей реализации кода сокращения Марка Харриса, предполагая, что вы синхронизируете только ядра. Если мы умножим на 4, чтобы учесть общий коэффициент производительности, он будет довольно близок к вашему измерению в 150 мкс.cudaMalloc
(~200us) и звонок на cudaFree
(~200us). Это неудивительно, потому что если вы изучите cub::DeviceReduce
методология, который, очевидно, используется тягой, он потребует начального вызова для временного распределения. Поскольку тяга обеспечивает автономный вызов thrust::reduce
, он должен выполнить этот вызов, а также выполнить операцию cudaMalloc
и cudaFree
для указанного временного хранилища.Так можно ли что-нибудь сделать?
Разработчики тяги знали об этой ситуации. Чтобы получить (ближе к) сравнение между яблоками и яблоками между простым измерением длительности(ей) ядра реализации CUDA C++ и использованием тяги для выполнения того же самого, вы можете использовать профилировщик для сравнения измерений или же взять под контроль временные ассигнования самостоятельно.
Один из способов сделать это — переключиться с тяги на детеныш.
Самый простой способ сделать это — использовать настраиваемый распределитель.
В методологии могут быть и другие отличия в деталях, которые влияют на ваши измерения. Например, вызов тяги по своей сути копирует результат сокращения обратно в память хоста. Вы можете или не можете рассчитать время для этого шага в другом подходе, который вы не показали. Но, согласно моим измерениям профилировщика, это составляет всего несколько микросекунд.
Примечание: будьте осторожны с
high_resolution_clock
. Если вы не будете осторожны, вы можете найти его поведение неудачным. Например, его можно реализовать как псевдонимsystem_clock
и получить результаты, когда время идет вспять. Хуже того, рабочий стол Windows не предназначен для обеспечения высокой точности синхронизации. Иногда задача может быть завершена за микросекунды, и все равно сообщается о полном тике таймера, близком к 16 мс, потому что именно так часто обновляются часы.