Есть ли способ на устройствах CUDA 2.0 отключить кеш L1 только для одной конкретной переменной?
Я знаю, что можно отключить кеш L1 во время компиляции, добавив флаг -Xptxas -dlcm=cg в nvcc для всех операций с памятью.
Тем не менее, я хочу отключить кеш только для чтения в памяти по определенной глобальной переменной, чтобы вся остальная часть памяти читала, чтобы пройти через кеш L1.
На основе поиска, который я сделал в Интернете, возможно решение с помощью кода сборки PTX.
Как уже упоминалось выше, вы можете использовать встроенный PTX, вот пример:
__device__ __inline__ double ld_gbl_cg(const double *addr) {
double return_value;
asm("ld.global.cg.f64 %0, [%1];" : "=d"(return_value) : "l"(addr));
return return_value;
}
Вы можете легко изменить это, заменив .f64 на .f32 (float) или .s32 (int) и т.д., ограничение return_value "= d" для "= f" (float) или "= r" (int ) и т.д.. Обратите внимание, что последнее ограничение перед (addr) - "l" означает 64-разрядную адресацию, если вы используете 32-разрядную адресацию, это должно быть "r".
Inline PTX может использоваться для загрузки и сохранения переменной. В инструкциях ld.cg и st.cg хранятся только данные кэша в L2. Операторы кэша описаны в разделе 8.7.8.1 Операторы кэширования документа PTX ISA 2.3. Инструкции или интерес: ld и st. Встроенный PTX описан в Использование встроенной сборки PTX в CUDA.
Если объявить переменную как volatile, тогда она будет кэшироваться только в кеше L2 на графических процессорах Fermi. Обратите внимание, что некоторые оптимизаторы компилятора, такие как удаление повторяющихся нагрузок, не выполняются для переменных volatile, потому что компилятор предполагает, что они могут быть записаны другим потоком.