Использовать общую память для соседних элементов массива?
Я хотел бы обработать изображение с CUDA. Новое значение каждого пикселя рассчитывается на основе двух соседних пикселей в одной строке. Будет ли смысл использовать __shared__ память для значений пикселей, так как каждое значение будет использоваться только дважды? Разве плитки не являются неправильным способом сделать это, так как это не соответствует структуре проблемы? Мой подход состоял бы в том, чтобы запустить поток для каждого пикселя и загружать значения соседних пикселей каждый раз для каждого потока.
Решение
Все поддерживаемые в настоящее время архитектуры CUDA имеют кэши.
Начиная с вычислительной возможности 3.5 и выше они особенно эффективны для данных только для чтения (поскольку данные для чтения и записи кэшируются только в L2, кэш L1 ограничен данными только для чтения). Если вы пометите указатель на входные данные как const __restrict__скорее всего компилятор его загрузит через кеш текстуры L1. Вы также можете форсировать это, явно используя __ldg() встроенная.
Хотя можно явно управлять повторным использованием данных из соседних пикселей через разделяемую память, вы, вероятно, обнаружите, что это не дает никаких преимуществ по сравнению с использованием только кэша.
Конечно, независимо от того, используете ли вы разделяемую память или нет, вы хотите максимизировать размер блоков в x-направлении и использовать blockSize.y, равный 1, для оптимальной локализации доступа.
Другие решения
Объедините использование общей памяти с преимуществами объединенного доступа к памяти. Все, что вам нужно сделать, это убедиться, что изображение хранится построчно. Каждый блок будет обрабатывать кусок линейного массива. Из-за повторного использования данных (каждый пиксель, кроме первого и последнего будет участвовать в обработке три раза), было бы полезно, если бы в начале вашего ядра вы копировали значения всех пикселей, которые будут обрабатываться, в общую память.