89 lines
3.0 KiB
Plaintext
89 lines
3.0 KiB
Plaintext
// gpu_log_ring.cu — unified GPU diagnostic log ring (device-side).
|
|
//
|
|
// Provides `log_record()` (a device __forceinline__ helper) and the
|
|
// `gpu_log_tick` kernel that increments the step counter once per
|
|
// captured-graph step.
|
|
//
|
|
// Constraints:
|
|
// - No atomicAdd (slot reservation by deterministic step/kid/rt math).
|
|
// - No host branches; thread-id gating only.
|
|
// - Header-last commit (magic field written after payload + other header).
|
|
// - Mapped-pinned ring; host reads via volatile, no DtoH memcpy.
|
|
//
|
|
// Per `feedback_nvidia_grade_perf_for_kernels`: single-thread writer for
|
|
// the slot's 64 B cacheline, one coalesced store.
|
|
|
|
#include "gpu_log_ids.h"
|
|
|
|
#include <cstdint>
|
|
|
|
struct LogHeader {
|
|
uint32_t magic;
|
|
uint32_t step;
|
|
uint8_t kernel_id;
|
|
uint8_t record_type;
|
|
uint8_t payload_words;
|
|
uint8_t reserved;
|
|
uint32_t _padding;
|
|
};
|
|
|
|
struct LogRecord {
|
|
LogHeader header;
|
|
float payload[GPU_LOG_PAYLOAD_F32];
|
|
};
|
|
|
|
struct LogRing {
|
|
LogRecord records[GPU_LOG_N_SLOTS];
|
|
};
|
|
|
|
// Tick kernel — 1 thread, 1 block. Increments step counter; runs FIRST in
|
|
// every captured-graph replay so all subsequent kernels see the new step.
|
|
extern "C" __global__ void gpu_log_tick(int* step_counter) {
|
|
if (threadIdx.x == 0 && blockIdx.x == 0) {
|
|
step_counter[0] = step_counter[0] + 1;
|
|
}
|
|
}
|
|
|
|
// log_record — emit a diagnostic record from a kernel.
|
|
//
|
|
// Single-writer per (step, kernel_id, record_type) slot. Thread (0,0) of
|
|
// block 0 writes; all other threads no-op. Header-last commit ensures
|
|
// torn-record reads on the host side are detectable via magic mismatch.
|
|
//
|
|
// Caller responsibility:
|
|
// - g_log_ring == nullptr → no-op (logging disabled).
|
|
// - g_step_counter is the device step counter (incremented by gpu_log_tick).
|
|
// - payload_words ≤ GPU_LOG_PAYLOAD_F32 (12); kernel must respect this.
|
|
__device__ __forceinline__ void log_record(
|
|
LogRing* g_log_ring,
|
|
const int* g_step_counter,
|
|
uint8_t kernel_id,
|
|
uint8_t record_type,
|
|
const float* payload,
|
|
int payload_words
|
|
) {
|
|
if (g_log_ring == nullptr) return;
|
|
if (blockIdx.x != 0 || threadIdx.x != 0) return;
|
|
|
|
const int step = g_step_counter[0];
|
|
const int slot = (step * GPU_LOG_N_RECORDS_PER_STEP
|
|
+ (int)kernel_id * GPU_LOG_N_RT_PER_KERNEL
|
|
+ (int)record_type) & (GPU_LOG_N_SLOTS - 1);
|
|
LogRecord* rec = &g_log_ring->records[slot];
|
|
|
|
const int n = payload_words < GPU_LOG_PAYLOAD_F32 ? payload_words : GPU_LOG_PAYLOAD_F32;
|
|
#pragma unroll
|
|
for (int i = 0; i < GPU_LOG_PAYLOAD_F32; ++i) {
|
|
rec->payload[i] = (i < n) ? payload[i] : 0.0f;
|
|
}
|
|
rec->header.step = (uint32_t)step;
|
|
rec->header.kernel_id = kernel_id;
|
|
rec->header.record_type = record_type;
|
|
rec->header.payload_words = (uint8_t)n;
|
|
rec->header.reserved = 0;
|
|
rec->header._padding = 0;
|
|
|
|
__threadfence(); // make payload + header fields visible …
|
|
rec->header.magic = GPU_LOG_MAGIC; // … BEFORE the commit sentinel.
|
|
}
|