Starten Cuda Anruf von struct

Question

eine einfache Struktur Aufgrund der CUDA-Code wickeln einen jedoch so etwas wie

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

schreiben kann, würde Ich mag den Block, Gitter, Shared-Memory-Berechnung in die Struktur setzen und rufen Sie die Kernel wie

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

Während die erste arbeitet, gibt der zweite mir eine illegal Speicherzugriffsfehler.

Ein minimales Beispiel mein Problem zu reproduzieren ist

#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; 
} 

Der Ausgang

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 

entspricht, sollten Sie nicht in beide Richtungen? Gibt es Alternativen, die auch die Shm, Grid-Berechnungen innerhalb der Struktur tun?

Answer 1

Es sei denn, Sie verwenden managed memory (die Sie nicht sind), ist es nicht legal ist Kernel-Parameter durch Verweis übergeben:

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

Als ich das Ampersand entfernen, Ihr Code ausgeführt wird, ohne Laufzeitfehler für mich gibt und vernünftige Ausgabe:

$ 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 
$ 

Beachten Sie, dass dies nicht wirklich Sinn machen:

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

und wirft eine Compiler-Warnung für mich.

Meinten Sie:

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