我不能从 __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. 使用可重定位设备代码编译
3. 链接到 cuda 设备运行时库

此外，您的代码还有一些其他编码错误，因为我们没有将指针的地址传递给cudaFree，而只是传递了指针本身。

此特定功能的注意事项：

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

2. 编程指南的 CDP 部分支持的设备运行时 API 调用列表中似乎存在文档错误（上面的链接）。 cudaGetDeviceProperty这个函数是不存在的，不过我相信应该是指cudaDeviceGetAttribute。我已为此文档错误提交了一个内部 NVIDIA 错误。