CUDA: время доступа к памяти текстур похоже на объединенную глобальную память? - PullRequest
Новая
       28

CUDA: время доступа к памяти текстур похоже на объединенную глобальную память?

3 голосов
/ 24 февраля 2012

Мои потоки ядра объединяются в линейный массив символов.Если я сопоставлю массив с текстурой, я не вижу никакого ускорения.Время работы практически одинаково.Я работаю над Tesla C2050 с вычислительными возможностями 2.0 и где-то читал, что глобальный доступ кешируется.Это правда?Возможно, именно поэтому я не вижу разницы во времени выполнения.

Массив в основной программе: 

char *dev_database = NULL;
cudaMalloc( (void**) &dev_database, JOBS * FRAGMENTSIZE * sizeof(char) );

, и я связываю его с текстурой texture<char> texdatabase с помощью 

cudaBindTexture(NULL, texdatabase, dev_database, JOBS * FRAGMENTSIZE * sizeof(char) );

Каждый поток затем читает символ ch = tex1Dfetch(texdatabase, p + id)где id равен threadIdx.x + blockIdx.x * blockDim.x, а p - это смещение.

Я связываюсь только один раз, а dev_database - большой массив.На самом деле я обнаружил, что если размер слишком велик, связывание не выполняется.Есть ли ограничение на размер связываемого массива?Большое спасибо.

1 Ответ

3 голосов
/ 24 февраля 2012

Существует несколько причин, почему вы не видите никакой разницы в производительности, но наиболее вероятно, что этот доступ к памяти не является вашим узким местом .Если это не ваше узкое место, то его ускорение не повлияет на производительность.

Что касается кэширования: в этом случае, поскольку вы читаете только байты, каждая деформация будет считывать 32 байта, что означает каждую группу из 4Деформации будут отображаться в каждой строке кэша.Таким образом, при условии небольшого количества конфликтов в кэше, вы получите до 4x повторного использования из кэша.Таким образом, если этот доступ к памяти является узким местом, вполне возможно, что кеш текстуры не принесет вам больше пользы, чем кеш общего назначения.

Сначала вы должны определить, ограничены ли вы пропускной способностью и является ли этот доступ к даннымвиновник.Как только вы это сделали, оптимизируйте доступ к памяти.Другая тактика, которую следует рассмотреть, - это доступ к 4-16 символам на поток на загрузку (используя структуру char4 или int4 с байтовой упаковкой / распаковкой), а не одну на поток для увеличения количества транзакций памяти в полете за раз - это может помочьнасыщать шину глобальной памяти.

Есть хорошая презентация Паулюса Микикявичюса из GTC 2010 , которую вы, возможно, захотите посмотреть.Он охватывает как оптимизацию на основе анализа, так и конкретную концепцию транзакций памяти в полете.

...