正在运行的CUDA内核的名称可以通过其线程获取吗？

Question

假设某些内核（名为__global__的foo函数）正在CUDA设备上运行。并假设内核调用了__device__函数bar，该函数有时在其他内核中被调用，即bar的代码在编译时不知道内核是foo还是其他。

foo中运行bar的线程能否获得名称“ foo”，签名或内核的其他某些标识符，最好是人类可读的标识符？

[如有必要，假定代码已使用--debug，--device-debug和/或--lineinfo中的任何一个进行编译。

Answer 1

内核可以读取特殊寄存器%gridid。 ％gridid在每次启动时都是唯一的。如果有性能，那么一个简单的内核序言可以从每个内核启动中使用一个线程，使用func和％gridid输出gridid全局函数映射。或者，可以使用CUPTI SDK活动API来收集此信息。 CUpti_ActivityKernel2事件包含每次启动的元数据，包括gridId和CUfunction名称。

这里是阅读％gridid的示例。

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <stdint.h>

cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size);

static __device__ __inline__ uint64_t __gridid()
{
    uint64_t gridid;
    asm volatile("mov.u64 %0, %%gridid;" : "=l"(gridid));
    return gridid;
}

__device__ void devPrintName()
{
    static const char* name = __func__;
    printf("%llu %s\n", __gridid(), name);
}

__global__ void globPrintName()
{
    static const char* name = __func__;
    printf("%llu %s\n", __gridid(), name);
    devPrintName();
}

int main()
{
    for (int i = 0; i < 4; ++i)
    {
        globPrintName<<<1,1,0>>>();
        cudaDeviceReset();
    }
    return 0;
}

此示例输出

1 globPrintName
1 devPrintName
2 globPrintName
2 devPrintName
3 globPrintName
3 devPrintName
4 globPrintName
4 devPrintName