c++ - 在cuda中用作模板参数

Question

我正在尝试在 cuda 中编写一个归约函数(这是一个练习，我知道我正在做其他人做得更好的事情)，它采用二元关联运算符和一个数组并归约数组使用运算符。

我在传递函数时遇到困难。我已经将 hostOp() 编写为基于主机的示例，它运行良好。

deviceOp() 适用于显式调用 fminf() 的第一条语句，但是当我调用函数参数时，出现非法内存访问错误。

#include <iostream>
#include <cstdio>
#include <cmath>
using namespace std; //for brevity

__device__  float g_d_a = 9, g_d_b = 5;
float g_h_a = 9, g_h_b = 5;

template<typename argT, typename funcT>
__global__
void deviceOp(funcT op){    
    argT result = fminf(g_d_a, g_d_b);                  //works fine
    printf("static function result: %f\n", result);
    result = op(g_d_a,g_d_b);                           //illegal memory access
    printf("template function result: %f\n", result);
}

template<typename argT, typename funcT>                 
void hostOp(funcT op){
    argT result = op(g_h_a, g_h_b);
    printf("template function result: %f\n", result);
}

int main(int argc, char* argv[]){
    hostOp<float>(min<float>);                          //works fine
    deviceOp<float><<<1,1>>>(fminf);

    cudaDeviceSynchronize(); 
    cout<<cudaGetErrorString(cudaGetLastError())<<endl;
}

输出:

host function result: 5.000000
static function result: 5.000000
an illegal memory access was encountered

假设我没有做一些非常愚蠢的事情，我应该如何将 fminf 传递给 deviceOp 以便不存在非法内存访问？

如果我正在做一些非常愚蠢的事情，什么是更好的方法？

Answer 1

要在设备上调用的函数必须用 __device__(或 __global__，如果您希望它是内核)装饰。 nvcc 编译器驱动程序将分离主机和设备代码，并在从设备代码调用(即编译)时使用函数的设备编译版本，否则使用主机版本。

这个构造是有问题的:

deviceOp<float><<<1,1>>>(fminf);

虽然可能并不明显，但这基本上都是主机代码。是的，它正在启动一个内核(通过来自主机代码的库调用的底层序列)，但它在技术上是主机代码。因此，此处“捕获”的 fminf 函数地址将是 fminf 函数的 host 版本，即使设备版本可用(通过 CUDA math.h，您实际上并未包括在内)。

解决此问题的典型方法是在设备代码中“捕获”设备地址，，然后将其作为参数传递给内核。

如果您传递的函数地址可以在编译时推导出来，您也可以(稍微)缩短此过程，使用稍微不同的模板技术。 this answer 中涵盖了这些概念.

这是使用“设备代码中的捕获函数地址”方法修改代码的完整示例:

$ cat t1176.cu
#include <iostream>
#include <cstdio>
#include <cmath>
using namespace std; //for brevity

__device__  float g_d_a = 9, g_d_b = 5;
float g_h_a = 9, g_h_b = 5;

template<typename argT, typename funcT>
__global__
void deviceOp(funcT op){
    argT result = fminf(g_d_a, g_d_b);                  //works fine
    printf("static function result: %f\n", result);
    result = op(g_d_a,g_d_b);                           //illegal memory access
    printf("template function result: %f\n", result);
}

__device__ float (*my_fminf)(float, float) = fminf;  // "capture" device function address

template<typename argT, typename funcT>
void hostOp(funcT op){
    argT result = op(g_h_a, g_h_b);
    printf("template function result: %f\n", result);
}

int main(int argc, char* argv[]){
    hostOp<float>(min<float>);                          //works fine
    float (*h_fminf)(float, float);
    cudaMemcpyFromSymbol(&h_fminf, my_fminf, sizeof(void *));
    deviceOp<float><<<1,1>>>(h_fminf);

    cudaDeviceSynchronize();
    cout<<cudaGetErrorString(cudaGetLastError())<<endl;
}
$ nvcc -o t1176 t1176.cu
$ cuda-memcheck ./t1176
========= CUDA-MEMCHECK
template function result: 5.000000
static function result: 5.000000
template function result: 5.000000
no error
========= ERROR SUMMARY: 0 errors
$