从结构启动 Cuda 调用答案

【问题标题】：Launching Cuda Call from struct从结构启动 Cuda 调用
【发布时间】：2017-07-07 01:08:19
【问题描述】：

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

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

但是，我想将块、网格、共享内存计算放入结构中并像调用内核一样

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、网格计算？

【问题讨论】：

标签： c++ struct cuda nvcc

【解决方案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));

【讨论】：