Динамическая общая память
Задавайте размер общей памяти во время запуска.
«Динамическая общая память» — бесплатный урок CUDA Academy на CoddyKit. Это урок 4 из 4. Ты можешь прочитать весь урок бесплатно ниже — а потом практиковать его прямо в браузере с встроенным редактором кода и ИИ-репетитором 24/7. Это часть пути обучения CUDA Academy, и твой прогресс синхронизируется между веб-версией и приложением CoddyKit. Курс CUDA Academy содержит 4 уроков всего.
Когда размер ещё неизвестен
Иногда размер массива в разделяемой памяти зависит от значения, известного только во время выполнения, поэтому задать его жёстко нельзя. CUDA решает эту задачу с помощью динамической разделяемой памяти.
Объявление extern
Объявите её с помощью extern __shared__ и пустых квадратных скобок. Размер не указывается в коде, поэтому компилятор оставляет его неопределённым.
extern __shared__ float sdata[];Размер задаётся при запуске
Передайте число байтов как третий параметр запуска — между числом потоков и закрывающими скобками. Так размер задаётся во время выполнения.
kern<<<blocks, threads, n * sizeof(float)>>>();Размер всегда указывается в байтах
Распространённая ошибка: третий аргумент задаёт размер в байтах, а не число элементов. Всегда умножайте количество элементов на sizeof, чтобы получить правильный размер.
По умолчанию — ноль байтов
Если пропустить третий аргумент, CUDA зарезервирует ноль байтов динамической разделяемой памяти. Чтение sdata в таком случае приводит к неопределённому поведению, поэтому не забывайте этот аргумент.
Статическая и динамическая память вместе
В одном ядре можно использовать оба вида. Статические массивы сохраняют фиксированные размеры, а динамическая часть меняется в соответствии с аргументом запуска.
Один буфер, много массивов
Динамическая разделяемая память предоставляет один линейный буфер. Чтобы разместить несколько массивов, самостоятельно разделите его с помощью смещений указателей. 🔪
Разделение с помощью смещений
Укажите начало второго массива в некоторой позиции внутри буфера. Только следите за тем, чтобы выровнять каждый участок и чтобы значение с плавающей точкой никогда не начиналось посреди слова.
extern __shared__ float buf[];
float* a = buf;
float* b = &buf[blockDim.x];Зачем нужна динамическая память
Она позволяет одному скомпилированному ядру работать со множеством размеров фрагментов. Вы настраиваете размер разделяемой памяти при каждом запуске, не перекомпилируя ядро для каждого случая.
Следите за аппаратным ограничением
Тем не менее превышать максимальный размер для блока нельзя. Если при запуске запросить слишком много байтов, ядро просто не запустится.
Разрешение больших выделений
Чтобы превысить стандартный предел на новых GPU, перед запуском необходимо opt in с помощью cudaFuncSetAttribute. В противном случае большие запросы отклоняются.
Быстрая проверка
Проверим, как задаётся размер динамической разделяемой памяти.
Итоги
Вы узнали, что массивы extern __shared__ получают размер в байтах при запуске, могут быть разделены на несколько массивов и ограничены лимитом на блок. Курс завершён! 🎉
Изучай C++ с ИИ-репетитором — бесплатно
Пиши и запускай код прямо в браузере, получай мгновенную помощь от ИИ-репетитора 24/7 и продолжи учиться на сайте или в приложении.
- Курсы
- 30
- Уроки
- 120
Часто задаваемые вопросы
Урок «Динамическая общая память» бесплатный?
Да — полный текст урока «Динамическая общая память» бесплатно доступен здесь в веб-версии. Чтобы практиковать его интерактивно (встроенный редактор кода и ИИ-репетитор 24/7) и разблокировать остальной курс CUDA Academy, подпишись на CoddyKit PRO. Курс CUDA Academy содержит 4 уроков всего.
Чему я научусь в уроке «Динамическая общая память»?
Задавайте размер общей памяти во время запуска. Ты практикуешь CUDA Academy с помощью реального кода, который запускаешь прямо в браузере, и ИИ-репетитор 24/7 отвечает на твои вопросы во время урока.
Нужен ли мне опыт, чтобы начать CUDA Academy?
Предыдущий опыт не требуется. CUDA Academy на CoddyKit структурирован для всех уровней — от новичков до продвинутых, поэтому ты можешь начать отсюда или с самого начала и учиться в своем темпе. Это урок 4 из 4.
Сколько времени занимает урок «Динамическая общая память»?
Большинство уроков CoddyKit занимают около 5–10 минут. Каждый из них компактный и интерактивный, поэтому ты постоянно делаешь прогресс и продолжаешь с того же места в веб-версии и приложении.
Можно ли писать и запускать код в этом уроке CUDA Academy?
Да. Каждый урок CUDA Academy включает встроенный редактор кода, поэтому ты пишешь и запускаешь реальный код прямо в браузере и получаешь моментальную обратную связь от AI — локальная установка не требуется.
Все уроки этого курса
- Объявление массивов __shared__
- Синхронизация с помощью __syncthreads
- Как избежать конфликтов банков памяти
- Динамическая общая память