AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

NPKit: improve clock calibration and fix GPU clock API (#683)

Browse Source 
* Improve clock calibration in NPKit

* Improve gfx macro

* Fix macro
This commit is contained in:
Ziyue Yang
2023-02-18 03:26:57 +08:00
committed by GitHub
parent f7a456122c
commit f4bf47f325
12 changed files with 201 additions and 239 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/include/npkit/npkit.h
+12 -5
View File
@@ -14,6 +14,12 @@
#include "npkit/npkit_event.h"
#include "npkit/npkit_struct.h"
#if defined(__GFX9__) || defined(__GFX10__)
#define NPKIT_GET_GPU_TIMESTAMP __builtin_amdgcn_s_memrealtime
#else
#define NPKIT_GET_GPU_TIMESTAMP clock64
#endif
class NpKit {
public:
static const uint64_t kNumGpuEventBuffers = 512;
@@ -43,7 +49,7 @@ class NpKit {
static void CollectCpuEvent(uint8_t type, uint32_t size, uint32_t rsvd, uint64_t timestamp, int channel_id);
static uint64_t* GetCpuTimestamp();
static uint64_t GetCpuTimestamp();
private:
static void CpuTimestampUpdateThread();
@@ -59,12 +65,13 @@ class NpKit {
static NpKitEventCollectContext* gpu_collect_contexts_;
static NpKitEventCollectContext* cpu_collect_contexts_;
static uint64_t* cpu_timestamp_;
static uint64_t base_cpu_timestamp_global_;
static uint64_t base_cpu_timestamp_local_;
static uint64_t base_gpu_timestamp_cpu_;
static uint64_t base_gpu_timestamp_gpu_;
static uint64_t rank_;
static std::thread* cpu_timestamp_update_thread_;
static volatile bool cpu_timestamp_update_thread_should_stop_;
};
#endif
src/include/npkit/npkit_event.h
+27 -30
View File
@@ -65,39 +65,36 @@
#define NPKIT_EVENT_NET_RECV_ENTRY 0x31
#define NPKIT_EVENT_NET_RECV_EXIT 0x32
#define NPKIT_EVENT_TIME_SYNC_GPU 0x33
#define NPKIT_EVENT_TIME_SYNC_CPU 0x34
#define NPKIT_EVENT_ALL_REDUCE_RING_SEND_ENTRY 0x33
#define NPKIT_EVENT_ALL_REDUCE_RING_SEND_EXIT 0x34
#define NPKIT_EVENT_ALL_REDUCE_RING_RECV_REDUCE_SEND_ENTRY 0x35
#define NPKIT_EVENT_ALL_REDUCE_RING_RECV_REDUCE_SEND_EXIT 0x36
#define NPKIT_EVENT_ALL_REDUCE_RING_DIRECT_RECV_REDUCE_COPY_SEND_ENTRY 0x37
#define NPKIT_EVENT_ALL_REDUCE_RING_DIRECT_RECV_REDUCE_COPY_SEND_EXIT 0x38
#define NPKIT_EVENT_ALL_REDUCE_RING_DIRECT_RECV_COPY_SEND_ENTRY 0x39
#define NPKIT_EVENT_ALL_REDUCE_RING_DIRECT_RECV_COPY_SEND_EXIT 0x3A
#define NPKIT_EVENT_ALL_REDUCE_RING_DIRECT_RECV_ENTRY 0x3B
#define NPKIT_EVENT_ALL_REDUCE_RING_DIRECT_RECV_EXIT 0x3C
#define NPKIT_EVENT_ALL_REDUCE_RING_SEND_ENTRY 0x35
#define NPKIT_EVENT_ALL_REDUCE_RING_SEND_EXIT 0x36
#define NPKIT_EVENT_ALL_REDUCE_RING_RECV_REDUCE_SEND_ENTRY 0x37
#define NPKIT_EVENT_ALL_REDUCE_RING_RECV_REDUCE_SEND_EXIT 0x38
#define NPKIT_EVENT_ALL_REDUCE_RING_DIRECT_RECV_REDUCE_COPY_SEND_ENTRY 0x39
#define NPKIT_EVENT_ALL_REDUCE_RING_DIRECT_RECV_REDUCE_COPY_SEND_EXIT 0x3A
#define NPKIT_EVENT_ALL_REDUCE_RING_DIRECT_RECV_COPY_SEND_ENTRY 0x3B
#define NPKIT_EVENT_ALL_REDUCE_RING_DIRECT_RECV_COPY_SEND_EXIT 0x3C
#define NPKIT_EVENT_ALL_REDUCE_RING_DIRECT_RECV_ENTRY 0x3D
#define NPKIT_EVENT_ALL_REDUCE_RING_DIRECT_RECV_EXIT 0x3E
#define NPKIT_EVENT_ALL_REDUCE_TREE_UPDOWN_REDUCE_ENTRY 0x3D
#define NPKIT_EVENT_ALL_REDUCE_TREE_UPDOWN_REDUCE_EXIT 0x3E
#define NPKIT_EVENT_ALL_REDUCE_TREE_UPDOWN_BROADCAST_ENTRY 0x3F
#define NPKIT_EVENT_ALL_REDUCE_TREE_UPDOWN_BROADCAST_EXIT 0x40
#define NPKIT_EVENT_ALL_REDUCE_TREE_UPDOWN_REDUCE_ENTRY 0x3F
#define NPKIT_EVENT_ALL_REDUCE_TREE_UPDOWN_REDUCE_EXIT 0x40
#define NPKIT_EVENT_ALL_REDUCE_TREE_UPDOWN_BROADCAST_ENTRY 0x41
#define NPKIT_EVENT_ALL_REDUCE_TREE_UPDOWN_BROADCAST_EXIT 0x42
#define NPKIT_EVENT_ALL_REDUCE_TREE_SPLIT_REDUCE_BROADCAST_ENTRY 0x41
#define NPKIT_EVENT_ALL_REDUCE_TREE_SPLIT_REDUCE_BROADCAST_EXIT 0x42
#define NPKIT_EVENT_ALL_REDUCE_TREE_SPLIT_REDUCE_ENTRY 0x43
#define NPKIT_EVENT_ALL_REDUCE_TREE_SPLIT_REDUCE_EXIT 0x44
#define NPKIT_EVENT_ALL_REDUCE_TREE_SPLIT_BROADCAST_ENTRY 0x45
#define NPKIT_EVENT_ALL_REDUCE_TREE_SPLIT_BROADCAST_EXIT 0x46
#define NPKIT_EVENT_ALL_REDUCE_TREE_SPLIT_REDUCE_BROADCAST_ENTRY 0x43
#define NPKIT_EVENT_ALL_REDUCE_TREE_SPLIT_REDUCE_BROADCAST_EXIT 0x44
#define NPKIT_EVENT_ALL_REDUCE_TREE_SPLIT_REDUCE_ENTRY 0x45
#define NPKIT_EVENT_ALL_REDUCE_TREE_SPLIT_REDUCE_EXIT 0x46
#define NPKIT_EVENT_ALL_REDUCE_TREE_SPLIT_BROADCAST_ENTRY 0x47
#define NPKIT_EVENT_ALL_REDUCE_TREE_SPLIT_BROADCAST_EXIT 0x48
#define NPKIT_EVENT_SEND_RECV_LOCAL_COPY_ENTRY 0x47
#define NPKIT_EVENT_SEND_RECV_LOCAL_COPY_EXIT 0x48
#define NPKIT_EVENT_SEND_RECV_SEND_ENTRY 0x49
#define NPKIT_EVENT_SEND_RECV_SEND_EXIT 0x4A
#define NPKIT_EVENT_SEND_RECV_RECV_ENTRY 0x4B
#define NPKIT_EVENT_SEND_RECV_RECV_EXIT 0x4C
#define NPKIT_EVENT_SEND_RECV_LOCAL_COPY_ENTRY 0x49
#define NPKIT_EVENT_SEND_RECV_LOCAL_COPY_EXIT 0x4A
#define NPKIT_EVENT_SEND_RECV_SEND_ENTRY 0x4B
#define NPKIT_EVENT_SEND_RECV_SEND_EXIT 0x4C
#define NPKIT_EVENT_SEND_RECV_RECV_ENTRY 0x4D
#define NPKIT_EVENT_SEND_RECV_RECV_EXIT 0x4E
#define NPKIT_PRIM_COLLECT_DATA_PROCESS_TIME 0x4F
#define NPKIT_PRIM_COLLECT_DATA_PROCESS_TIME 0x4D
#endif