时间戳

首页 > 解决方案 > 一个正在运行的 CUDA 内核的名字可以通过它的线程来获取吗？

reflection - 一个正在运行的 CUDA 内核的名字可以通过它的线程来获取吗？

问题描述

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

foo运行的线程是否可以bar获取名称“foo”、签名或内核的其他标识符，最好是人类可读的标识符？

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

标签： reflectioncudagpgpu

解决方案

内核可以读取特殊寄存器%gridid。%gridid 每次启动都是唯一的。如果性能好，那么一个简单的内核 prolog 可以让每个内核启动的一个线程使用func和 %gridid 输出 gridid 全局函数映射。或者，可以使用 CUPTI SDK Activity 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

推荐阅读