從結構

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