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

Question

Существенный вопрос: alias templates поддерживается компилятором CUDA?

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

У меня есть простой класс 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 еще не поддерживает шаблоны псевдонимов. Или я сделал что-то глупое, чего не могу понять?

Answer 1

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

§ 14.5.7 [temp.alias]/p2:

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

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

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.тип]/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.