Kann ich eine __host__ __device__ Funktion nicht von einer __device__ Funktion aufrufen?

Question

In der CUDA-Dokumentation habe ich festgestellt, dass cudaDeviceGetAttribute eine __host__ __device__-Funktion ist. Also dachte ich, ich könnte es in meiner __global__ Funktion nennen, um einige Attribute meines Gerätes zu bekommen. Leider scheint es etwas anderes zu bedeuten, weil ich ein Kompilierfehlerereignis erhalte, wenn ich es in eine __device__ Funktion bringe und diese aus meinem globalen herausrufe.

Ist es möglich, cudaDeviceGetAttribute auf meiner GPU aufzurufen? oder was bedeutet sonst __host__ __device__?

Hier ist meine Quellcode:

__device__ void GetAttributes(int* unique) 
{ 
    cudaDeviceAttr attr = cudaDevAttrMaxThreadsPerBlock; 
    cudaDeviceGetAttribute(unique, attr, 0); 
} 


__global__ void ClockTest(int* a, int* b, long* return_time, int* unique) 
{ 
    clock_t start = clock(); 

    //some complex calculations 

    *a = *a + *b; 
    *b = *a + *a; 

    GetAttributes(unique); 

    *a = *a + *b - *a; 

    clock_t end = clock(); 
    *return_time = end - start; 
} 

int main() 
{ 
    int a = 2; 
    int b = 3; 
    long time = 0; 
    int uni; 

    int* dev_a; 
    int* dev_b; 
    long* dev_time; 
    int* unique; 

    for (int i = 0; i < 10; ++i) { 

     cudaMalloc(&dev_a, sizeof(int)); 
     cudaMalloc(&dev_b, sizeof(int)); 
     cudaMalloc(&dev_time, sizeof(long)); 
     cudaMalloc(&unique, sizeof(int)); 

     cudaMemcpy(dev_a, &a, sizeof(int), cudaMemcpyHostToDevice); 
     cudaMemcpy(dev_b, &b, sizeof(int), cudaMemcpyHostToDevice); 

     ClockTest <<<1,1>>>(dev_a, dev_b, dev_time, unique); 

     cudaMemcpy(&a, dev_a, sizeof(int), cudaMemcpyDeviceToHost); 
     cudaMemcpy(&time, dev_time, sizeof(long), cudaMemcpyDeviceToHost); 
     cudaMemcpy(&uni, unique, sizeof(int), cudaMemcpyDeviceToHost); 

     cudaFree(&dev_a); 
     cudaFree(&dev_b); 
     cudaFree(&dev_time); 
     cudaFree(&unique); 

     printf("%d\n", time); 
     printf("unique: %d\n", uni); 

     cudaDeviceReset(); 
    } 

    return 0; 
} 

Answer 1

EDIT: sorry, meine Antwort war nicht korrekt. Es scheint ein Problem in nvcc (siehe unten) zu sein.

cudaDeviceGetAttribute korrekt in Gerätecode arbeiten, hier ist ein Beispiel arbeitete auf K20X, CUDA 8.0.61:

$ cat t1305.cu 
#include <stdio.h> 

__global__ void tkernel(){ 

    int val; 
    cudaError_t err = cudaDeviceGetAttribute(&val, cudaDevAttrMaxThreadsPerBlock, 0); 
    printf("err = %d, %s\n", err, cudaGetErrorString(err)); 
    printf("val = %d\n", val); 

} 


int main(){ 

    tkernel<<<1,1>>>(); 
    cudaDeviceSynchronize(); 
} 


$ nvcc -arch=sm_35 -o t1305 t1305.cu -rdc=true -lcudadevrt 
$ cuda-memcheck ./t1305 
========= CUDA-MEMCHECK 
err = 0, no error 
val = 1024 
========= ERROR SUMMARY: 0 errors 
$ 

Es gibt various runtime API functions supported for use in device code. Für die Laufzeit-API-Funktionen unterstützt, ist es in der Regel notwendig:

1. Kompilierung für eine cc 3.5 oder höher Gerät
2. mit Code relocatable Gerät kompilieren
3. Link gegen die Laufzeit cuda Gerätebibliothek

Außerdem hat Ihr Code einige andere Codierungsfehler, da wir die Adresse des Zeigers nicht an cudaFree übergeben, sondern nur an den Zeiger selbst.

Caveats für diese besondere Funktion:

1. Es gibt ein Problem in dem CUDA Compiler zu sein scheint, dass, wenn dieses Gerät-Laufzeit-API-Aufruf ohne weiteren Laufzeit-API-Aufruf in dem Kernel-Code verwendet wird, dann ist die Code- Generation wird nicht korrekt passieren. Die Problemumgehung zu diesem Zeitpunkt besteht darin, sicherzustellen, dass Ihr Kernel mindestens einen weiteren cuda-Laufzeit-API-Aufruf enthält. In meinem obigen Beispiel habe ich cudaGetErrorString verwendet, aber Sie könnten z.B. Verwenden Sie cudaDeviceSynchronize() oder etwas anderes, denke ich. Ich habe einen internen NVIDIA-Fehler gemeldet, um dieses Problem zu melden.

2. In der Liste der API-Aufrufe der Gerätelaufzeit, die im CDP-Abschnitt des Programmierhandbuchs (Link oben) unterstützt werden, liegt ein Dokumentationsfehler vor. Die Funktion cudaGetDeviceProperty existiert nicht, aber ich glaube, sie sollte sich auf cudaDeviceGetAttribute beziehen. Ich habe einen internen NVIDIA-Fehler für diesen Dokumentationsfehler eingereicht.