CUPTI / CUDA prof_trigger 未按预期工作

Posted

技术标签:

【中文标题】CUPTI / CUDA prof_trigger 未按预期工作【英文标题】:CUPTI / CUDA prof_trigger not working as expected 【发布时间】:2012-07-01 02:14:50 【问题描述】:

根据 NVidia,__prof_trigger() 调用是在 warp 级别执行的,或者换句话说,每个 warp 都会将指定的触发器增加 1。

所以我写了一个小内核来测试一下:

__global__ void kernel(int *arr) 
   __prof_trigger(00);

   // from here, it's irrevelant to the question
   int id = threadIdx.x + blockDim.x * blockIdx.x;
   if (id >= N) return;
   __prof_trigger(01);
   if (arr[id] < 4) __prof_trigger(02);
   else             __prof_trigger(03);

我使用以下方法调用程序: ./prof_trigger_test [block_size] [event_name]

此时输入数组无关紧要(我只是在测试 prof_trigger_0,所以它甚至没有被使用)

我只测试了一个块(据我的理解,不同数量的块不应该影响这个问题,还是我错了?)并且事件名称是 CUPTI 给出的名称,或 prof_trigger_XX,它被翻译为运行时的事件 ID。

因此,鉴于经纱大小为 32,我应该期待类似:

./prof_trigger_test 1  prof_trigger_00   // expected to return 1
./prof_trigger_test 33 prof_trigger_00   // expected to return 2

prof_trigger_00 应该在我每次将线程数增加 32 时增加(这需要一个新的 warp)

事实并非如此。在我的笔记本电脑上运行时,使用 NVidia 9600M GT 而不是 32,我只需增加 4 即可看到计数器的增量值。并在远程集群上运行,使用 Tesla M2070,所需增量为 8

很明显我错过了一些东西。 GPU 是否出于某种原因创建了更小的扭曲(性能,我认为,虽然不确定如何)?

这里是完整的代码,根据要求:

#include <cuda.h>
#include <cupti.h>
#include <stdio.h>

#define N 10

#define CHECK_CU_ERROR(err, cufunc)                                     \
    if (err != CUDA_SUCCESS)                                           \
        printf("%s:%d: error %d for CUDA Driver API function '%s'\n",   \
                __FILE__, __LINE__, err, cufunc);                       \
        exit(-1);                                                       \
    

#define CHECK_CUPTI_ERROR(err, cuptifunc)                               \
    if (err != CUPTI_SUCCESS)                                          \
        const char *errstr;                                             \
        cuptiGetResultString(err, &errstr);                             \
        printf("%s:%d:Error %s for CUPTI API function '%s'\n",          \
                __FILE__, __LINE__, errstr, cuptifunc);                 \
        exit(-1);                                                       \
    

typedef struct cupti_eventData_st 
    CUpti_EventGroup eventGroup;
    CUpti_EventID eventId;
 cupti_eventData;

// Structure to hold data collected by callback
typedef struct RuntimeApiTrace_st 
    cupti_eventData *eventData;
    uint64_t eventVal;
 RuntimeApiTrace_t;

void CUPTIAPI getEventValueCallback(
                        void *userdata,
                        CUpti_CallbackDomain domain,
                        CUpti_CallbackId cbid,
                        const CUpti_CallbackData *cbInfo) 

    CUptiResult cuptiErr;
    RuntimeApiTrace_t *traceData = (RuntimeApiTrace_t*) userdata;
    size_t bytesRead;

    // This callback is enabled for launch so we shouldn't see anything else.
    if (cbid != CUPTI_RUNTIME_TRACE_CBID_cudaLaunch_v3020) 
        printf("%s:%d: unexpected cbid %d\n", __FILE__, __LINE__, cbid);
        exit(-1);
    

    switch(cbInfo->callbackSite) 
        case CUPTI_API_ENTER:
            cudaThreadSynchronize();
            cuptiErr = cuptiSetEventCollectionMode(cbInfo->context, CUPTI_EVENT_COLLECTION_MODE_KERNEL);
            CHECK_CUPTI_ERROR(cuptiErr, "cuptiSetEventCollectionMode");
            cuptiErr = cuptiEventGroupEnable(traceData->eventData->eventGroup);
            CHECK_CUPTI_ERROR(cuptiErr, "cuptiEventGroupEnable");
            break;

        case CUPTI_API_EXIT:
            bytesRead = sizeof(uint64_t);
            cudaThreadSynchronize();
            cuptiErr = cuptiEventGroupReadEvent(traceData->eventData->eventGroup, CUPTI_EVENT_READ_FLAG_NONE, traceData->eventData->eventId, &bytesRead, &traceData->eventVal);
            CHECK_CUPTI_ERROR(cuptiErr, "cuptiEventGroupReadEvent");
            cuptiErr = cuptiEventGroupDisable(traceData->eventData->eventGroup);
            CHECK_CUPTI_ERROR(cuptiErr, "cuptiEventGroupDisable");
            break;
    


static void displayEventVal(RuntimeApiTrace_t *trace, char *eventName) 
    printf("Event Name: %s \n", eventName);
    printf("Event Value: %llu\n", (unsigned long long) trace->eventVal);


__global__ void kernel(int *arr) 
    __prof_trigger(00);

    int id = threadIdx.x + blockDim.x * blockIdx.x;
    if (id >= N) return;

    __prof_trigger(01);

    if (arr[id] < 4) __prof_trigger(02);
    else             __prof_trigger(03);


int main(int argc, char **argv) 
    int deviceCount;
    CUcontext context = 0;
    CUdevice dev = 0;
    char deviceName[32];
    char *eventName;
    CUptiResult cuptiErr;
    CUpti_SubscriberHandle subscriber;
    cupti_eventData cuptiEvent;
    RuntimeApiTrace_t trace;
    int cap_major, cap_minor;

    CUresult err = cuInit(0);
    CHECK_CU_ERROR(err, "cuInit");

    err = cuDeviceGetCount(&deviceCount);
    CHECK_CU_ERROR(err, "cuDeviceGetCount");

    if (deviceCount == 0) 
        printf("There is no device supporting CUDA.\n");
        return -2;
    

    if (argc < 3) 
        printf("Usage: ./a.out <num_threads> <event_name>\n");
        return -2;
    

    err = cuDeviceGet(&dev, 0);
    CHECK_CU_ERROR(err, "cuDeviceGet");

    err = cuDeviceGetName(deviceName, 32, dev);
    CHECK_CU_ERROR(err, "cuDeviceGetName");

    err = cuDeviceComputeCapability(&cap_major, &cap_minor, dev);
    CHECK_CU_ERROR(err, "cuDeviceComputeCapability");

    printf("CUDA Device Name: %s\n", deviceName);
    printf("CUDA Capability: %d.%d\n", cap_major, cap_minor);

    err = cuCtxCreate(&context, 0, dev);
    CHECK_CU_ERROR(err, "cuCtxCreate");

    cuptiErr = cuptiEventGroupCreate(context, &cuptiEvent.eventGroup, 0);
    CHECK_CUPTI_ERROR(cuptiErr, "cuptiEventGroupCreate");

    int threads = atoi(argv[1]);
    eventName = argv[2];

    cuptiErr = cuptiEventGetIdFromName(dev, eventName, &cuptiEvent.eventId);
    if (cuptiErr != CUPTI_SUCCESS) 
        printf("Invalid eventName: %s\n", eventName);
        return -1;
    

    cuptiErr = cuptiEventGroupAddEvent(cuptiEvent.eventGroup, cuptiEvent.eventId);
    CHECK_CUPTI_ERROR(cuptiErr, "cuptiEventGroupAddEvent");

    trace.eventData = &cuptiEvent;

    cuptiErr = cuptiSubscribe(&subscriber, (CUpti_CallbackFunc)getEventValueCallback, &trace);
    CHECK_CUPTI_ERROR(cuptiErr, "cuptiSubscribe");

    cuptiErr = cuptiEnableCallback(1, subscriber, CUPTI_CB_DOMAIN_RUNTIME_API, CUPTI_RUNTIME_TRACE_CBID_cudaLaunch_v3020);
    CHECK_CUPTI_ERROR(cuptiErr, "cuptiEnableCallback");


    int host_arr[] = 0, 1, 2, 3, 4, 5, 6, 7, 8, 9;
    int *dev_arr;


    cudaMalloc(&dev_arr, sizeof(int) * N);
    cudaMemcpy(dev_arr, &host_arr, sizeof(int) * N, cudaMemcpyHostToDevice);
    kernel<<< threads, 1 >>>(dev_arr);

    displayEventVal(&trace, eventName);
    trace.eventData = NULL;

    cuptiErr = cuptiEventGroupRemoveEvent(cuptiEvent.eventGroup, cuptiEvent.eventId);
    CHECK_CUPTI_ERROR(cuptiErr, "cuptiEventGroupRemoveEvent");

    cuptiErr = cuptiEventGroupDestroy(cuptiEvent.eventGroup);
    CHECK_CUPTI_ERROR(cuptiErr, "cuptiEventGroupDestroy");

    cuptiErr = cuptiUnsubscribe(subscriber);
    CHECK_CUPTI_ERROR(cuptiErr, "cuptiUnsubscribe");

    cudaDeviceSynchronize();

以及 SASS 代码:

Fatbin ptx code:
================
arch = sm_10
code version = [1,4]
producer = cuda
host = linux
compile_size = 64bit
identifier = cuda.cu

Fatbin elf code:
================
arch = sm_10
code version = [1,2]
producer = cuda
host = linux
compile_size = 64bit
identifier = cuda.cu

    code for sm_10
        Function : _Z6kernelPi
    /*0000*/ ;
    /*0008*/     /*0xf0000001e0000004*/     MOV.U16 R0H, g [0x1].U16;
    /*0010*/     /*0x100042050023c780*/     I2I.U32.U16 R1, R0L;
    /*0018*/     /*0xa000000504000780*/     IMAD.U16 R0, g [0x6].U16, R0H, R1;
    /*0020*/     /*0x60014c0100204780*/     ISET.S32.C0 o [0x7f], R0, c [0x1] [0x0], LE;
    /*0028*/     /*0x308001fd6c40c7c8*/     RET C0.EQU;
    /*0030*/ ;
    /*0038*/     /*0x3000000300000500*/     SHL R0, R0, 0x2;
    /*0040*/     /*0xf0000401e0000004*/     IADD R0, g [0x4], R0;
    /*0048*/     /*0x30020001c4100780*/     GLD.U32 R0, global14 [R0];
    /*0050*/     /*0x2000c80104200780*/     ISET.S32.C0 o [0x7f], R0, c [0x1] [0x1], GT;
    /*0058*/     /*0xd00e000180c00780*/     BRA C0.NE, 0x70;
    /*0060*/ ;
    /*0068*/     /*0x308101fd6c4107c8*/     RET;
    /*0070*/ ;
    /*0078*/     /*0x1000e00300000280*/     NOP;
        ............................

【问题讨论】:

如果我们不知道命令行选项的含义,那么为您的程序提供命令行是没有用的。您也没有向您的程序提供有关输入数据的信息。更完整的代码(主机代码)会有所帮助。 抱歉,这篇文章吃掉了我解释命令行用法的部分(禁止字符)。我现在修复了它,但我正在使用 ./prof_trigger_test [block_size] [event_name] 调用程序。输入数组无关紧要,但它是一个包含 10 个元素的数组(N=10) 主机代码仅包括创建数组、CUPTI 调用以检索计数器值(基于官方样本)和内核调用 请同时发布主机代码、SASS 代码 (cuobjdump -sass) 和驱动程序版本。编译器倾向于移动和消除 __prof_triggers。旧版本的 CUPTI 只收集 1 个 SM 的结果。最新版本可以收集所有 SM 的结果。 原始帖子已编辑,根据要求添加主机代码和 SASS 转储。也忘了提,但是除了所需的链接器和包含 CUDA 和 CUPTI 的路径之外,所有东西都是在没有优化标志或其他任何东西的情况下编译的 【参考方案1】:

prof_trigger_XX 的集合因架构而异。在计算能力为 1.* 的设备上,计数器值仅从 1 个 SM 收集。在计算能力 >= 2 时,从所有 SM 收集计数器值。

GeForce 9600M GPU 是具有 4 个 SM 的 1.1 设备。

示例程序正在启动 [block_size] 块,每个块有 1 个线程。这有点令人困惑,因为我认为您的意思是启动 N 个线程的示例,每个块最多 1 个扭曲。这不是示例代码中实现的内容。

[block_size]  prof_trigger_00 expected range
1-3           0-1
4-7           1-2
8-11          2-3

预期值是一个范围,因为无法保证观察到的 SM 是否会是第一个接收到被调度线程块的 SM。

【讨论】:

以上是关于CUPTI / CUDA prof_trigger 未按预期工作的主要内容,如果未能解决你的问题,请参考以下文章

TensorFlow couldn t open CUDA library cupti64 80 dll Intern

anaconda env:cupti64_100.dll找不到

cuda11,cudnn8.1,cuda10.2

CUDA - CUDA 驱动程序不足以支持 CUDA 运行时版本 6.0

CUDA 核函数调试:cuda-gdb,cuda printf, cuda assert 读书笔记

cuda11.7显卡要求