Или как я научился не беспокоиться из-за cudaGetLastError() и просто признал её нежелательной
Любому, кто сталкивался с проектированием API, приходилось мучиться над вопросом, как лучше всего распространять ошибки к вызывателям функций. Эта проблема появилась ещё раньше, чем термин «API». Даже несколько десятилетий назад (скажем, 30+) проектировщики интерфейсов умели отделять возвращаемые ошибки от прочей полезной нагрузки в рамках тех функций, которые доставляют эту информацию вызвавшему их узлу.
Иногда полезно знать, как лучше не делать. Мой любимый пример антипаттерна в рассматриваемой области — прославленная функция atoi() из среды выполнения C, которая преобразует строку в целое число:
int atoi (const char * str);Поскольку не существует такого гарантированно недействительного значения int, которое показало бы вызывающей стороне, что функции была передана недопустимая строка (такая, которая не может быть интерпретирована как целое число), этот интерфейс принципиально плох. В некоторых реализациях даже подразумевается, что в случае ошибки возвращается 0, поэтому возникает вопрос: а допустима ли строка «0» в качестве ввода?
Среда выполнения C кишит подобными неаккуратными интерфейсами и, пожалуй, лучше всего доказывает, что умение проектировать языки и умение проектировать интерфейсы вместе практически не встречаются.
Хорошо, мы разобрали, как лучше не делать. Давайте рассмотрим, какими ещё способами можно обрабатывать ошибки. Они подразделяются на три обширные категории:
- Исключения: функции в составе API могут сигнализировать об ошибках в целевом языке. Любой элемент стека вызовов может выделить исключение, заключив вызов(ы) API в условие try.
- Немедленный возврат ошибки: определяется унифицированный тип ошибок, специфичный для данного API, а конкретные ошибки различаются путём указания кода.
- Получение последней ошибки: в случае возникновения ошибки устанавливается глобальное состояние ошибки, которое впоследствии может быть запрошено другой функцией, относящейся к данному API.
Для начала давайте отсеем исключения. Даже если целевой язык известен, и мы знаем, что в этом языке поддерживаются исключения (в этой категории сразу вспоминаются C++ и Python), порой выдача исключения — не лучший вариант. Исключения присутствуют в языке C++ вот уже 30 лет по-прежнему воспринимаются неоднозначно, поскольку сложно проверить код и определить, где и почему могут сигнализироваться какие исключения. Раймонд Чен, легендарный разработчик из Microsoft, занимавшийся Windows более 30 лет, несколько раз высказывался на эту тему в своём блоге уже почти 20 лет назад. Ему потребовалось написать продолжение к “Cleaner, More Elegant, and Wrong”, статью “Cleaner, More Elegant, and Harder To Recognize”, так как сложно согласиться с неизбирательным порицанием исключений. Но при написании библиотек, например, среды выполнения для CUDA или предметно-ориентированных библиотек CUDA, например, cuBLAS, cuSOLVER и им подобных, легко обосновать обратное делегирование исключений, так, чтобы они обрабатывались на вызывающей стороне. Действительно, клиент всегда может по своему усмотрению преобразовывать возвращённые ошибки в исключения. С другой стороны, чтобы инкорпорировать обработку ошибок в такую библиотеку (например, определить тип исключения и документировать, какие именно функции будут сигнализировать о каких исключениях, и в каких именно обстоятельствах), нужно навязать клиенту соответствующую политику. Если при этом вдобавок учитывать, что специфика обработки исключений зависит от языка, то оказывается, что проще вообще воздержаться от обработки исключений в библиотеках, которые проектируются как не зависящие от языка.
Об исключениях поговорили. Следующий вариант распространения ошибки — это её немедленный возврат. Именно такая стратегия была выбрана в CUDA с самого начала, но с одной оговоркой, на которой мы остановимся в следующем разделе. Изначально в CUDA немедленный возврат ошибок был оформлен как тип CUresult для драйверного API и как cudaError_t для среды выполнения CUDA. Практически все функции API возвращают какой-то код состояния, означающий ошибку. Если они выделяют ресурс по поручению вызывающей стороны, то принимают параметр, через который функция может передать обратно дескриптор данного ресурса. Вот почему cudaMalloc() возвращает cudaError_t, а не void *, как функция malloc() из среды выполнения C. Пусть даже cudaMalloc() могла бы эмулировать malloc() среды выполнения C — для этого в случае отказа следовало бы возвращать гарантированно недопустимое значение. Но разработчики сочли, что важнее сделать среду выполнения CUDA самодостаточной, чем оставить в ней древний отголосок API, существовавшего в среде исполнения для C. (Если вы уже встрепенулись, готовясь защищать среду выполнения C, предлагаю вам пересмотреть семантику realloc(). Я могу этим заниматься целый день!)
Главная проблема, возникающая при немедленном возврате ошибки – в том, что ошибка может быть проверена только там, где её вызывали. Если клиент вызывает функцию, та вызывает ошибку, а вызывающая сторона решает её проигнорировать, то подробности о возвращённой ошибке просто теряются. Но то, что на первый взгляд кажется багом, на самом деле может быть фичей – иными словами, часто код пишут именно так, что он явно игнорирует возможные возвраты ошибок. Рассмотрим функцию, которая по поручению вызывающей стороны выделяет два буфера:
cudaError_t allocateTwoBuffers( void **bufferA, size_t nA, void **buffer,
size_t nB )
{
cudaError_t err = cudaMalloc( bufferA, nA );
if ( cudaSuccess == err ) {
err = cudaMalloc( bufferB, nB );
if ( cudaSuccess != err ) {
cudaFree( bufferA );
}
}
return err;
}
Здесь возвращаемое значение cudaFree() игнорируется, так как мы хотим распространить информацию о выделении ресурса, а не те данные, которые возвращает cudaFree().
А теперь – хорошая возможность отдать должное среде выполнения C, хотя бы после того, какую большую работу в ней проделали различные комитеты по стандартизации. Дело в том, что free() возвращает void; NULL считается допустимым вводом для free(); а free(NULL) определяется как фиктивная операция (no-op). В ретроспективе можно сказать, что благоразумно было бы заранее определить подобную семантику для семейства API CUDA, которые высвобождают ресурсы — не только для cudaFree(), но и для cudaStreamDestroy(), cudaEventDestroy() и т.д. Но того же эффекта в приложениях можно добиться, просто игнорируя поступающие от этих функций возвращаемые значения — как было показано выше в примере с allocateTwoBuffers()
Функция allocateTwoBuffers() структурирована не так, как следует показывать в учебниках, поскольку идиома для обработки ошибок, построенная на основе goto, будет масштабироваться лучше при условии, что функция немного модифицирована — например, по поручению вызывающей стороны может выделить больше ресурсов.
Как получать последнюю ошибку
Итак, разобравшись с обработкой исключений и немедленным возвратом ошибок, переходим к идиомам с получением последней ошибки. Они встречаются в API различных платформ: у OpenGL это glGetError(), у Win32 — GetLastError(), а у CUDA — cudaGetLastError() и cudaPeekLastError(). Их можно встретить даже в среде исполнения C, в форме errno.
Точная семантика паттернов возвращения последней ошибки может варьироваться от API к API. Определённо, не каждая функция Win32 обновит ошибку, которая будет возвращена от GetLastError(), и не каждая функция среды выполнения C установит errno. Но в целом идея такова: если произойдёт ошибка, то устанавливается глобальная (а в многопоточных системах — локальная в пределах потока) переменная, которую впоследствии можно проверить и узнать, не произошла ли ошибка. Плюс такого подхода в том, что ошибки «прилипают» к функциям, то есть, ошибка может быть проверена не только на вызывающей стороне и не только там, где функция API вернула ошибку.
Я предпочитаю не пользоваться в моих API идиомами возврата последней ошибки, и у меня на это несколько причин:
- Эти идиомы не потокобезопасны – «последняя ошибка» одного потока ЦП определённо отличается от последней ошибки, возвращённой другим потоком, так что «последнюю ошибку» необходимо сохранять в ячейке TLS.
- Этим идиомам сложно обучать: как часто следует проверять последнюю ошибку? После каждого вызова? Да, мне известно одно знаменитое приложение для OpenGL, в котором glGetLastError() срабатывает после каждого вызова API. Но ни один разработчик API, заслуживающий своей зарплаты, не будет выступать за то, чтобы удвоить количество вызовов API для своих клиентов. Итак, как часто?
- Семантика, связанная с установкой, проверкой и очисткой последних ошибок нуждается в разъяснении. Подтверждение тому —CUDA, где сосуществуют отдельные функции cudaGetLastError() и cudaPeekLastError(). Функция cudaGetLastError() очищает код ошибки, поэтому её можно лишь один раз вызвать в программе ниже той точки, в которой ошибка произошла. В результате возникают примерно такие же проблемы, как и с кодами при немедленном возврате ошибки — в частности, проверить ошибку можно только на вызывающей стороне. Функция cudaPeekLastError() обеспечивает больше контроля и пространства для творчества, но при работе с большими объёмами кода следует тщательно определять собственные протоколы, в которых, например, описано, как именно должны вызываться различные функции для обработки ошибок.
- Наиболее важно, в чём можно убедиться на примере приведённого выше фрагмента allocateTwoBuffers() — что иногда в приложениях бывает целесообразно явно игнорировать потенциальные коды ошибок, но в то же время передавать коды ошибок, возвращённые ранее.
Итак, в CUDA применяется немедленный возврат ошибок. Все функции CUDA возвращают коды ошибок, а в хорошо написанных приложениях для CUDA проверяются все до единого возвращаемые значения. Редко случается так, что функция CUDA не может оказать. Я как проектировщик API активно выступаю за разработку именно таких функций, которые не могут отказать, но они сложны с инженерной точки зрения. Обычно такая функция подразумевает, что ресурсы будут выделяться заранее, так, что можно не сомневаться, что ресурс действительно существует уже потому, что существует код для его обработки. (Вот почему в семействе функций ядра для синхронизации и выделения потоков от Windows NT такие функции как KeInitializeMutex() и KeInitializeSemaphore() возвращают void. Они часто вызываются для работы со статически выделенной памятью, поэтому, если код в принципе работает, мы знаем, что загрузчику удалось выделить память там, где располагается мьютекс или семафор). Также в CUDA заранее выделяются ресурсы, чтобы снизить вероятность отказа во время выполнения. Когда CUDA был совсем новым языком, требовалось удостовериться, что все контексты CUDA заблаговременно выделили промежуточные буферы, необходимые для копирования памяти. Если выделение этих промежуточных буферов (или любых других бесчисленных ресурсов, необходимых контексту CUDA) не прошло, то откажет и сам этот контекст CUDA. Таким образом, если у вас в распоряжении есть действующий дескриптор контекста CUDA, то вы можете безотказно копировать память, даже если по каким-то причинам не удалось выделить промежуточный буфер. Но функции копирования памяти в CUDA могут отказывать и по другим причинам, например, если приложение передаёт на асинхронное копирование ту часть памяти хоста, которая зависит от подкачки. Среди других функций, которые обычно не отказывают — cuEventQuery() и cuStreamSynchronize(); но эти функции могут провоцировать передачу подвисших задач для обработки на аппаратный уровень. Это также может приводить к отказам, которых мы коснёмся в следующем разделе. Таким образом, вызыватели функций должны проверять их возвращаемые значения.
Учитывая, что CUDA уже старше 15 лет, просто удивительно, насколько разнообразные идиомы обработки ошибок там применяются. Даже не вдаваясь в подробности о том, что именно вы хотите делать при возникновении ошибки — например, вывести на экран информацию и выйти, продвинуть ошибку к вызывающей стороне, просигнализировать об исключении и т.д. — нет общего мнения о том, воспользоваться ли шаблонизированной служебной функцией или просто выполнить макрос препроцессора. Большинство разработчиков предпочитают определять именованные макросы, например, SAFE_CUDA или CUDART_CHECK.
Даже в примерах из CUDA SDK не приведён стандартизированный метод обработки ошибок.
Неявное получение последней ошибки
Как следует изучив API CUDA для немедленного возврата ошибок, убеждаешься, что в этой операции скрывается семантика получения последней ошибки. В cudaDeviceSynchronize() и многих других API CUDA содержатся аннотации примерно в таком духе: «Обратите внимание: эта функция также может возвращать коды ошибок от более ранних асинхронных запусков». Дело в том, что при каждом вызове ядра могут возникать ошибки времени выполнения — например, разыменование недействительного указателя — восстановиться после которых практически невозможно. Если не обеспечивается более детализированная проверка ошибок, то асинхронные API (например, запуск ядра, который всегда был асинхронным, даже в CUDA 1.0) просто требуют действовать в парадигме получения последней ошибки. Для проверки ошибок требуется синхронизировать CPU и GPU, что плохо влияет на производительность. Сразу просится пример о том, как именно распространяются значения NaN и INF при вычислениях с плавающей точкой. Проверять каждую арифметическую операцию слишком затратно, поэтому удобно наладить механизм, благодаря которому ошибки прилипают к результатам и могут быть проверены позднее, с более грубой детализацией. Комитет IEEE по стандартизации признал такой компромисс приемлемым.
Насколько мне известно, вполне возможно писать совершенно корректные программы на CUDA, вообще не вызывая cudaGetLastError(). Если у вас будут какие-то функции CUDA, которые устанавливают последнюю ошибку, но не возвращают код ошибки немедленно, то можно заключить, что NVIDIA просто усложнила работу всем программистам без какой-то явной причины, однако ничто не доказывает, что NVIDIA действительно так сделала.
Создаётся впечатление, как будто эта функция была добавлена просто для более явной демонстрации семантики последней ошибки, но эта семантика уже и так включена в среду выполнения CUDA.
Действенное средство
Будет не так сложно добавить более детализированную обработку ошибок, которая почти не скажется на производительности: достаточно будет просто немного перегрузить события CUDA. В CUDA уже есть события, рассчитанные на блокирующее или неблокирующее ожидание; почему бы не добавить механизм для сообщения об ошибках? С его помощью удалось бы хотя бы вычленить, какое именно ядро привело к ошибке. Новая функция могла бы запрашивать статус ошибки для данного сообщения и возвращать код состояния для последней произошедшей ошибки на тот момент, когда поступил сигнал о событии:
CUresult cuEventCheckError( CUresult *p, CUevent hEvent, uint32_t Flags );Такой гипотетический API мог бы обязывать, что поступивший сигнал о событии характеризует это событие как допустимый ввод. Также можно было бы ставить флаг, означающий, что функция обязательно должна работать синхронно (например, мы бы дожидались выполнения условий события, и только затем возвращали бы статус последней ошибки). Детали здесь не важны. Этот API всё равно мог бы напоминать функцию get-last-error, поскольку, если до события происходили бы множественные пуски ядра или вызовы к API CUDA, то было бы точно так же непонятно, какая именно операция спровоцировала ту ошибку, о которой необходимо просигнализировать. Но степень детализации контролировал бы сам разработчик приложения, и программа была бы ограничена по скорости лишь аппаратными возможностями CUDA.
