我不能从__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
$

有支持的设备代码中使用的各种运行时API函数.对于受支持的运行时API函数,通常需要:

1. 编译cc 3.5或更高版本的设备
2. 使用可重定位的设备代码进行编译
3. 链接到cuda设备运行时库

此外,您的代码还有一些其他编码错误,因为我们不会将指针的地址传递给cudaFree指针本身.

这个特殊功能的注意事项:

1. 在CUDA编译器中似乎存在一个问题,即如果在内核代码中使用此设备运行时API调用而没有任何其他运行时API调用,则代码生成将不会正确发生.此时的解决方法是确保您的内核至少包含一个其他cuda运行时API调用.在我上面使用的例子中,我想cudaGetErrorString,你可以使用cudaDeviceSynchronize()或其他任何东西.我已提交内部NVIDIA错误报告此问题.

2. 编程指南的CDP部分支持的设备运行时API调用列表中似乎存在文档错误(上面的链接).该功能cudaGetDeviceProperty不存在,但我相信它应该参考cudaDeviceGetAttribute.我已针对此文档错误提交了内部NVIDIA错误.