c++ - 从结构启动 Cuda Call

Question

给定一个简单的结构来包装 cuda 代码，我们可以编写类似的东西

func<float> s;
s.val = 3.f;
start_correct<<<1, 2>>>(s);

不过，我想把block，grid，shared memory computation放到struct里面调用kernel like

func<float> s;
s.val = 3.f;
s.launch();

当第一个工作正常时，第二个给了我一个非法内存访问错误。

重现我的问题的一个最小例子是

#include <stdio.h>

template<typename T>
struct func;

template<typename T>
__global__ void start(const func<T>& s){
  printf("host access val %f \n",s.val);
  s();
}

template<typename T>
struct func
{
  T val;

  __device__ void operator()() const{
    printf("device access val %f [%d]\n",val,threadIdx.x);
  }

  enum{ C_N = 2 };

  void launch()
  {
    start<<<1, C_N>>>(*this);
  }

};

template<typename T>
__global__ void start_correct(const func<T> s){
  printf("host access val %f \n", s.val);
  s();
}

int main(int argc, char const *argv[])
{
  cudaError_t err;

  func<float> s;
  s.val = 3.f;

  // launch cuda kernel <-- WORKS
  start_correct<<<1, 2>>>(s);
  cudaDeviceSynchronize();
  if (err != cudaSuccess) printf("Error: %s\n", cudaGetErrorString(err));


  // launch cuda kernel <-- DOES NOT WORK
  s.launch();
  cudaDeviceSynchronize();
  err = cudaGetLastError();
  if (err != cudaSuccess) printf("Error: %s\n", cudaGetErrorString(err));


  return 0;
}

输出是

host access val 3.000000 
host access val 3.000000 
device access val 3.000000 [0]
device access val 3.000000 [1]
host access val 0.000000 
host access val 0.000000 
device access val 0.000000 [0]
device access val 0.000000 [1]
Error: an illegal memory access was encountered

这两种方式不应该是等价的吗？是否有任何替代方案，也可以在结构内部进行 shm、网格计算？

Answer 1

除非您使用 managed memory (你不是)，通过引用传递内核参数是不合法的:

__global__ void start(const func<T>& s){
                                   ^

当我删除该符号时，您的代码运行时对我来说没有任何运行时错误，并提供合理的输出:

$ cuda-memcheck ./t355
========= CUDA-MEMCHECK
host access val 3.000000
host access val 3.000000
device access val 3.000000 [0]
device access val 3.000000 [1]
host access val 3.000000
host access val 3.000000
device access val 3.000000 [0]
device access val 3.000000 [1]
========= ERROR SUMMARY: 0 errors
$

请注意，这实际上没有意义:

  cudaDeviceSynchronize();
  if (err != cudaSuccess) printf("Error: %s\n", cudaGetErrorString(err));

并为我抛出一个编译器警告。

也许你的意思是:

  err = cudaDeviceSynchronize();
  if (err != cudaSuccess) printf("Error: %s\n", cudaGetErrorString(err));