CUDA отключает кэш L1 только для одной переменной
Есть ли способ на устройствах CUDA 2.0 отключить кэш L1 только для одной конкретной переменной? Я знаю, что можно отключить кэш L1 во время компиляции, добавив флаг -Xptxas -dlcm=cg в nvcc для всех операций с памятью. Тем не менее, я хочу отключить кэш только для чтения памяти по определенной глобальной переменной, чтобы вся остальная часть чтения памяти проходила через кэш L1.
Основываясь на поиске, который я сделал в Интернете, возможное решение - через код сборки PTX.
3 ответа
Как уже упоминалось выше, вы можете использовать встроенный 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".
Встроенный PTX может использоваться для загрузки и сохранения переменной. Инструкции ld.cg и st.cg кэшируют данные только в L2. Операторы кэширования описаны в разделе 8.7.8.1 Операторы кэширования документа PTX ISA 2.3. Инструкции или проценты лд и ст. Inline PTX описан в разделе Использование встроенной сборки PTX в CUDA.
Если вы объявите переменную volatileтогда он будет кэшироваться только в кеше L2 на графических процессорах Fermi. Обратите внимание, что некоторые оптимизации компилятора, такие как удаление повторяющихся загрузок, не выполняются с переменными переменными, потому что компилятор предполагает, что они могут быть записаны другим потоком.