在cuda中作为模板参数

function as templated parameter in cuda

本文关键字：参数 cuda      更新时间：2023-10-16

我正在尝试在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: %fn", result);
    result = op(g_d_a,g_d_b);                           //illegal memory access
    printf("template function result: %fn", result);
}
template<typename argT, typename funcT>                 
void hostOp(funcT op){
    argT result = op(g_h_a, g_h_b);
    printf("template function result: %fn", 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，这样就不会有非法的内存访问?

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

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

这个结构有问题:

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

虽然可能不太明显，但这基本上是所有的主机代码。是的，它正在启动一个内核(通过宿主代码的底层库调用序列)，但从技术上讲，它是宿主代码。因此，这里"捕获"的fminf函数地址将是fminf函数的主机版本，即使设备版本可用(通过CUDA math.h，您实际上没有包括)。

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

如果您传递的函数地址可以在编译时推断出来，那么您也可以使用稍微不同的模板技术(在某种程度上)缩短此过程。这些概念将在本回答中讨论。

下面是一个使用"capture function address in device code"方法修改代码的完整示例:

$ 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: %fn", result);
    result = op(g_d_a,g_d_b);                           //illegal memory access
    printf("template function result: %fn", 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: %fn", 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
$
相关文章：