我無法從__device__函數調用__host__ __device__函數嗎？

Question

在CUDA文檔中，我發現cudaDeviceGetAttribute是__host__ __device__函數。所以我想我可以在我的__global__函數中調用它來獲取我的設備的一些屬性。可悲的是，它似乎意味着一些不同，因爲如果我將它放入__device__函數並從我的全局函數中調用此函數，我會收到一個編譯錯誤事件。

是否可以在我的GPU上調用cudaDeviceGetAttribute？或者__host__ __device__是什麼意思？

這裏是我的源代碼：

__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

編輯：對不起，我以前的答案是不正確的。在nvcc似乎有問題（見下文）。

cudaDeviceGetAttribute可以在設備代碼中正常工作，這裏是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 
$ 

有various runtime API functions supported for use in device code。 對於支持的運行時的API函數，這是通常必要的：

1. 編譯用於CC 3.5或更高的器件
2. 與重定位裝置代碼編譯針對CUDA設備運行庫

鏈路

此外，您的代碼還有一些其他編碼錯誤，因爲我們不會將指針的地址傳遞給cudaFree，而只是指針本身。

對於該特定功能注意事項：

1. 似乎有在，如果該設備運行時的API調用，而不在內核代碼的任何其它運行時API調用中使用的CUDA編譯器的一個問題，則代碼世代不會正確發生。此時的解決方法是確保您的內核至少包含一個其他cuda運行時API調用。在我上面的示例中，我使用cudaGetErrorString，但您可以使用cudaDeviceSynchronize()或其他任何東西，我想。我提交了一個內部的NVIDIA錯誤報告這個問題。

2. 在編程指南的CDP部分支持的設備運行時API調用列表中似乎存在文檔錯誤（上面的鏈接）。函數cudaGetDeviceProperty不存在，但我相信它應該參考cudaDeviceGetAttribute。我已經提交了一個內部NVIDIA錯誤的文檔錯誤。