AI/rocm-systems
2
0
Çatalla 0
Kod Konular Değişiklik İstekleri İşlemler Paketler Projeler Sürüm Viki Aktivite

PC-Sampling - gfx94x Hosttrap method support

Kaynağa Gözat 
Supports PC-Sampling on gfx94x in both CPX and SPX mode

Change-Id: Ife1e50ab08155678ea4aa2b80475b9974812c40e


[ROCm/ROCR-Runtime commit: fda2a395a3]
Bu işleme şunda yer alıyor:
Shweta.Khatri
2024-08-16 15:40:17 +00:00
işlemeyi yapan: David Yat Sin
ebeveyn 84135d4f49
işleme ce616cce30
3 değiştirilmiş dosya ile 175 ekleme ve 34 silme
Yama Dosyasını İndir Diff Dosyasını İndir Tüm dosyaları genişlet Tüm dosyaları daralt
projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_gpu_pm4.h
+5
Dosyayı Görüntüle
@@ -51,6 +51,7 @@
#define PM4_HDR_IT_OPCODE_ACQUIRE_MEM 0x58
#define PM4_HDR_IT_OPCODE_ATOMIC_MEM 0x1E
#define PM4_HDR_IT_OPCODE_PRED_EXEC 0x23
#define PM4_HDR_IT_OPCODE_WRITE_DATA 0x37
#define PM4_HDR_IT_OPCODE_WAIT_REG_MEM 0x3C
#define PM4_HDR_IT_OPCODE_COPY_DATA 0x40
@@ -97,6 +98,10 @@
#define PM4_ATOMIC_MEM_DW4_SRC_DATA_LO(x) (((x) & 0xFFFFFFFF) << 0)
#define PM4_ATOMIC_MEM_DW5_SRC_DATA_HI(x) (((x) & 0xFFFFFFFF) << 0)
#define PM4_PRED_EXEC_DW1_HEADER(x) (((x) & 0xFFFFFFFF) << 0)
#define PM4_PRED_EXEC_DW2_EXEC_COUNT(x) (((x) & 0x3FFF) << 0)
#define PM4_PRED_EXEC_DW2_VIRTUALXCCID_SELECT(x) (((x) & 0xFF) << 24)
#define PM4_COPY_DATA_DW1(x) (((x) & 0xFFFFFFFF) << 0)
# define PM4_COPY_DATA_SRC_SEL_ATOMIC_RETURN_DATA (6 << 0)
# define PM4_COPY_DATA_DST_SEL_TC_12 (2 << 8)
projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp
+31 -10
Dosyayı Görüntüle
@@ -52,6 +52,7 @@
#include <memory>
#include <utility>
#include <iomanip>
#include <cmath>
#include "core/inc/amd_aql_queue.h"
#include "core/inc/amd_blit_kernel.h"
@@ -2846,7 +2847,7 @@ hsa_status_t GpuAgent::PcSamplingFlushHostTrapDeviceBuffers(
uint32_t next_buffer;
uint64_t reset_write_val;
uint32_t to_copy, copy_bytes;
uint32_t to_copy = 0, copy_bytes;
const uint32_t atomic_ex_cmd_sz = 9;
const uint32_t wait_reg_mem_cmd_sz = 7;
@@ -2854,6 +2855,8 @@ hsa_status_t GpuAgent::PcSamplingFlushHostTrapDeviceBuffers(
const uint32_t copy_data_cmd_sz = 6;
const uint32_t write_data_cmd_sz = 5;
uint32_t pred_exec_cmd_sz = 0;
uint8_t* host_buffer_begin = ht_data.host_buffer;
uint8_t* host_buffer_end = ht_data.host_buffer + ht_data.host_buffer_size;
@@ -2872,6 +2875,15 @@ hsa_status_t GpuAgent::PcSamplingFlushHostTrapDeviceBuffers(
next_buffer = (which_buffer + 1) % 2;
reset_write_val = (uint64_t)next_buffer << 63;
unsigned int i = 0;
memset(cmd_data, 0, cmd_data_sz);
if (properties_.NumXcc > 1) {
pred_exec_cmd_sz = 2;
cmd_data[i++] = PM4_HDR(PM4_HDR_IT_OPCODE_PRED_EXEC, pred_exec_cmd_sz, isa_->GetMajorVersion());
cmd_data[i++] = PM4_PRED_EXEC_DW2_EXEC_COUNT(0xF) | PM4_PRED_EXEC_DW2_VIRTUALXCCID_SELECT(0x1);
}
/*
* ATOMIC_MEM, perform atomic_exchange
* We use a double-buffer mechanism so that trap handlers calls are writing to one buffer while
@@ -2884,8 +2896,7 @@ hsa_status_t GpuAgent::PcSamplingFlushHostTrapDeviceBuffers(
* i.e., what value to wait for in buf_written_val to know all previous trap entries were
* done.
*/
unsigned int i = 0;
memset(cmd_data, 0, cmd_data_sz);
cmd_data[i++] = PM4_HDR(PM4_HDR_IT_OPCODE_ATOMIC_MEM, atomic_ex_cmd_sz, isa_->GetMajorVersion());
cmd_data[i++] = PM4_ATOMIC_MEM_DW1_ATOMIC(PM4_ATOMIC_MEM_GL2_OP_ATOMIC_SWAP_RTN_64);
cmd_data[i++] = PM4_ATOMIC_MEM_DW2_ADDR_LO(buf_write_val);
@@ -2905,8 +2916,8 @@ hsa_status_t GpuAgent::PcSamplingFlushHostTrapDeviceBuffers(
HSA::hsa_signal_store_screlease(exec_pm4_signal, 1);
queues_[QueuePCSampling]->ExecutePM4(
cmd_data, (atomic_ex_cmd_sz + copy_data_cmd_sz) * sizeof(uint32_t), HSA_FENCE_SCOPE_NONE,
HSA_FENCE_SCOPE_SYSTEM, &exec_pm4_signal);
cmd_data, (pred_exec_cmd_sz + atomic_ex_cmd_sz + copy_data_cmd_sz) * sizeof(uint32_t),
HSA_FENCE_SCOPE_NONE, HSA_FENCE_SCOPE_SYSTEM, &exec_pm4_signal);
do {
hsa_signal_value_t val = HSA::hsa_signal_wait_scacquire(
exec_pm4_signal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_BLOCKED);
@@ -2931,6 +2942,18 @@ hsa_status_t GpuAgent::PcSamplingFlushHostTrapDeviceBuffers(
ht_data.host_buffer_wrap_pos = ht_data.host_write_ptr;
ht_data.host_write_ptr = host_buffer_begin;
}
i = 0;
memset(cmd_data, 0, cmd_data_sz);
if (properties_.NumXcc > 1) {
const uint32_t n = ceil(to_copy / (32 * 1024 * 1024));
pred_exec_cmd_sz = 2;
cmd_data[i++] = PM4_HDR(PM4_HDR_IT_OPCODE_PRED_EXEC, pred_exec_cmd_sz, isa_->GetMajorVersion());
cmd_data[i++] =
PM4_PRED_EXEC_DW2_EXEC_COUNT(0x13 + 7 * n) | PM4_PRED_EXEC_DW2_VIRTUALXCCID_SELECT(0x1);
}
/*
* Do the WAIT_REG_MEM, DMA_DATA(s) and WRITE_DATA
*
@@ -2940,8 +2963,6 @@ hsa_status_t GpuAgent::PcSamplingFlushHostTrapDeviceBuffers(
* 3. Reset buf_written_val so that we start writing to beginning of this buffer on the next
* buffer swap.
*/
i = 0;
memset(cmd_data, 0, cmd_data_sz);
/* WAIT_REG_MEM, wait on buf_written_val */
cmd_data[i++] =
@@ -2958,6 +2979,7 @@ hsa_status_t GpuAgent::PcSamplingFlushHostTrapDeviceBuffers(
unsigned int num_copy_command = 0;
uint8_t* buffer_temp = buffer[which_buffer];
for (copy_bytes = CP_DMA_DATA_TRANSFER_CNT_MAX; 0 < to_copy; to_copy -= copy_bytes) {
num_copy_command++;
@@ -2977,7 +2999,6 @@ hsa_status_t GpuAgent::PcSamplingFlushHostTrapDeviceBuffers(
} else {
cmd_data[i++] = PM4_DMA_DATA_DW6(PM4_DMA_DATA_BYTE_COUNT(copy_bytes) | PM4_DMA_DATA_DIS_WC);
}
buffer_temp += copy_bytes;
ht_data.host_write_ptr += copy_bytes;
}
@@ -2990,8 +3011,8 @@ hsa_status_t GpuAgent::PcSamplingFlushHostTrapDeviceBuffers(
cmd_data[i++] = PM4_WRITE_DATA_DW3_DST_MEM_ADDR_HI((buf_written_val[which_buffer]) >> 32);
cmd_data[i++] = PM4_WRITE_DATA_DW4_DATA(0);
unsigned int cmd_sz =
wait_reg_mem_cmd_sz + (num_copy_command * dma_data_cmd_sz) + write_data_cmd_sz;
unsigned int cmd_sz = pred_exec_cmd_sz + wait_reg_mem_cmd_sz +
(num_copy_command * dma_data_cmd_sz) + write_data_cmd_sz;
HSA::hsa_signal_store_screlease(exec_pm4_signal, 1);
queues_[QueuePCSampling]->ExecutePM4(cmd_data, cmd_sz * sizeof(uint32_t), HSA_FENCE_SCOPE_NONE,
projects/rocr-runtime/runtime/hsa-runtime/core/runtime/trap_handler/trap_handler.s
+139 -24
Dosyayı Görüntüle
@@ -3,7 +3,7 @@
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2022, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2014-2024, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
@@ -91,6 +91,10 @@
.set SQ_WAVE_IB_STS_REPLAY_W64H_MASK , 0x2000000
.endif
// Defining TTMP_REG1 and TTMP_REG2 for clarity in comments
// TTMP_REG1 means ttmp6 register if gfx>=940 and means ttmp13 register if gfx<940
// TTMP_REG2 means ttmp11 register if gfx>=940 and means ttmp6 register if gfx<940
.if .amdgcn.gfx_generation_number == 9 && .amdgcn.gfx_generation_minor >= 4
.set TTMP11_TTMPS_SETUP_SHIFT , 31
@@ -99,6 +103,98 @@
.set TTMP13_PCS_IS_STOCHASTIC , 24
.endif
.if (.amdgcn.gfx_generation_number == 9)
.macro S_LOAD_DWORD_PCS_TTMP_REG1 base, offset
.if (.amdgcn.gfx_generation_minor >= 4)
s_load_dword ttmp6, \base, \offset
.else
s_load_dword ttmp13,\base, \offset
.endif
.endm
.macro S_BITSET0_B32_PCS_TTMP_REG2 bit_index
.if (.amdgcn.gfx_generation_minor >= 4)
s_bitset0_b32 ttmp11, \bit_index
.else
s_bitset0_b32 ttmp6, \bit_index
.endif
.endm
.macro S_BITSET1_B32_PCS_TTMP_REG2 bit_index
.if (.amdgcn.gfx_generation_minor >= 4)
s_bitset1_b32 ttmp11, \bit_index
.else
s_bitset1_b32 ttmp6, \bit_index
.endif
.endm
.macro S_CMP_GE_U32_PCS_TTMP_REG1 src0
.if (.amdgcn.gfx_generation_minor >= 4)
s_cmp_ge_u32 \src0, ttmp6
.else
s_cmp_ge_u32 \src0, ttmp13
.endif
.endm
.macro S_MOV_B32_SRC_PCS_TTMP_REG1 src0
.if (.amdgcn.gfx_generation_minor >= 4)
s_mov_b32 ttmp6, \src0
.else
s_mov_b32 ttmp13, \src0
.endif
.endm
.macro S_MOV_B32_DST_PCS_TTMP_REG1 dst
.if (.amdgcn.gfx_generation_minor >= 4)
s_mov_b32 \dst, ttmp6
.else
s_mov_b32 \dst, ttmp13
.endif
.endm
.macro S_LSHR_B32_PCS_TTMP_REG1_REG2 src1
.if (.amdgcn.gfx_generation_minor >= 4)
s_lshr_b32 ttmp6, ttmp11, \src1
.else
s_lshr_b32 ttmp13, ttmp6, \src1
.endif
.endm
.macro S_STORE_DWORD_PCS_TTMP_REG1 base, offset
.if (.amdgcn.gfx_generation_minor >= 4)
s_store_dword ttmp6, \base, \offset
.else
s_store_dword ttmp13, \base, \offset
.endif
.endm
.macro S_MULK_I32_PCS_TTMP_REG1 const_val
.if (.amdgcn.gfx_generation_minor >= 4)
s_mulk_i32 ttmp6, \const_val
.else
s_mulk_i32 ttmp13, \const_val
.endif
.endm
.macro S_ADD_U32_PCS_TTMP_REG1 dst, src0
.if (.amdgcn.gfx_generation_minor >= 4)
s_add_u32 \dst, \src0, ttmp6
.else
s_add_u32 \dst, \src0, ttmp13
.endif
.endm
.macro S_CMP_LG_U32_PCS_TTMP_REG1 src0
.if (.amdgcn.gfx_generation_minor >= 4)
s_cmp_lg_u32 \src0, ttmp6
.else
s_cmp_lg_u32 \src0, ttmp13
.endif
.endm
.endif
// ABI between first and second level trap handler:
// ttmp0 = PC[31:0]
// ttmp8 = WorkgroupIdX
@@ -131,8 +227,8 @@ trap_entry:
s_bfe_u32 ttmp2, ttmp1, SQ_WAVE_PC_HI_TRAP_ID_BFE
s_cbranch_scc0 .no_skip_debugtrap
.if (.amdgcn.gfx_generation_number == 9 && .amdgcn.gfx_generation_minor < 4) // PC_SAMPLING_GFX9
// ttmp[14:15] is TMA2; Available: ttmp[2:3], ttmp[4:5], ttmp7, ttmp13
.if (.amdgcn.gfx_generation_number == 9) // PC_SAMPLING_GFX9
// ttmp[14:15] is TMA2; Available: ttmp[2:3], ttmp[4:5], ttmp7, TTMP_REG1
// Check if this is a host-trap. For now, if so, that means we are sampling
//
// TMA2 layout:
@@ -175,7 +271,8 @@ trap_entry:
// Ignore llvm.debugtrap.
s_branch .exit_trap
.if (.amdgcn.gfx_generation_number == 9 && .amdgcn.gfx_generation_minor < 4) // PC_SAMPLING_GFX9
.if (.amdgcn.gfx_generation_number == 9) // PC_SAMPLING_GFX9
// tma->host_trap_buffers Offsets:
// [0x00] uint64_t buf_write_val;
// [0x08] uint32_t buf_size;
@@ -219,25 +316,31 @@ trap_entry:
.profile_trap_handlers_gfx9:
s_mov_b64 ttmp[2:3], 1 // atomic increment buf_write_val
s_atomic_add_x2 ttmp[2:3], ttmp[14:15], glc // ttmp[2:3] = packed local_entry
s_load_dword ttmp13, ttmp[14:15], 0x8 // ttmp13 = tma->buf_size
S_LOAD_DWORD_PCS_TTMP_REG1 ttmp[14:15], 0x8 // TTMP_REG1 = tma->buf_size
s_waitcnt lgkmcnt(0)
s_lshr_b32 ttmp7, ttmp3, 31 // ttmp7 = buf_to_use
s_bitset0_b32 ttmp6, 31 // clear out ttmp6 bit31
S_BITSET0_B32_PCS_TTMP_REG2 31 // clear out TTMP_REG2 bit31
s_cmp_eq_u32 ttmp7, 0 // store off buf_to_use ...
s_cbranch_scc1 .skip_ttmp6_set_gfx9 // into bit31 of ttmp6
s_bitset1_b32 ttmp6, 31
.skip_ttmp6_set_gfx9:
s_cbranch_scc1 .skip_ttmp_set_gfx9 // into bit31 of TTMP_REG2
S_BITSET1_B32_PCS_TTMP_REG2 31
.skip_ttmp_set_gfx9:
s_bfe_u64 ttmp[2:3], ttmp[2:3], (63<<16) // ttmp[2:3] = new local_entry
s_cmp_lg_u32 ttmp3, 0 // if entry >= 2^32, always lost
s_cbranch_scc1 .pc_sampling_exit
s_cmp_ge_u32 ttmp2, ttmp13 // if local_entry >= buf_size
S_CMP_GE_U32_PCS_TTMP_REG1 ttmp2 // if local_entry >= buf_size
s_cbranch_scc1 .pc_sampling_exit
// ttmp2=local_entry, ttmp7=buf_to_use (also in bit31 of ttmp6), ttmp13=buf_size
// ttmp2=local_entry, ttmp7=buf_to_use (also in bit31 of TTMP_REG2), TTMP_REG1=buf_size
// ttmp[14:15] is tma->host_trap_buffers. Available: ttmp3, ttmp[4:5]
.if (.amdgcn.gfx_generation_number == 9 && .amdgcn.gfx_generation_minor == 4)
s_mul_i32 ttmp6, ttmp6, ttmp7 // ttmp[4:5]=buf_size if ...
s_mul_i32 ttmp4, ttmp6, 0x40 // buf_to_use=1, 0 otherwise
s_mul_hi_u32 ttmp5, ttmp6, 0x40
.else
s_mul_i32 ttmp13, ttmp13, ttmp7 // ttmp[4:5]=buf_size if ...
s_mul_i32 ttmp4, ttmp13, 0x40 // buf_to_use=1, 0 otherwise
s_mul_hi_u32 ttmp5, ttmp13, 0x40
.endif
s_add_u32 ttmp4, ttmp4, 0x40 // now ttmp[4:5]=offset from ...
s_addc_u32 ttmp5, ttmp5, 0 // tma to start of target buffer;
@@ -287,16 +390,24 @@ trap_entry:
s_and_b32 ttmp1, ttmp1, 0xffff // clear out extra data from PC_HI
s_store_dwordx2 ttmp[0:1], ttmp[2:3] // store PC
s_waitcnt lgkmcnt(0) // wait for timestamp
s_mov_b32 ttmp13, exec_lo
s_store_dword ttmp13, ttmp[2:3], 0x8 // store EXEC_LO
s_mov_b32 ttmp13, exec_hi
s_store_dword ttmp13, ttmp[2:3], 0xc // store EXEC_HI
S_MOV_B32_SRC_PCS_TTMP_REG1 exec_lo
S_STORE_DWORD_PCS_TTMP_REG1 ttmp[2:3], 0x8 // store EXEC_LO
S_MOV_B32_SRC_PCS_TTMP_REG1 exec_hi
S_STORE_DWORD_PCS_TTMP_REG1 ttmp[2:3], 0xc // store EXEC_HI
s_store_dwordx2 ttmp[8:9], ttmp[2:3], 0x10 // store wg_id_x and wg_id_y
s_store_dword ttmp10, ttmp[2:3], 0x18 // store wg_id_z
s_store_dwordx2 ttmp[4:5], ttmp[2:3], 0x30 // store timestamp
.if (.amdgcn.gfx_generation_number == 9 && .amdgcn.gfx_generation_minor >= 4)
s_getreg_b32 ttmp4, hwreg(HW_REG_XCC_ID) //store XCC_ID
s_lshl_b32 ttmp4, ttmp4, 8
s_and_b32 ttmp5, ttmp11, 0x3f
s_or_b32 ttmp4, ttmp4, ttmp5
s_store_dword ttmp4, ttmp[2:3], 0x1c // store wave_in_wg
.else
s_and_b32 ttmp4, ttmp11, 0x3f
s_store_dword ttmp4, ttmp[2:3], 0x1c // store wave_in_wg
.endif
// Get HW_ID using S_GETREG_B32 with size=32 (F8 in upper bits), offset=0, and HW_ID = 4 (0x4)
s_getreg_b32 ttmp4, hwreg(HW_REG_HW_ID)
s_store_dword ttmp4, ttmp[2:3], 0x20 // store HW_ID
@@ -336,7 +447,12 @@ trap_entry:
s_mov_b32 exec_hi, ttmp5 // do not care about message[63:32]
s_and_b32 ttmp5, exec_lo, DOORBELL_ID_MASK // doorbell now in ttmp5
s_mov_b32 exec_lo, ttmp4 // exec mask restored
.if (.amdgcn.gfx_generation_number == 9 && .amdgcn.gfx_generation_minor >= 4)
s_bfe_u32 ttmp4, ttmp11, (6 | 25 << 16) // extract dispatch ID from ttmp11
.else
s_and_b32 ttmp4, ttmp6, 0x1ffffff // extract low 25 bits from ttmp6 (DispatchPktIndx[24:0])
.endif
s_store_dwordx2 ttmp[4:5], ttmp[2:3], 0x38 // ttmp[4:5] is correlation ID. Store correlation_id to sample
// get_correlation_id() -- end //
@@ -347,16 +463,15 @@ trap_entry:
// ttmp[2:3], ttmp[4:5], ttmp7, and ttmp13 are free
// ttmp[14:15] = tma->host_trap_buffers; ttmp6.b31 is buf_to_use, 0 or 1
s_lshr_b32 ttmp13, ttmp6, 31 // ttmp13 is buf_to_use
s_mulk_i32 ttmp13, 0x10
// written_val0 to written_val_X
s_add_u32 ttmp14, ttmp14, ttmp13 // now ttmp[14:15] points to ...
S_LSHR_B32_PCS_TTMP_REG1_REG2 31 // TTMP_REG1 is buf_to_use
S_MULK_I32_PCS_TTMP_REG1 0x10 // written_val0 to written_val_X
S_ADD_U32_PCS_TTMP_REG1 ttmp14, ttmp14 // now ttmp[14:15] points to ...
s_addc_u32 ttmp15, ttmp15, 0x0 // buf_written_valX-0x10
s_mov_b32 ttmp7, 1 // atomic increment buf_written_valX
s_atomic_add ttmp7, ttmp[14:15], 0x10 glc // ttmp7 will contain 'done'
s_load_dword ttmp13, ttmp[14:15], 0x14 // ttmp13 will hold watermark
S_LOAD_DWORD_PCS_TTMP_REG1 ttmp[14:15], 0x14 // TTMP_REG1 will hold watermark
s_waitcnt lgkmcnt(0)
s_cmp_lg_u32 ttmp7, ttmp13 // if 'done' not at watermark, exit
S_CMP_LG_U32_PCS_TTMP_REG1 ttmp7 // if 'done' not at watermark, exit
s_cbranch_scc1 .pc_sampling_exit
// ttmp[2:3], [4:5], ttmp7, and ttmp13 are free
@@ -390,12 +505,12 @@ trap_entry:
s_cbranch_scc1 .pc_sampling_exit // event_id zero means no interrupt
s_store_dword ttmp7, ttmp[4:5] glc // send event ID to the mailbox
s_waitcnt lgkmcnt(0)
s_mov_b32 ttmp13, m0 // save off m0
S_MOV_B32_SRC_PCS_TTMP_REG1 m0 // save off m0
s_mov_b32 m0, ttmp7 // put ID into message payload
s_nop 0x0 // Manually inserted wait states
s_sendmsg sendmsg(MSG_INTERRUPT) // send interrupt message
s_waitcnt lgkmcnt(0) // wait for message to be sent
s_mov_b32 m0, ttmp13 // restore m0
S_MOV_B32_DST_PCS_TTMP_REG1 m0 // restore m0
// send_signal(...) - end //
.pc_sampling_exit:
// We can receive regular exceptions while doing PC-Sampling so we need to make sure we