Định thời các phần khác nhau trong hạt nhân CUDA

Question

Tôi có hạt nhân CUDA gọi ra một chuỗi các chức năng của thiết bị.

Cách tốt nhất để có được thời gian thực hiện cho từng chức năng thiết bị là gì?

Cách tốt nhất để có được thời gian thực hiện cho một phần mã trong một trong các chức năng của thiết bị là gì?

Answer 1

Trong mã của riêng tôi, tôi sử dụng hàm clock() để nhận được thời gian chính xác. Để thuận tiện, tôi có macro

enum { 
    tid_this = 0, 
    tid_that, 
    tid_count 
    }; 
__device__ float cuda_timers[ tid_count ]; 
#ifdef USETIMERS 
#define TIMER_TIC clock_t tic; if (threadIdx.x == 0) tic = clock(); 
#define TIMER_TOC(tid) clock_t toc = clock(); if (threadIdx.x == 0) atomicAdd(&cuda_timers[tid] , (toc > tic) ? (toc - tic) : (toc + (0xffffffff - tic))); 
#else 
#define TIMER_TIC 
#define TIMER_TOC(tid) 
#endif 

Những sau đó có thể được sử dụng để công cụ mã thiết bị như sau:

__global__ mykernel (...) { 

    /* Start the timer. */ 
    TIMER_TIC 

    /* Do stuff. */ 
    ... 

    /* Stop the timer and store the results to the "timer_this" counter. */ 
    TIMER_TOC(tid_this); 

    } 

Sau đó bạn có thể đọc các cuda_timers trong các mã chủ.

Một vài lưu ý:

• Các tính giờ làm việc trên một cơ sở cho mỗi khối, ví dụ: nếu bạn có 100 khối thực hiện các hạt nhân tương tự, tổng của tất cả các lần họ sẽ được lưu trữ.
• Có nói rằng, bộ đếm thời gian giả định rằng chuỗi zeroth đang hoạt động, vì vậy hãy đảm bảo bạn không gọi các macro này trong một phần có thể khác nhau của mã.
• Bộ hẹn giờ đếm số lượng đồng hồ đếm. Để nhận được số mili giây, hãy chia số này cho số GHz trên thiết bị của bạn và nhân với 1000.
• Bộ hẹn giờ có thể làm chậm mã của bạn một chút, đó là lý do tại sao tôi bọc chúng trong #ifdef USETIMERS để bạn có thể tắt chúng dễ dàng.
• Mặc dù clock() trả về giá trị nguyên kiểu clock_t, tôi lưu trữ giá trị tích lũy là float, nếu không giá trị sẽ được bao quanh cho hạt dài hơn vài giây (tích lũy trên tất cả các khối).
• Lựa chọn (toc > tic) ? (toc - tic) : (toc + (0xffffffff - tic))) là cần thiết trong trường hợp bộ đếm đồng hồ kết thúc tốt đẹp xung quanh.

P.S. Đây là bản sao câu trả lời của tôi cho this question, không nhận được nhiều điểm ở đó vì thời gian bắt buộc dành cho toàn bộ hạt nhân.