CUDA 中的 C++11 别名模板
C++11 alias templates in CUDA
基本问题是 alias templates 是否受 CUDA 编译器支持?
我在 Ubuntu 上使用 CUDA 7.5 和 gcc-4.8。我的所有模板 classes 都在头文件中定义,并在编译期间 #include
d 到单个翻译单元中。
我有一个简单的 cuda_array
class,它提供了一个围绕 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
别名模板转换为模板 class.
template <typename T, const size_t M, const size_t N>
class cuda_matrix : public cuda_array<T, M * N> {};
并且编译错误消失了!因此,CUDA 似乎还不支持别名模板。还是我做了一些我想不通的蠢事?
你必须记住:
§ 14.5.7 [temp.alias]/p2:
When a template-id refers to the specialization of an alias template, it is equivalent to the associated type
obtained by substitution of its template-arguments for the template-parameters in the type-id of the alias
template. [ Note: An alias template name is never deduced. — end note ]
这意味着不执行扣除:
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:
If, in the declaration of a function template with a non-type template parameter, the non-type template
parameter is used in a subexpression in the function parameter list, the expression is a non-deduced context
as specified above.
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
.
的独立非类型模板参数
基本问题是 alias templates 是否受 CUDA 编译器支持?
我在 Ubuntu 上使用 CUDA 7.5 和 gcc-4.8。我的所有模板 classes 都在头文件中定义,并在编译期间 #include
d 到单个翻译单元中。
我有一个简单的 cuda_array
class,它提供了一个围绕 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
别名模板转换为模板 class.
template <typename T, const size_t M, const size_t N>
class cuda_matrix : public cuda_array<T, M * N> {};
并且编译错误消失了!因此,CUDA 似乎还不支持别名模板。还是我做了一些我想不通的蠢事?
你必须记住:
§ 14.5.7 [temp.alias]/p2:
When a template-id refers to the specialization of an alias template, it is equivalent to the associated type obtained by substitution of its template-arguments for the template-parameters in the type-id of the alias template. [ Note: An alias template name is never deduced. — end note ]
这意味着不执行扣除:
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:
If, in the declaration of a function template with a non-type template parameter, the non-type template parameter is used in a subexpression in the function parameter list, the expression is a non-deduced context as specified above.
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
.