cuda - CUPTI / CUDA prof_trigger 未按预期工作

Question

根据 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

每次我将线程数增加 32 时，prof_trigger_00 都应该增加（这需要一个新的扭曲）

不是这种情况。在我的笔记本电脑上运行时，使用 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;
        ............................

score 2 · Accepted Answer

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。

cuda - CUPTI / CUDA prof_trigger 未按预期工作

1 回答 1

Related

Reference