cuda - CUPTI / CUDA の prof_trigger が期待どおりに機能しない

Question

NVidia によると、__prof_trigger() 呼び出しはワープレベルで実行されます。つまり、各ワープは指定されたトリガーを 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 によって指定された名前であり、ランタイム。

ワープサイズが 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 増やすたびに増やす必要があります (これには新しいワープが必要です)

これはそうではありません。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デバイスです。

サンプルプログラムは、それぞれ1スレッドの[block_size]ブロックを起動しています。これは、各ブロックが最大で1ワープのNスレッドを起動するサンプルを意味していると思うので、少し混乱します。これは、サンプルコードで実装されたものではありません。

[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