c++ - 将内部函数作为模板参数传递

Question

我正在尝试将 atomicAdd 函数作为模板参数传递给另一个函数。

这是我的 Kernel1:

template<typename T, typename TAtomic>
__global__ void myfunc1(T *address, TAtomic atomicFunc) {
    atomicFunc(address, 1);
}

尝试 1:

myfunc1<<<1,1>>>(val.dev_ptr, atomicAdd);

它不起作用，因为编译器无法匹配预期的函数签名。

尝试 2:首先，我将 atomicAdd 包装到一个名为 MyAtomicAdd 的自定义函数中。

template<typename T>
__device__ void MyAtomicAdd(T *address, T val) {
    atomicAdd(address, val);
}

然后，我定义了一个名为“TAtomic”的函数指针，并将 TAtomic 声明为模板参数。

typedef void (*TAtomic)(float *,float);

template<typename T, TAtomic atomicFunc>
__global__ void myfunc2(T *address) {
    atomicFunc(address, 1);
}

myfunc2<float, MyAtomicAdd><<<1,1>>>(dev_ptr);
CUDA_CHECK(cudaDeviceSynchronize());

实际上，尝试 2 有效。但是，我不想使用 typedef。我需要更通用的东西。

尝试 3:只需将 MyAtomicAdd 传递给 myfunc1。

myfunc1<<<1,1>>>(dev_ptr, MyAtomicAdd<float>);
CUDA_CHECK(cudaDeviceSynchronize());

编译器可以编译代码。但是当我运行程序时，报错:

"ERROR in /home/liang/groute-dev/samples/framework/pagerank.cu:70: invalid program counter (76)"

我只是想知道，为什么尝试 3 不起作用？是否存在任何简单或温和的方法来实现此要求？谢谢。

Answer 1

尝试 3 不起作用，因为您试图在主机代码中获取 __device__ 函数的地址，这在 CUDA 中是非法的:

myfunc1<<<1,1>>>(dev_ptr, MyAtomicAdd<float>);
                          ^
                          effectively a function pointer - address of a __device__ function

CUDA 中的此类使用尝试将解析为某种“地址”——但它是垃圾，因此当您尝试将其用作设备代码中的实际函数入口点时，您会遇到遇到的错误:无效的程序计数器(或者在某些情况下，只是非法地址)。

您可以通过将内在函数包装在仿函数中而不是裸露的 __device__ 函数中来使您的 Try 3 方法工作(没有 typedef):

$ cat t48.cu
#include <stdio.h>

template<typename T>
__device__ void MyAtomicAdd(T *address, T val) {
    atomicAdd(address, val);
}


template <typename T>
struct myatomicadd
{
  __device__ T operator()(T *addr, T val){
    return atomicAdd(addr, val);
  }
};

template<typename T, typename TAtomic>
__global__ void myfunc1(T *address, TAtomic atomicFunc) {
    atomicFunc(address, (T)1);
}


int main(){

  int *dev_ptr;
  cudaMalloc(&dev_ptr, sizeof(int));
  cudaMemset(dev_ptr, 0, sizeof(int));
//  myfunc1<<<1,1>>>(dev_ptr, MyAtomicAdd<int>);
  myfunc1<<<1,1>>>(dev_ptr, myatomicadd<int>());
  int h = 0;
  cudaMemcpy(&h, dev_ptr, sizeof(int), cudaMemcpyDeviceToHost);
  printf("h = %d\n", h);
  return 0;
}
$ nvcc -arch=sm_35 -o t48 t48.cu
$ cuda-memcheck ./t48
========= CUDA-MEMCHECK
h = 1
========= ERROR SUMMARY: 0 errors
$

我们也可以实现一个稍微简单的版本，让仿函数模板类型从内核模板类型中推断出来:

$ cat t48.cu
#include <stdio.h>

struct myatomicadd
{
template <typename T>
  __device__ T operator()(T *addr, T val){
    return atomicAdd(addr, val);
  }
};

template<typename T, typename TAtomic>
__global__ void myfunc1(T *address, TAtomic atomicFunc) {
    atomicFunc(address, (T)1);
}


int main(){

  int *dev_ptr;
  cudaMalloc(&dev_ptr, sizeof(int));
  cudaMemset(dev_ptr, 0, sizeof(int));
  myfunc1<<<1,1>>>(dev_ptr, myatomicadd());
  int h = 0;
  cudaMemcpy(&h, dev_ptr, sizeof(int), cudaMemcpyDeviceToHost);
  printf("h = %d\n", h);
  float *dev_ptrf;
  cudaMalloc(&dev_ptrf, sizeof(float));
  cudaMemset(dev_ptrf, 0, sizeof(float));
  myfunc1<<<1,1>>>(dev_ptrf, myatomicadd());
  float hf = 0;
  cudaMemcpy(&hf, dev_ptrf, sizeof(float), cudaMemcpyDeviceToHost);
  printf("hf = %f\n", hf);
  return 0;
}
$ nvcc -arch=sm_35 -o t48 t48.cu
$ cuda-memcheck ./t48
========= CUDA-MEMCHECK
h = 1
hf = 1.000000
========= ERROR SUMMARY: 0 errors
$

有关 CUDA 中设备函数指针使用的更多处理链接到 this answer .