这里有一个比较好的例子，在cuda 的复数运算中，复数的加法 在 float2 类型的复数，与double2 类型的复数，函数名不同，但是已经实现过了，可以直接拿过来使用。例如 加法：cuCaddf(cuFloatComplex a, cuFloatComplex b)
cuCadd(cuDoubleComplex a, cuDoubleComplex b)
多了个 f，这是因为其为 c语言风格的函数，不能重名。
在cpp 风格的模版开发中，希望不同的函数名/cuda kernel 名 能够覆盖相同功能的运算，但这是后的。
这里的示例给出了通过 特化，实现封装的过程
hello.cu
#include "cuComplex.h" #include <stdio.h> template <typename T1, typename T2> __global__ void add(T1 *a, T1 *b, T1 *c, T2 alpha, int N) { } template<> __global__ void add<cuFloatComplex, float>(cuFloatComplex* A, cuFloatComplex* B, cuFloatComplex* C, float alpha, int N) { int i = blockIdx.x*blockDim.x + threadIdx.x; if(i<N) { C[i] = cuCaddf(C[i], cuCaddf(A[i], B[i])); C[i].x *= alpha; } } template<> __global__ void add<cuDoubleComplex, double>(cuDoubleComplex* A, cuDoubleComplex* B, cuDoubleComplex* C, double alpha, int N) { int i = blockIdx.x * blockDim.x + threadIdx.x; if(i<N) { C[i] = cuCadd(C[i], cuCadd(A[i], B[i])); C[i].x *= alpha; } } void init(cuFloatComplex* A, int N, float seed) { for(int i=0; i<N; i++) A[i] = make_cuFloatComplex(seed, seed); } void print_vector(cuFloatComplex* A, int N) { for(int i=0; i<N; i++) printf("(%3.2f,%3.2f)", A[i].x, A[i].y); printf("\n"); } int main() { cuFloatComplex* Ah = nullptr; cuFloatComplex* Bh = nullptr; cuFloatComplex* Ch = nullptr; cuFloatComplex* Ad = nullptr; cuFloatComplex* Bd = nullptr; cuFloatComplex* Cd = nullptr; int N =64; Ah = (cuFloatComplex*)malloc(N * sizeof(cuFloatComplex)); Bh = (cuFloatComplex*)malloc(N * sizeof(cuFloatComplex)); Ch = (cuFloatComplex*)malloc(N * sizeof(cuFloatComplex)); init(Ah, N, 1.0f); init(Bh, N, 2.0f); init(Ch, N, 7.0f); cudaMalloc((void**)&Ad, N*sizeof(cuFloatComplex)); cudaMalloc((void**)&Bd, N*sizeof(cuFloatComplex)); cudaMalloc((void**)&Cd, N*sizeof(cuFloatComplex)); cudaMemcpy(Ad, Ah, N*sizeof(cuFloatComplex), cudaMemcpyHostToDevice); cudaMemcpy(Bd, Bh, N*sizeof(cuFloatComplex), cudaMemcpyHostToDevice); cudaMemcpy(Cd, Ch, N*sizeof(cuFloatComplex), cudaMemcpyHostToDevice); add<cuFloatComplex, float><<<1, 64>>>(Ad, Bd, Cd, .1, N); cudaDeviceSynchronize(); cudaMemcpy(Ch, Cd, N*sizeof(cuFloatComplex), cudaMemcpyDeviceToHost); print_vector(Ch, N); return 0; }
