我不能从 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 函数，通常需要：

为 cc 3.5 或更高版本的设备编译
使用可重定位设备代码编译
link 针对 cuda 设备运行时库

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

此特定功能的注意事项：

CUDA 编译器中似乎存在问题，如果在内核代码中使用此设备运行时 API 调用而没有任何其他运行时 API 调用，则代码生成不会正确发生。此时的解决方法是确保您的内核至少包含一个其他 cuda 运行时 API 调用。在上面的示例中，我使用了 cudaGetErrorString，但您可以例如使用 cudaDeviceSynchronize() 或其他任何东西，我想。我已提交内部 NVIDIA 错误来报告此问题。
编程指南的 CDP 部分（link 以上）支持的设备运行时 API 调用列表中似乎存在文档错误。函数 cudaGetDeviceProperty 不存在，但我相信它应该指的是 cudaDeviceGetAttribute。我已针对此文档错误提交了 NVIDIA 内部错误。

我不能从 device 函数调用 host device 函数吗？

Can't I call a host device function from a device function?

cuda

gpgpu

我不能从 __device__ 函数调用 __host__ __device__ 函数吗？

Can't I call a __host__ __device__ function from a __device__ function?

cuda

gpgpu

我不能从 device 函数调用 host device 函数吗？

Can't I call a host device function from a device function?