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

PC Sampling: Add gfx9 2nd trap handler for PC Sampling

Browse Source 
Code is valid for gfx9 GPUs excluding gfx94x.

1st level trap handler will use TTMP13[22] to indicate host trap and
TTMP13[21] to indicate stochastic trap.

For each PC sampling method (hosttrap and stochastic), we use a double
buffering mechanism to transfer data between GPU and host.
The GPU will dump data into one buffer while CPU may be reading data
from the other buffer. There are 2 separate signals, one for each
buffer.
When signal != 0, the buffer belongs to the GPU and the GPU can write
to it. Once the buffer has reached the high watermark, the GPU will
set the signal to 0 to wake up the host and so that the host can try
to switch the buffers and read the data.

Co-authored-by: David Yat Sin <David.YatSin@amd.com>
Change-Id: If3eb0913e52fb4788059a71e5feca334612f3d5d


[ROCm/ROCR-Runtime commit: 431a70471e]
This commit is contained in:
Joseph Greathouse
2023-08-02 02:52:42 +00:00
committed by David Yat Sin
parent c3f9368b8f
commit 5c61936483
2 changed files with 305 additions and 5 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_gpu_agent.h
+16
View File
@@ -703,6 +703,22 @@ class GpuAgent : public GpuAgentInt {
std::function<void(void*)> finegrain_deallocator_;
/* PC Sampling fields - begin */
/* 2nd level Trap handler code is based on the offsets within this structure */
typedef struct {
uint64_t buf_write_val;
uint32_t buf_size;
uint32_t reserved0;
uint32_t buf_written_val0;
uint32_t buf_watermark0;
hsa_signal_t done_sig0;
uint32_t buf_written_val1;
uint32_t buf_watermark1;
hsa_signal_t done_sig1;
uint8_t reserved1[16];
/* pc_sample_t buffer0[buf_size]; */
/* pc_sample_t buffer1[buf_size]; */
} pcs_hosttrap_sampling_data_t;
typedef struct {
os::Thread thread;
pcs::PcsRuntime::PcSamplingSession* session;
projects/rocr-runtime/runtime/hsa-runtime/core/runtime/trap_handler/trap_handler.s
+289 -5
View File
@@ -54,6 +54,7 @@
.set SQ_WAVE_TRAPSTS_XNACK_ERROR_SHIFT , 28
.set SQ_WAVE_TRAPSTS_MATH_EXCP , 0x7F
.set SQ_WAVE_MODE_EXCP_EN_SHIFT , 12
.set SQ_WAVE_MODE_EXCP_EN_SIZE , 8
.set TRAP_ID_ABORT , 2
.set TRAP_ID_DEBUGTRAP , 3
.set DOORBELL_ID_SIZE , 10
@@ -92,21 +93,34 @@
.if .amdgcn.gfx_generation_number == 9 && .amdgcn.gfx_generation_minor >= 4
.set TTMP11_TTMPS_SETUP_SHIFT , 31
// Bit to indicate that this is a hosttrap trap instead of stochastic trap
// Currently not used
.set TTMP13_PCS_IS_STOCHASTIC , 24
.endif
// ABI between first and second level trap handler:
// ttmp0 = PC[31:0]
// ttmp0 = PC[31:0]
// ttmp8 = WorkgroupIdX
// ttmp9 = WorkgroupIdY
// ttmp10 = WorkgroupIdZ
// ttmp12 = SQ_WAVE_STATUS
// ttmp14 = TMA[31:0]
// ttmp15 = TMA[63:32]
// gfx9:
// ttmp1 = 0[2:0], PCRewind[3:0], HostTrap[0], TrapId[7:0], PC[47:32]
// gfx906/gfx908/gfx90a:
// ttmp11 = SQ_WAVE_IB_STS[20:15], 0[1:0], DebugEnabled[0], 0[15:0], NoScratch[0], WaveIdInWG[5:0]
// all gfx9 (except gfx940, gfx941, gfx942):
// ttmp6 = 0[6:0], DispatchPktIndx[24:0]
// ttmp11 = SQ_WAVE_IB_STS[20:15], 0[1:0], DebugEnabled[0], 0[15:0], NoScratch[0], WaveInWg[5:0]
// Note: Once stochastic sampling is implemented, L2 Trap Handler will use Bit 23
// (TTMP11_PCS_IS_STOCHASTIC) to differentiate between stochastic and hosttrap
// gfx940/gfx941/gfx942:
// ttmp11 = 0[0], DispatchPktIndx[24:0], WaveIdInWg[5:0]
// ttmp13 = SQ_WAVE_IB_STS[20:15], 0[1:0], DebugEnabled[0], 0[22:0]
// gfx10:
// ttmp1 = 0[0], PCRewind[5:0], HostTrap[0], TrapId[7:0], PC[47:32]
// gfx10/gfx11:
// ttmp6 = 0[6:0], DispatchPktIndx[24:0]
// gfx1010:
// ttmp11 = SQ_WAVE_IB_STS[25], SQ_WAVE_IB_STS[21:15], DebugEnabled[0], 0[15:0], NoScratch[0], WaveIdInWG[5:0]
// gfx1030/gfx1100:
@@ -117,6 +131,31 @@ 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
// Check if this is a host-trap. For now, if so, that means we are sampling
//
// TMA2 layout:
// [0x00] out_buf_t* host_trap_buffers;
// [0x08] out_buf_t* stochastic_trap_buffers;
//
// --- Start profile trap handlers GFX9 --- //
// if (host_trap) {
// if (stochastic) // Not implemented yet
// ttmp11.bit23 = 1; // Not implemented yet
// profiling_trap_handler(tma->host_trap_buffers);
// }
s_bitcmp1_b32 ttmp1, SQ_WAVE_PC_HI_HT_SHIFT
s_cbranch_scc0 .not_host_trap_gfx9
s_load_dwordx2 ttmp[14:15], ttmp[14:15], 0 glc // ttmp[14:15]=&host_trap_buffers
// TODO: When implementing stochastic sampling, need to set TTMP11_PCS_IS_STOCHASTIC
// or TTMP13_PCS_IS_STOCHASTIC to differentiate between hosttrap and stochastic sampling
s_waitcnt lgkmcnt(0)
s_branch .profile_trap_handlers_gfx9 // Off to the profile handlers
.not_host_trap_gfx9:
.endif // PC_SAMPLING_GFX9
// If caused by s_trap then advance PC.
s_bitcmp1_b32 ttmp1, SQ_WAVE_PC_HI_HT_SHIFT
s_cbranch_scc1 .not_s_trap
@@ -136,7 +175,252 @@ trap_entry:
// Ignore llvm.debugtrap.
s_branch .exit_trap
.if (.amdgcn.gfx_generation_number == 9 && .amdgcn.gfx_generation_minor < 4) // PC_SAMPLING_GFX9
// tma->host_trap_buffers Offsets:
// [0x00] uint64_t buf_write_val;
// [0x08] uint32_t buf_size;
// [0x0c] uint32_t reserved0;
// [0x10] uint32_t buf_written_val0;
// [0x14] uint32_t buf_watermark0;
// [0x18] hsa_signal_t done_sig0;
// [0x20] uint32_t buf_written_val1;
// [0x24] uint32_t buf_watermark1;
// [0x28] hsa_signal_t done_sig1;
// [0x30] uint8_t reserved1[16];
// [0x40] sample_t buffer0[buf_size];
// [0x40+(buf_size*sizeof(sample_t))]sample_t buffer1[buf_size];
//
//__global__ void profiling_trap_handler(out_buf_t* tma) {
// uint64_t local_entry = atomicAdd(&tma->buf_write_val, 1);
// int buf_to_use = local_entry >> 63;
// local_entry &= (ULLONG_MAX >> 1);
//
// if (local_entry < tma->buf_size) {
// sample_t *buf_base = buf_to_use ? tma->buffer1 : tma->buffer0;
// fill_sample(&buf_base[local_entry]); // reads TTMP11 as well
//
// uint32_t * written = buf_to_use ? &(tma->buf_written_val1) :
// &(tma->buf_written_val0);
//
// uint64_t done = __atomic_fetch_add(&written, 1,
// memory_order_release, memory_scope_system);
//
// uint32_t watermark = buf_to_use ? tma->buf_watermark0 :
// tma->buf_watermark1;
// if (done == watermark) {
// hsa_signal_t done_sig = buf_to_use ? tma->done_sig1 :
// tma->done_sig0;
// send_signal(done_sig);
// }
// }
//}
// ttmp[14:15] is tma->host_trap_buffers; Available: ttmp[2:3], ttmp[4:5], ttmp7, ttmp13
.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_waitcnt lgkmcnt(0)
s_lshr_b32 ttmp7, ttmp3, 31 // ttmp7 = buf_to_use
s_bitset0_b32 ttmp6, 31 // clear out ttmp6 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_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_cbranch_scc1 .pc_sampling_exit
// ttmp2=local_entry, ttmp7=buf_to_use (also in bit31 of ttmp6), ttmp13=buf_size
// ttmp[14:15] is tma->host_trap_buffers. Available: ttmp3, ttmp[4:5]
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
s_add_u32 ttmp4, ttmp4, 0x40 // now ttmp[4:5]=offset from ...
s_addc_u32 ttmp5, ttmp5, 0 // tma to start of target buffer;
s_add_u32 ttmp4, ttmp14, ttmp4 // ttmp[4:5] now points to ...
s_addc_u32 ttmp5, ttmp15, ttmp5 // buffer0 or buffer1
s_mov_b32 ttmp7, ttmp2
// ttmp7 contains local_entry, ttmp[4:5] contains "&bufferX",
// ttmp[14:15] holds 'tma->host_trap_buffers' pointer
// ttmp[2:3] and ttmp13 are available for gathering perf sample info
// ttmp[14:15] is live out
// fill_sample(...) - begin //
// typedef struct {
// [0x00] uint64_t pc;
// [0x08] uint64_t exec_mask;
// [0x10] uint32_t workgroup_id_x;
// [0x14] uint32_t workgroup_id_y;
// [0x18] uint32_t workgroup_id_z;
// [0x1c] uint32_t wave_in_wg : 6;
// uint32_t chiplet : 3; // Currently not used
// uint32_t reserved : 23;
// [0x20] uint32_t hw_id;
// [0x24] uint32_t reserved0;
// [0x28] uint64_t reserved1;
// [0x30] uint64_t timestamp;
// [0x38] uint64_t correlation_id;
// } perf_sample_hosttrap_v1_t;
//
// __device__ void fill_sample_hosttrap_v1(perf_sample_hosttrap_v1_t* buf) {
// buf->pc = ((ttmp1 & 0xffff) << 32) | ttmp0;
// buf->exec_mask = EXEC;
// buf->workgroup_id_x = ttmp8;
// buf->workgroup_id_y = ttmp9;
// buf->workgroup_id_z = ttmp10;
// buf->chiplet_and_wave_id = ttmp11 & 0x3f;
// buf->hw_id = s_getreg_b32(HW_REG_HW_ID);
// buf->timestamp = s_memrealtime;
// buf->correlation_id = get_correlation_id();
// }
s_mul_i32 ttmp2, ttmp7, 0x40 // offset into buffer for 64B objects
s_mul_hi_u32 ttmp3, ttmp7, 0x40 // ttmp[2:3] will contain byte ...
s_add_u32 ttmp2, ttmp2, ttmp4
s_addc_u32 ttmp3, ttmp3, ttmp5 // ttmp[2:3]=&bufferX[local_entry]
s_memrealtime ttmp[4:5]
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_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
s_and_b32 ttmp4, ttmp11, 0x3f
s_store_dword ttmp4, ttmp[2:3], 0x1c // store wave_in_wg
// 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
// ttmp[2:3] = &buffer[local_entry]; ttmp[4:5], ttmp7, and ttmp13 are free
// ttmp[14:15] = tma->host_trap_buffers and is live out; ttmp6.b31 is buf_to_use, 0 or 1
// get_correlation_id() -- begin //
// Returns a value to use as a correlation ID.
// Returns a 64bit number made up of the 9-bit queue ID and the
// 25-bit dispatch_pkt concatenated together as:
// Upper 32 bits: {23 0s}{9b queue_id}
// Lower 32 bits: { 7 0s}{25b dispatch_pkt}
// __device__ uint64_t get_correlation_id() {
// uint64_t output;
// // Get bottom 10 bits of queue's doorbell, in doorbell region.
// // Doorbell is 8B (3b per); region is 8K (13b total) so 10 bits.
// output = s_sendmsg(MSG_GET_DOORBELL);
// output &= 0x3ff;
// output <<= 32;
// // TTMP6 contains this packet dispatch ID modulus the queue size
// output |= TTMP6;
// return output;
// }
// ttmp[2:3] = &buffer[local_entry]
// ttmp[4:5], ttmp7, and ttmp13 are free
// ttmp[14:15] = tma->host_trap_buffers and is live out
// ttmp6.b31 is buf_to_use, 0 or 1 and is live out
s_mov_b64 ttmp[4:5], exec // back up EXEC mask
s_mov_b32 exec_lo, 0x80000000 // prepare EXEC for doorbell spin
s_sendmsg sendmsg(MSG_GET_DOORBELL) // message 10, puts doorbell in EXEC
.wait_for_doorbell:
s_nop 0x7 // wait a bit for message to return
s_bitcmp0_b32 exec_lo, 0x1f // returned message will 0 bit 31
s_cbranch_scc0 .wait_for_doorbell // wait some more if no data yet
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
s_and_b32 ttmp4, ttmp6, 0x1ffffff // extract low 25 bits from ttmp6 (DispatchPktIndx[24:0])
// ttmp[4:5] is correlation ID
s_store_dwordx2 ttmp[4:5], ttmp[2:3], 0x38 // store correlation_id to sample
// get_correlation_id() -- end //
// complete stores before returning
s_dcache_wb
s_waitcnt lgkmcnt(0)
// fill_sample(...) - end //
// 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_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_waitcnt lgkmcnt(0)
s_cmp_lg_u32 ttmp7, ttmp13 // if 'done' not at watermark, exit
s_cbranch_scc1 .pc_sampling_exit
// ttmp[2:3], [4:5], ttmp7, and ttmp13 are free
// ttmp[14:15] = buf_written_valX-0x10
// send_signal(...) - begin //
//__device__ void send_signal(hsa_signal_t* signal) {
//
// amd_signal_t *sig = (amd_signal_t *)signal->handle;
// __atomic_store(&(sig->value), 0, memory_order_relaxed, memory_scope_system);
// if (sig->event_mailbox_ptr != NULL && sig->event_id != NULL) {
// uint32_t id = sig->event_id;
// __atomic_store(sig->event_mailbox_ptr, id,
// memory_order_relaxed, memory_scope_system);
// __builtin_amdgcn_s_sendmsg(1, id);
// }
//}
// We jump to the trap handler exit after this, so no live-out registers except
// those that must survive the trap handler
s_load_dwordx2 ttmp[2:3], ttmp[14:15], 0x18 // load done_sig into ttmp[2:3]
s_waitcnt lgkmcnt(0) // it's actually an amd_signal_t*
s_load_dwordx2 ttmp[4:5], ttmp[2:3], 0x10 // load event mailbox ptr into 4:5
s_load_dword ttmp7, ttmp[2:3], 0x18 // load event_id into ttmp7
s_mov_b64 ttmp[14:15], 0
s_store_dwordx2 ttmp[14:15], ttmp[2:3], 0x8 glc // zero out signal value
s_waitcnt lgkmcnt(0) // wait for value store to complete
s_cmp_eq_u64 ttmp[4:5], 0
s_cbranch_scc1 .pc_sampling_exit // null mailbox means no interrupt
s_cmp_eq_u32 ttmp7, 0
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 m0, ttmp7 // put ID into message payload
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
// send_signal(...) - end //
.pc_sampling_exit:
// We can receive regular exceptions while doing PC-Sampling so we need to make sure we
// handle these exceptions here
s_getreg_b32 ttmp2, hwreg(HW_REG_TRAPSTS)
s_getreg_b32 ttmp3, hwreg(HW_REG_MODE, SQ_WAVE_MODE_EXCP_EN_SHIFT, SQ_WAVE_MODE_EXCP_EN_SIZE) // ttmp3[7:0] = MODE.EXCP_EN
// Set bits corresponding to TRAPSTS.MEM_VIOL, TRAPSTS.ILLEGAL_INST and TRAPSTS.XNACK_ERROR
s_or_b32 ttmp3, ttmp3, (1 << SQ_WAVE_TRAPSTS_MEM_VIOL_SHIFT | 1 << SQ_WAVE_TRAPSTS_ILLEGAL_INST_SHIFT | 1 << SQ_WAVE_TRAPSTS_XNACK_ERROR_SHIFT)
s_getreg_b32 ttmp2, hwreg(HW_REG_TRAPSTS)
s_and_b32 ttmp2, ttmp2, ttmp3
// SCC will be 1 if either a maskable instruction was set, or one of MEM_VIOL, ILL_INST, XNACK_ERROR
s_cbranch_scc1 .no_skip_debugtrap // if any of those are set, handle exceptions
// Check for maskable exceptions
s_getreg_b32 ttmp3, hwreg(HW_REG_MODE, SQ_WAVE_MODE_EXCP_EN_SHIFT, SQ_WAVE_MODE_EXCP_EN_SIZE)
s_and_b32 ttmp3, ttmp2, ttmp3
s_cbranch_scc1 .no_skip_debugtrap
// Since we are in PC sampling, it is safe to ignore watch1/2/3 and single step
// as those should only be enabled by the debugger.
// We could add them for completeness, i.e. check MODE.DEBUG_EN (bit 11)
// and "MODE.EXCP_EN.WATCH (bit 19) && (TRAPSTS.EXCP_HI.ADDR_WATCH1 (bit 12) || TRAPSTS.EXCP_HI.ADDR_WATCH2 (bit 13) || TRAPSTS.EXCP_HI.ADDR_WATCH3 (bit 14)).
s_branch .exit_trap
.endif // PC_SAMPLING_GFX9
.no_skip_debugtrap:
// Save trap id and halt status in ttmp6.
s_andn2_b32 ttmp6, ttmp6, (TTMP6_SAVED_TRAP_ID_MASK | TTMP6_SAVED_STATUS_HALT_MASK)
@@ -279,8 +563,8 @@ trap_entry:
.endif
// Restore SQ_WAVE_STATUS.
s_and_b64 exec, exec, exec // Restore STATUS.EXECZ, not writable by s_setreg_b32
s_and_b64 vcc, vcc, vcc // Restore STATUS.VCCZ, not writable by s_setreg_b32
s_and_b64 exec, exec, exec // restore STATUS.EXECZ, not writable by s_setreg_b32
s_and_b64 vcc, vcc, vcc // restore STATUS.VCCZ, not writable by s_setreg_b32
s_setreg_b32 hwreg(HW_REG_STATUS), ttmp12
// Return to original (possibly modified) PC.