从结构

Question

启动Cuda的调用中给出一个简单的结构来包装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，grid的计算？

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));