Объединенная глобальная память записывает с использованием хэша
Мой вопрос касается объединенных глобальных записей в динамически изменяющийся набор элементов массива в CUDA. Рассмотрим следующее ядро:
__global__ void
kernel (int n, int *odata, int *idata, int *hash)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n)
odata[hash[i]] = idata[i];
}
Здесь первый n элементы массива hash содержат индексы odata обновляться с первого раза n элементы idata, Очевидно, это приводит к ужасному, ужасному отсутствию слияния. В моем коде хэш на одном вызове ядра совершенно не связан с хешем на другом (и другие ядра обновляют данные другими способами), поэтому просто переупорядочить данные для оптимизации этого конкретного kenrel не вариант.
Есть ли какая-то особенность в CUDA, которая позволила бы мне улучшить производительность в этой ситуации? Я слышал много разговоров о памяти текстур, но мне не удалось перевести прочитанное в решение этой проблемы.
Решение
Текстурирование является механизмом только для чтения, поэтому оно не может напрямую улучшить производительность отдельных записей в GMEM. Если бы вы «хэшировали» вот так:
odata[i] = idata[hash[i]];
(возможно, ваш алгоритм может быть преобразован?)
Тогда может быть некоторая выгода для рассмотрения Механизм текстуры. (Ваш пример, кажется, 1D в природе).
Вы также можете убедиться, что ваша разделяемая память / разделение L1 оптимизированы для кэширования. Это не сильно поможет с разрозненными записями.
Другие решения
Можете ли вы ограничить диапазон результатов хеширования? Например, вы можете знать, что первые 1К итераций потоков будут иметь доступ в диапазоне от 0 до 8К odata только.
Если это так, вы можете использовать общую память. Вы можете выделить блок совместно используемой памяти и временно выполнить быструю точную запись в разделяемую память. Затем запишите блок разделяемой памяти в глобальную память в объединенных транзакциях.