Шаблоны псевдонимов C++11 в CUDA

Важный вопрос: поддерживаются ли шаблоны псевдонимов компилятором CUDA?

Я использую CUDA 7.5 на Ubuntu с gcc-4.8. Все мои классы шаблонов определены в файлах заголовков и #included в единую единицу перевода во время компиляции.

У меня есть простой класс cuda_array, который обеспечивает тонкую оболочку вокруг класса std::vector. По сути, это очень простая версия thrust::host_vector в сочетании с thrust::device_vector. Его декларация

template <typename T, const size_t N>
class cuda_array {
    std::vector<T> host;
    T *device;
public:
    // lots of type aliases to meet container requirements
    void push() { /* cudaMemcpy(...,H2D); */ }
    void pull() { /* cudaMemcpy(...,D2H); */ }
    // a few others that aren't relevant here
};

Чтобы сделать матрицу, я просто сделал быстрый псевдоним шаблона.

template <typename T, const size_t M, const size_t N>
using cuda_matrix = cuda_array<T, M * N>;

Я хочу сопоставить свое ядро ​​​​CUDA, умножающее матрицу и вектор, на перегруженный operator* для безопасности типов и простоты использования (вызывающий должен убедиться, что push и pull вызываются правильно).

template <typename T, const size_t rows, const size_t cols>
__global__ void matrix_vector_mul(T *A, T *b, T *result) {
     __shared__ T shared_b[cols];
    // rest of it
}

template <typename T, const size_t M, const size_t N>
__host__ cuda_array<T, M> operator*(cuda_matrix<T, M, N> &m, cuda_array<T, N> &v) {
    cuda_array<T, M> result;
    matrix_vector_mul<T, M, N><<<16, 32>>>(m.device_data(), v.device_data(), result.device_data());
    return result;
}

В моем «main.cpp» у меня есть

cuda_matrix<int,16,32> A;
cuda_array<int,32> b;
auto result = A * b;

В последней строке выдает ошибку

error: no operator "*" matches these operands
        operand types are: cuda_matrix<int, 16UL, 32UL> * cuda_array<int, 32UL>

Я выследил всех обычных подозреваемых в ошибках вывода типа шаблона, которые только мог придумать, но ничего не сработало. В отчаянии я преобразовал свой шаблон псевдонима cuda_matrix в класс шаблонов.

template <typename T, const size_t M, const size_t N>
class cuda_matrix : public cuda_array<T, M * N> {};

И ошибка компиляции исчезает! Поэтому кажется, что CUDA еще не поддерживает шаблоны псевдонимов. Или я сделал что-то глупое, что я не могу понять?


c++ templates cuda c++11 template-aliases
person Tim    schedule 08.10.2015    source источник
comment
M N в cuda_array<T, M * N> невыводимы, так как являются частью подвыражения, это вообще не связано с CUDA   -  person Piotr Skotnicki    schedule 08.10.2015

Ответы (1)


arrow_upward
4
arrow_downward

Вы должны помнить, что:

§ 14.5.7 [временный псевдоним]/p2:

Когда идентификатор-шаблона относится к специализации шаблона псевдонима, он эквивалентен связанному типу, полученному путем замены его аргументов-шаблона параметрами-шаблона в идентификаторе-типа< /em> шаблона псевдонима. [ Примечание: имя шаблона псевдонима никогда не выводится. — конец примечания ]

Это означает, что вычет не производится для:

template <typename T, const size_t M, const size_t N>
__host__ cuda_array<T, M> operator*(cuda_matrix<T, M, N> &m, cuda_array<T, N> &v)

но для:

template <typename T, const size_t M, const size_t N>
__host__ cuda_array<T, M> operator*(cuda_array<T, M * N> &m, cuda_array<T, N> &v)
//                                  ~~~~~~~~~~~~~~~~~~~^

Так что:

§ 14.8.2.5 [temp.deduct.type]/p16:

Если в объявлении шаблона функции с параметром шаблона, не являющимся типом, параметр шаблона, не являющийся типом, используется в подвыражении в списке параметров функции, выражение является невыведенным контекстом как указан выше.

M находится в невыводимом контексте, поэтому этот operator* не считается жизнеспособной перегрузкой.

В качестве одного из обходных путей вы можете вместо этого проверить выведенное значение для самого cuda_array:

template <typename T, std::size_t MN, std::size_t N>
auto operator*(const cuda_array<T, MN>& m, const cuda_array<T, N>& v)
    -> typename std::enable_if<(MN/N)*N==MN, cuda_array<T, MN/N>>::type;

или используйте трюк с наследованием, который у вас уже есть; тогда M и N являются отдельными нетиповыми параметрами шаблона cuda_matrix.

person Piotr Skotnicki    schedule 08.10.2015
comment
Я вижу сейчас. Я не знал, что типы шаблонов псевдонимов невыводимы. Удаление частей, специфичных для CUDA, и запуск их через gcc-4.9 дает сообщение об ошибке, которое делает этот момент более ясным. - person Tim; 09.10.2015