Отсутствует встроенная буква ограничения CUDA PTX для 8-битных переменных, чтобы отключить кэш L1 для 8-битной переменной (bool)

Question

ВСТУПЛЕНИЕ

В этом вопросе мы можем узнать, как отключить кэш L1 для одной переменной. Вот принятый ответ:

Как уже упоминалось выше, вы можете использовать встроенный 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" для "=d" (float) или "=r" (int) и т. Д. Обратите внимание, что последнее ограничение перед (addr) - "l" - обозначает 64-битную адресацию, если вы используете 32-битную адресацию, это должно быть "r".

Теперь я хочу загрузить логическое значение (1 байт), а не число с плавающей запятой. Итак, я подумал, что мог бы сделать что-то вроде этого (для архитектуры>=sm_20):

__device__ inline bool ld_gbl_cg(const bool* addr){
  bool return_value;
  asm("ld.global.cg.u8 %0, [%1];" : "=???"(return_value) : "l"(addr));
  return return_value;
}

, где "???" должна быть соответствующей буквой ограничения для логического значения, соответственно для 8-битного незапятнанного целого числа (из этого вопроса я вывел это, поскольку отмечается, что для>=sm_20 для логического значения используется "u8"). Однако я не могу найти подходящую букву ограничения в документе nvidias " Использование встроенной сборки PTX в CUDA" (на странице 6 перечислены некоторые буквы ограничения). Итак, мой вопрос:

ВОПРОС

1. Существует ли какое-либо письмо с ограничением для встроенного PTX CUDA для любого из типов:

   • логический
   • 8-разрядное целое число без знака
   • или 8-битная переменная evtl

2. Если нет, что я могу сделать в моем случае (объяснено во введении)? - Могут ли помочь параметры "b0", "b1" и т. Д., Которые вкратце обсуждаются здесь?

Заранее большое спасибо за любую помощь или комментарии!

ОБНОВИТЬ

Мне также нужно чтение функции хранилища из кэша L2 вместо глобальной памяти - то есть функция хранилища, которая дополняет вышеприведенную функцию ld_gbl_cg (только когда у меня есть эта функция, я могу полностью проверить, что ответ njuffa работает). Моя лучшая догадка, основанная на ответе njuffa ниже:

__device__ __forceinline__ void st_gbl_cg (const bool *addr, bool t)
{
#if defined(__LP64__) || defined(_WIN64)
    asm ("st.global.cg.u8 [%0], %1;" : "=l"(addr) : "h"((short)t));
#else
    asm ("st.global.cg.u8 [%0], %1;" : "=r"(addr) : "h"((short)t));
#endif
}

Однако компилятор выдает предупреждение "параметр" addr "был установлен, но никогда не использовался", и во время выполнения программы происходит сбой с "неопределенной ошибкой запуска". Я также попробовал с.u16 вместо.u8, так как я не знаю, к чему именно это относится. И все же результат тот же.

(Дополнительная информация) Следующий пункт в документации по PTX 3.1, кажется, важен для этого вопроса:

5.2.2 Ограниченное использование размеров подслов. Типы команд.u8, .s8 и.b8 ограничены инструкциями ld, st и cvt. Тип с плавающей точкой.f16 допускается только при преобразованиях в типы.f32 и.f64 и обратно. Все инструкции с плавающей точкой работают только для типов.f32 и.f64. Для удобства команды ld, st и cvt позволяют операндам источника и назначения быть шире, чем размер типа инструкции, так что узкие значения могут быть загружены, сохранены и преобразованы с использованием регистров обычной ширины. Например, 8-битные или 16-битные значения могут храниться непосредственно в 32-битных или 64-битных регистрах при загрузке, сохранении или преобразовании в другие типы и размеры.

Answer 1

Согласно документу "Использование встроенного PTX в CUDA", нет никаких ограничений для операндов размером в байт. Насколько я могу судить, наиболее близким к желаемой функциональности является перемещение данных через промежуточный "короткий". Это приводит к одной дополнительной инструкции SASS для преобразования из "short" в "bool".

__device__ __forceinline__ bool ld_gbl_cg (const bool *addr)
{
    short t;
#if defined(__LP64__) || defined(_WIN64)
    asm ("ld.global.cg.u8 %0, [%1];" : "=h"(t) : "l"(addr));
#else
    asm ("ld.global.cg.u8 %0, [%1];" : "=h"(t) : "r"(addr));
#endif
    return (bool)t;
}