diff --git a/csrc/cuda/distributed/deterministic_collective.cu b/csrc/cuda/distributed/deterministic_collective.cu index 7a3cdf0e..54324ae2 100644 --- a/csrc/cuda/distributed/deterministic_collective.cu +++ b/csrc/cuda/distributed/deterministic_collective.cu @@ -30,7 +30,7 @@ constexpr int kThreads = 256; constexpr int kMaxBlocks = 4096; constexpr int kStagingFrames = 3; constexpr int kFusedStagingSlots = 2; -constexpr int64_t kSequenceHeaderBytes = 3 * sizeof(uint64_t); +constexpr int64_t kSequenceHeaderBytes = 4 * sizeof(uint64_t); // Small tensor collectives are launch-bound. Keep the existing DMA path for // larger transfers, where replacing cudaMemcpyAsync with a one-block copy // would reduce bandwidth and hurt overlap. @@ -216,6 +216,13 @@ __device__ __forceinline__ T fixed_tree_reduce( const PeerPointers& peers, int64_t index); +#if (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__)) +template +__device__ __forceinline__ nv_bfloat162 fixed_tree_reduce_bf16x2( + const PeerPointers& peers, + int64_t pair_index); +#endif + template __global__ void deterministic_all_reduce_fast_kernel( PeerPointers peers, @@ -228,8 +235,153 @@ __global__ void deterministic_all_reduce_fast_kernel( } __syncthreads(); - for (int64_t index = threadIdx.x; index < element_count; index += blockDim.x) { - output[index] = fixed_tree_reduce(peers, index); + if constexpr (std::is_same_v) { +#if (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__)) + const int64_t pair_count = element_count / 2; + auto* pair_output = reinterpret_cast(output); + for (int64_t pair_index = threadIdx.x; + pair_index < pair_count; + pair_index += blockDim.x) { + pair_output[pair_index] = + fixed_tree_reduce_bf16x2(peers, pair_index); + } + if ((element_count & 1) != 0 && threadIdx.x == 0) { + output[element_count - 1] = + fixed_tree_reduce(peers, element_count - 1); + } +#endif + } else { + for (int64_t index = threadIdx.x; index < element_count; index += blockDim.x) { + output[index] = fixed_tree_reduce(peers, index); + } + } + __syncthreads(); + + if (threadIdx.x == 0) { + __threadfence_system(); + store_release_system( + local_done_sequence, + load_acquire_system(local_stage_sequence)); + } +} + +// Preserve the graph-safe single-slot protocol while removing the launch +// boundary between staging and reduction. The sequence and fixed-tree order +// are identical to stage_payload_fast_kernel followed by +// deterministic_all_reduce_fast_kernel. +template +__global__ void deterministic_all_reduce_graph_safe_fused_fast_kernel( + PeerPointers peers, + uint64_t* local_stage_sequence, + uint64_t* local_done_sequence, + const uint8_t* input, + uint8_t* payload, + T* output, + int64_t element_count, + int64_t input_bytes) { + if (threadIdx.x == 0) { + const uint64_t sequence = load_acquire_system(local_stage_sequence); + for (int peer = 0; peer < WorldSize; ++peer) { + while (load_acquire_system(peers.done_sequences[peer]) < sequence) { + __nanosleep(64); + } + } + } + __syncthreads(); + + if (((reinterpret_cast(input) | + reinterpret_cast(payload) | + static_cast(input_bytes)) & 15u) == 0u) { + const auto* source = reinterpret_cast(input); + auto* destination = reinterpret_cast(payload); + const int64_t vector_count = input_bytes / sizeof(uint4); + for (int64_t index = threadIdx.x; index < vector_count; index += blockDim.x) { + destination[index] = source[index]; + } + } else { + for (int64_t index = threadIdx.x; index < input_bytes; index += blockDim.x) { + payload[index] = input[index]; + } + } + __syncthreads(); + + if (threadIdx.x == 0) { + __threadfence_system(); + store_release_system( + local_stage_sequence, + load_acquire_system(local_stage_sequence) + 1); + wait_for_stage_sequence(peers, WorldSize, local_stage_sequence); + } + __syncthreads(); + + if constexpr (std::is_same_v) { +#if (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__)) + const int64_t pair_count = element_count / 2; + auto* pair_output = reinterpret_cast(output); + for (int64_t pair_index = threadIdx.x; + pair_index < pair_count; + pair_index += blockDim.x) { + pair_output[pair_index] = + fixed_tree_reduce_bf16x2(peers, pair_index); + } + if ((element_count & 1) != 0 && threadIdx.x == 0) { + output[element_count - 1] = + fixed_tree_reduce(peers, element_count - 1); + } +#endif + } else { + for (int64_t index = threadIdx.x; index < element_count; index += blockDim.x) { + output[index] = fixed_tree_reduce(peers, index); + } + } + __syncthreads(); + + if (threadIdx.x == 0) { + __threadfence_system(); + store_release_system( + local_done_sequence, + load_acquire_system(local_stage_sequence)); + } +} + +// The caller has already written the GEMM result into this rank's IPC +// payload. Publish those bytes and run the same fixed tree without copying +// them through an intermediate tensor. +template +__global__ void deterministic_all_reduce_staged_fast_kernel( + PeerPointers peers, + uint64_t* local_stage_sequence, + uint64_t* local_done_sequence, + T* output, + int64_t element_count) { + if (threadIdx.x == 0) { + __threadfence_system(); + store_release_system( + local_stage_sequence, + load_acquire_system(local_stage_sequence) + 1); + wait_for_stage_sequence(peers, WorldSize, local_stage_sequence); + } + __syncthreads(); + + if constexpr (std::is_same_v) { +#if (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__)) + const int64_t pair_count = element_count / 2; + auto* pair_output = reinterpret_cast(output); + for (int64_t pair_index = threadIdx.x; + pair_index < pair_count; + pair_index += blockDim.x) { + pair_output[pair_index] = + fixed_tree_reduce_bf16x2(peers, pair_index); + } + if ((element_count & 1) != 0 && threadIdx.x == 0) { + output[element_count - 1] = + fixed_tree_reduce(peers, element_count - 1); + } +#endif + } else { + for (int64_t index = threadIdx.x; index < element_count; index += blockDim.x) { + output[index] = fixed_tree_reduce(peers, index); + } } __syncthreads(); @@ -577,6 +729,87 @@ void launch_all_reduce_fast( } } +template +void launch_all_reduce_graph_safe_fused_fast( + const PeerPointers& peers, + uint64_t* local_stage_sequence, + uint64_t* local_done_sequence, + const uint8_t* input, + uint8_t* payload, + T* output, + int64_t element_count, + int64_t input_bytes, + int64_t world_size, + cudaStream_t stream) { + switch (world_size) { + case 1: + deterministic_all_reduce_graph_safe_fused_fast_kernel + <<<1, kThreads, 0, stream>>>( + peers, local_stage_sequence, local_done_sequence, + input, payload, output, element_count, input_bytes); + break; + case 2: + deterministic_all_reduce_graph_safe_fused_fast_kernel + <<<1, kThreads, 0, stream>>>( + peers, local_stage_sequence, local_done_sequence, + input, payload, output, element_count, input_bytes); + break; + case 4: + deterministic_all_reduce_graph_safe_fused_fast_kernel + <<<1, kThreads, 0, stream>>>( + peers, local_stage_sequence, local_done_sequence, + input, payload, output, element_count, input_bytes); + break; + case 8: + deterministic_all_reduce_graph_safe_fused_fast_kernel + <<<1, kThreads, 0, stream>>>( + peers, local_stage_sequence, local_done_sequence, + input, payload, output, element_count, input_bytes); + break; + default: + TORCH_CHECK(false, "unsupported deterministic collective world size ", world_size); + } +} + +template +void launch_all_reduce_staged_fast( + const PeerPointers& peers, + uint64_t* local_stage_sequence, + uint64_t* local_done_sequence, + T* output, + int64_t element_count, + int64_t world_size, + cudaStream_t stream) { + switch (world_size) { + case 1: + deterministic_all_reduce_staged_fast_kernel + <<<1, kThreads, 0, stream>>>( + peers, local_stage_sequence, local_done_sequence, + output, element_count); + break; + case 2: + deterministic_all_reduce_staged_fast_kernel + <<<1, kThreads, 0, stream>>>( + peers, local_stage_sequence, local_done_sequence, + output, element_count); + break; + case 4: + deterministic_all_reduce_staged_fast_kernel + <<<1, kThreads, 0, stream>>>( + peers, local_stage_sequence, local_done_sequence, + output, element_count); + break; + case 8: + deterministic_all_reduce_staged_fast_kernel + <<<1, kThreads, 0, stream>>>( + peers, local_stage_sequence, local_done_sequence, + output, element_count); + break; + default: + TORCH_CHECK(false, "unsupported deterministic collective world size ", world_size); + } +} + template void launch_all_reduce_fused_fast( const FusedPeerPointers& peer_slots, @@ -942,6 +1175,79 @@ class DeterministicCollectiveState { } } + void prepare_staged(torch::Tensor& input, cudaStream_t stream) { + check_tensor(input, "direct staging input"); + const int64_t input_bytes = input.numel() * input.element_size(); + TORCH_CHECK( + input_bytes <= kSingleBlockFastPathMaxBytes, + "direct staging supports at most ", + kSingleBlockFastPathMaxBytes, + " bytes, got ", + input_bytes); + TORCH_CHECK( + input.data_ptr() == peers_.values[rank_], + "direct staging tensor must start at this rank's IPC payload"); + wait_for_previous_done_kernel<<<1, 1, 0, stream>>>( + peers_, world_size_, local_stage_sequence_); + AT_CUDA_CHECK(cudaGetLastError()); + } + + void all_reduce_staged( + torch::Tensor& input, + torch::Tensor& output, + cudaStream_t stream) { + check_tensor(input, "direct staging input"); + check_tensor(output, "output"); + const int64_t input_bytes = input.numel() * input.element_size(); + TORCH_CHECK( + input_bytes <= kSingleBlockFastPathMaxBytes, + "direct staged all-reduce supports at most ", + kSingleBlockFastPathMaxBytes, + " bytes, got ", + input_bytes); + TORCH_CHECK( + input.data_ptr() == peers_.values[rank_], + "direct staging tensor must start at this rank's IPC payload"); + TORCH_CHECK( + output.scalar_type() == input.scalar_type(), + "direct staged all-reduce output dtype must match the input dtype"); + TORCH_CHECK( + output.numel() == input.numel(), + "direct staged all-reduce output size must match the input size"); + TORCH_CHECK( + output.data_ptr() != input.data_ptr(), + "direct staged all-reduce output must not alias the IPC payload"); + + switch (input.scalar_type()) { + case at::ScalarType::Float: + launch_all_reduce_staged_fast( + peers_, local_stage_sequence_, local_done_sequence_, + static_cast(output.data_ptr()), output.numel(), + world_size_, stream); + break; + case at::ScalarType::Half: + launch_all_reduce_staged_fast( + peers_, local_stage_sequence_, local_done_sequence_, + static_cast(output.data_ptr()), output.numel(), + world_size_, stream); + break; +#if (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__)) + case at::ScalarType::BFloat16: + launch_all_reduce_staged_fast( + peers_, local_stage_sequence_, local_done_sequence_, + static_cast(output.data_ptr()), output.numel(), + world_size_, stream); + break; +#endif + default: + TORCH_CHECK( + false, + "deterministic all-reduce supports float32, float16, and bfloat16; got ", + input.scalar_type()); + } + AT_CUDA_CHECK(cudaGetLastError()); + } + void stage(torch::Tensor& input, cudaStream_t stream) { check_tensor(input, "input"); const int64_t input_bytes = input.numel() * input.element_size(); @@ -1182,9 +1488,46 @@ class DeterministicCollectiveState { output.numel() == input.numel(), "all-reduce output size must match the input size"); - // Use the graph-safe staged protocol for every message size. The fused - // two-slot protocol is intentionally kept available in the extension for - // experiments, but is not safe to replay across vLLM's many graph shapes. + // Fuse the hot small-message path without changing the graph-safe + // single-slot sequence protocol or the fixed-tree reduction order. The + // experimental two-slot protocol remains disabled across graph shapes. + if (input_bytes <= kSingleBlockFastPathMaxBytes) { + auto* payload = const_cast( + static_cast(peers_.values[rank_])); + switch (input.scalar_type()) { + case at::ScalarType::Float: + launch_all_reduce_graph_safe_fused_fast( + peers_, local_stage_sequence_, local_done_sequence_, + static_cast(input.data_ptr()), payload, + static_cast(output.data_ptr()), output.numel(), + input_bytes, world_size_, stream); + break; + case at::ScalarType::Half: + launch_all_reduce_graph_safe_fused_fast( + peers_, local_stage_sequence_, local_done_sequence_, + static_cast(input.data_ptr()), payload, + static_cast(output.data_ptr()), output.numel(), + input_bytes, world_size_, stream); + break; +#if (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__)) + case at::ScalarType::BFloat16: + launch_all_reduce_graph_safe_fused_fast( + peers_, local_stage_sequence_, local_done_sequence_, + static_cast(input.data_ptr()), payload, + static_cast(output.data_ptr()), output.numel(), + input_bytes, world_size_, stream); + break; +#endif + default: + TORCH_CHECK( + false, + "deterministic all-reduce supports float32, float16, and bfloat16; got ", + input.scalar_type()); + } + AT_CUDA_CHECK(cudaGetLastError()); + return; + } + stage(input, stream); all_reduce(output, stream, /*allow_owner_path=*/false); } @@ -1225,6 +1568,79 @@ class DeterministicCollectiveState { AT_CUDA_CHECK(cudaGetLastError()); } + void all_gather_many( + const std::vector& inputs, + const std::vector& outputs, + cudaStream_t stream) { + TORCH_CHECK(!inputs.empty(), "all_gather_many requires at least one input"); + TORCH_CHECK( + inputs.size() == outputs.size(), + "all_gather_many inputs and outputs must have the same length"); + TORCH_CHECK( + !has_staged_input_, + "cannot run all_gather_many with a pending stage()"); + + std::vector offsets(inputs.size()); + std::vector input_bytes(inputs.size()); + int64_t total_bytes = 0; + for (size_t index = 0; index < inputs.size(); ++index) { + const auto& input = inputs[index]; + const auto& output = outputs[index]; + check_tensor(input, "input"); + check_tensor(output, "output"); + TORCH_CHECK( + output.scalar_type() == input.scalar_type(), + "all_gather_many output dtype must match its input dtype"); + const int64_t bytes = input.numel() * input.element_size(); + TORCH_CHECK( + output.numel() * output.element_size() == bytes * world_size_, + "all_gather_many output must contain one input per rank"); + total_bytes = (total_bytes + 15) & ~int64_t{15}; + offsets[index] = total_bytes; + input_bytes[index] = bytes; + total_bytes += bytes; + } + TORCH_CHECK( + total_bytes <= capacity_bytes_, + "all_gather_many inputs require ", total_bytes, + " bytes but staging capacity is ", capacity_bytes_); + + wait_for_previous_done_kernel<<<1, 1, 0, stream>>>( + peers_, world_size_, local_stage_sequence_); + AT_CUDA_CHECK(cudaGetLastError()); + auto* local_payload = const_cast( + static_cast(peers_.values[rank_])); + for (size_t index = 0; index < inputs.size(); ++index) { + if (input_bytes[index] == 0) continue; + AT_CUDA_CHECK(cudaMemcpyAsync( + local_payload + offsets[index], + inputs[index].data_ptr(), + input_bytes[index], + cudaMemcpyDeviceToDevice, + stream)); + } + publish_next_stage_sequence_kernel<<<1, 1, 0, stream>>>( + local_stage_sequence_); + AT_CUDA_CHECK(cudaGetLastError()); + wait_for_staged_peers(stream); + + for (size_t index = 0; index < inputs.size(); ++index) { + auto* output = static_cast(outputs[index].data_ptr()); + for (int peer = 0; peer < world_size_; ++peer) { + if (input_bytes[index] == 0) continue; + const auto* peer_payload = + static_cast(peers_.values[peer]); + AT_CUDA_CHECK(cudaMemcpyAsync( + output + static_cast(peer) * input_bytes[index], + peer_payload + offsets[index], + input_bytes[index], + cudaMemcpyDeviceToDevice, + stream)); + } + } + publish_done(stream); + } + void reduce_scatter(torch::Tensor& output, cudaStream_t stream) { check_tensor(output, "output"); TORCH_CHECK(has_staged_input_, "stage() must be called before reduce_scatter()"); @@ -1547,6 +1963,23 @@ void deterministic_collective_all_reduce(int64_t handle, torch::Tensor& output) state_from_handle(handle)->all_reduce(output, stream); } +void deterministic_collective_prepare_staged( + int64_t handle, + torch::Tensor& input) { + const c10::cuda::CUDAGuard device_guard(input.device()); + auto stream = c10::cuda::getCurrentCUDAStream().stream(); + state_from_handle(handle)->prepare_staged(input, stream); +} + +void deterministic_collective_all_reduce_staged( + int64_t handle, + torch::Tensor& input, + torch::Tensor& output) { + const c10::cuda::CUDAGuard device_guard(input.device()); + auto stream = c10::cuda::getCurrentCUDAStream().stream(); + state_from_handle(handle)->all_reduce_staged(input, output, stream); +} + void deterministic_collective_all_reduce_fused( int64_t handle, torch::Tensor& input, @@ -1565,6 +1998,16 @@ void deterministic_collective_all_gather_fused( state_from_handle(handle)->all_gather_fused(input, output, stream); } +void deterministic_collective_all_gather_many( + int64_t handle, + std::vector inputs, + std::vector outputs) { + TORCH_CHECK(!inputs.empty(), "all_gather_many requires at least one input"); + const c10::cuda::CUDAGuard device_guard(inputs.front().device()); + auto stream = c10::cuda::getCurrentCUDAStream().stream(); + state_from_handle(handle)->all_gather_many(inputs, outputs, stream); +} + void deterministic_collective_reduce_scatter(int64_t handle, torch::Tensor& output) { const c10::cuda::CUDAGuard device_guard(output.device()); auto stream = c10::cuda::getCurrentCUDAStream().stream(); diff --git a/csrc/cuda/gemm/det_gemm_kernel.cu b/csrc/cuda/gemm/det_gemm_kernel.cu index 1b535d2a..bd009bd1 100644 --- a/csrc/cuda/gemm/det_gemm_kernel.cu +++ b/csrc/cuda/gemm/det_gemm_kernel.cu @@ -213,7 +213,10 @@ constexpr int M_TILES = WARP_M / MMA_M; // 2 constexpr int N_TILES = BN / MMA_N; // 8 constexpr int K_TILES = BK / MMA_K; // 2 constexpr int KK_GROUPS = BK / 32; // 1 -constexpr int TREE_DEPTH = 16; +// Online mid-tree reduction needs ceil(log2(K / BK)) live levels. The +// training contract caps a rank's GEMM K at 32768, so ten levels cover every +// configured shape while avoiding six never-addressed per-thread stack slots. +constexpr int TREE_DEPTH = 10; __device__ __forceinline__ int mid_tree_merge_count(int leaf, int n) { int lo = 0, hi = n, count = 0; @@ -402,6 +405,7 @@ template bool launch_sm90(const nv_bf16* A, const nv_bf16* Bt, output_t* C, int M, int N, int K, cudaStream_t stream) { if (M % BM != 0 || N % BN != 0 || K % BK != 0) return false; // fall back + if (K / BK > (1 << TREE_DEPTH)) return false; CUtensorMap a_tmap, bt_tmap; det_gemm::init_tmap_noswizzle(&a_tmap, A, M, K, BM, BK); diff --git a/csrc/ops.cpp b/csrc/ops.cpp index 0a43baf1..603adceb 100644 --- a/csrc/ops.cpp +++ b/csrc/ops.cpp @@ -138,12 +138,19 @@ int64_t deterministic_collective_create( void deterministic_collective_destroy(int64_t handle); void deterministic_collective_stage(int64_t handle, torch::Tensor& input); void deterministic_collective_all_reduce(int64_t handle, torch::Tensor& output); +void deterministic_collective_prepare_staged(int64_t handle, torch::Tensor& input); +void deterministic_collective_all_reduce_staged( + int64_t handle, torch::Tensor& input, torch::Tensor& output); void deterministic_collective_all_reduce_fused( int64_t handle, torch::Tensor& input, torch::Tensor& output); void deterministic_collective_reduce_scatter(int64_t handle, torch::Tensor& output); void deterministic_collective_all_gather(int64_t handle, torch::Tensor& output); void deterministic_collective_all_gather_fused( int64_t handle, torch::Tensor& input, torch::Tensor& output); +void deterministic_collective_all_gather_many( + int64_t handle, + std::vector inputs, + std::vector outputs); #endif #if defined(KERNEL_ALIGN_WITH_ROCM) @@ -549,6 +556,14 @@ PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) { "deterministic_collective_all_reduce", &deterministic_collective_all_reduce, "Run the TP=8 deterministic fixed-tree all-reduce kernel"); + m.def( + "deterministic_collective_prepare_staged", + &deterministic_collective_prepare_staged, + "Reserve the local IPC payload for direct GEMM output"); + m.def( + "deterministic_collective_all_reduce_staged", + &deterministic_collective_all_reduce_staged, + "Reduce a GEMM result already resident in the local IPC payload"); m.def( "deterministic_collective_all_reduce_fused", &deterministic_collective_all_reduce_fused, @@ -565,6 +580,10 @@ PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) { "deterministic_collective_all_gather_fused", &deterministic_collective_all_gather_fused, "Run a fused small-message deterministic rank-ordered all-gather"); + m.def( + "deterministic_collective_all_gather_many", + &deterministic_collective_all_gather_many, + "Gather multiple tensors with one deterministic staging handshake"); #endif #if defined(KERNEL_ALIGN_WITH_ROCM) diff --git a/examples/__init__.py b/examples/__init__.py new file mode 100644 index 00000000..e69de29b diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/README.md b/examples/vime_qwen3_8b_tp4_cp2_200/README.md new file mode 100644 index 00000000..e51aceb6 --- /dev/null +++ b/examples/vime_qwen3_8b_tp4_cp2_200/README.md @@ -0,0 +1,194 @@ +# VIME Qwen3-8B TP4/CP2 200-round consistency experiment + +This example measures train/rollout numerical consistency at two independent +layers: VIME's framework-level reuse of rollout log-probabilities and +RL-Kernel's operator-level alignment of Attention, dense FFN, and linear logp. +It is designed for one 8×H100 node. Megatron uses TP4/CP2 across all eight +GPUs; two TP4 vLLM engines share those GPUs through VIME colocated offload. + +The optimization algorithm is explicitly fixed to GRPO with +`--advantage-estimator grpo`. DAPO-Math-17k is only the prompt/answer dataset; +it does not select the DAPO training algorithm. The rule reward is computed by +VIME's `deepscaler` reward implementation. + +The experiment is fail-closed. A run is accepted only when its Ray job +succeeds, every expected operator route has runtime execution evidence, no +fallback or Triton route is observed for an R/R arm, the requested number of +steps is present, and vLLM CUDA Graph evidence matches the manifest. + +## Ablation matrix + +| Group | VIME `--use-rollout-logprobs` | Attention / FFN / logp | Purpose | +|---|---:|---|---| +| G00 | off | P/P | Production baseline | +| G10 | on | P/P | Framework-level consistency only | +| G01 | off | R/R | RL-Kernel operator-level consistency only | +| G11 | on | R/R | Framework-level plus operator-level consistency | + +`P/P` selects the production implementation on training and rollout. For +Megatron linear logp this means that no external provider is configured and +VIME calls its native `calculate_log_probs_and_entropy` implementation +directly. `R/R` selects RL-Kernel on both sides and installs the strict +RL-Kernel linear-logp provider. All four groups use the same prompts, initial +checkpoint, sampling settings, seeds, TP4/CP2 topology, and batch sizes. + +Do not interpret G10/G11 as evidence that train and rollout recomputation is +bitwise equal: framework reuse changes which stored logp enters the RL loss. +The direct numerical claim comes from G01/G11 and the runtime comparison +metrics. + +## Required gates + +- NVIDIA H100 × 8; colocated actor/rollout GPUs 8; actor TP=4, CP=2, PP=1; + two rollout engines with TP=4 each. +- Keep the TP4 Megatron actor resident and offload rollout during training. + This avoids remapping live NCCL parameter buffers while still fitting Qwen3-8B + on each 80GB H100. +- Pin Megatron's production attention backend to Transformer Engine `fused` for + CP2/P2P. Backend auto-selection is host-dependent and would make G00/G10 + incomparable across environments. On hosts exposing multiple CUDA runtime + majors, use Transformer Engine 2.18 or newer and select the CUDA 12 runtime + explicitly with `CUDNN_FRONTEND_CUDART_LIB_NAME`. +- GRPO, BF16, `top_p=1.0`, temperature 1, no dropout, fixed training and rollout seeds. +- A 7168-token response budget, one prompt with eight GRPO samples per step, + and full uniform activation recomputation + (`recompute-num-layers=1`). Do not enable expandable CUDA allocator segments: + deterministic TP collectives require CUDA IPC-capable staging allocations. +- vLLM CUDA Graph mode `FULL_DECODE_ONLY`, not eager, with exact capture sizes + `1..(rollout_batch_size × n_samples_per_prompt)`. +- Megatron and vLLM integration readbacks with positive call counts for every + configured route. A production Megatron logp route instead requires VIME's + native-backend runtime marker and rejects any provider hook or provider + readback. +- Production routes reject provenance whose actual backend is RL-Kernel, even + if an outer integration layer labeled the call as production. +- R/R runs must report zero bitwise mismatches, zero max absolute logp + difference, CUDA execution, and no fallback or Triton provenance. +- Append-only run directories. A passing validator creates `COMPLETE`; failed + attempts remain available for audit and are not overwritten. + +The current VIME debug dump does not include training `log_probs` in +`rollout_data`. `validate_run.py` therefore uses VIME's runtime `torch.ne`, +maximum, and mean absolute-difference metrics. Counts are reconstructed from +the sample means and global batch size. The report marks offline tensor +comparison as unavailable instead of claiming it was performed. + +## Recommended phases + +The phase definitions are frozen in `experiment_matrix.json`. + +| Phase | Steps | Seeds | Decision | +|---|---:|---|---| +| short | 8 | 1234 | Catch state transition, weight-update, and cache issues | +| precision | 30 | 1234, 2345, 3456 | Estimate drift distribution before the long run | +| convergence | 200 | 1234 | Primary PR evidence and learning/performance curves | + +Use the paired 200-step runs for the main claim. A bitwise invariant does not +need seed averaging; the three paired 30-step seeds test repeatability and +provide uncertainty estimates for reward, throughput, and overhead. Run groups +in the same seed order and compare paired seeds; report the mean and a 95% +confidence interval. Never merge runs from different code revisions, +checkpoints, prompt hashes, or CUDA Graph settings in one estimate. + +## Prepare DAPO-Math-17k + +`prepare_dapo_data.py` downloads or converts the official Parquet file and +emits VIME `prompt`/`label` JSONL. It deduplicates by `extra_info.index` and +writes source/output hashes and row counts to a sibling manifest. + +```bash +python examples/vime_qwen3_8b_tp4_cp2_200/prepare_dapo_data.py \ + --download \ + --source /data/dapo-math-17k.parquet \ + --output /data/dapo-math-17k.vime.jsonl +``` + +The converter requires `pyarrow`. The small +`qwen3_8b_multiround_math.jsonl` file is a developer fixture and must not be +used for experiment or reward claims. + +## Run one arm + +The complete host setup, data and checkpoint preparation, exact historical +revision table, formal 200-step launch commands, Ray log capture, validation, +and performance-analysis commands are recorded in +[`REPRODUCTION.md`](REPRODUCTION.md). The short example below is schematic; +every expanded path and command is recorded in `manifest.json`. + +```bash +python examples/vime_qwen3_8b_tp4_cp2_200/run_arm.py \ + --group G01 \ + --num-rollout 8 \ + --seed 1234 \ + --rollout-seed 1234 \ + --output-root /data/vime-200/runs/short \ + --rl-kernel-root /path/to/RL-Kernel \ + --vime-root /path/to/vime \ + --megatron-root /path/to/Megatron-LM \ + --model-root /models/Qwen3-8B \ + --ref-load /models/Qwen3-8B_torch_dist \ + --prompt-data /data/dapo-math-17k.vime.jsonl \ + --python /path/to/python \ + --ray-bin /path/to/ray \ + --wait +``` + +`run_arm.py` refuses to reuse an existing run ID. It records repository +revisions, command line, environment, data hash, GPU inventory, topology, +seeds, batch parameters, and CUDA Graph contract before submission. + +After the Ray job finishes, save its combined log as `run.log` in the run +directory and validate it: + +```bash +python examples/vime_qwen3_8b_tp4_cp2_200/validate_run.py \ + --run-dir /data/vime-200/runs/short/ \ + --seal +``` + +## Aggregate and plot + +Only sealed runs are collected. `collect_results.py` writes one row per run, +one row per training step, and group-level summaries. + +```bash +python examples/vime_qwen3_8b_tp4_cp2_200/collect_results.py \ + --runs-root /data/vime-200/runs \ + --output-dir /data/vime-200/results + +python examples/vime_qwen3_8b_tp4_cp2_200/plot_results.py \ + --rounds-csv /data/vime-200/results/rounds.csv \ + --phase convergence \ + --output-dir /data/vime-200/results/figures +``` + +The plotting step requires Matplotlib. It produces: + +- `consistency.png`: mean/max absolute logp difference, mismatch rate, and + mismatch count per step; +- `learning.png`: raw reward with moving average, PPO KL, entropy, and response + truncation ratio; +- `optimization.png`: GRPO policy-gradient loss, clipped ratio fraction, PPO + KL, and gradient norm; +- `performance.png`: end-to-end step time, rollout time, and actor throughput. + +The summary table also reports total active-token exposure, cumulative +bitwise mismatch count, token-weighted mean absolute difference, maximum +absolute difference, reward, truncation, step time, and throughput. For R/R, +the strongest claim is `mismatch_count = 0` over the stated token exposure; +reward and speed are secondary quality and cost measurements. + +## Published convergence results + +The sealed 200-step G10/G11 results, per-step data, reproducible plotting +script, and consistency figures are published in +[`results/convergence_s1234_g10_g11`](results/convergence_s1234_g10_g11/README.md). +G00 and G01 were paused, so the publication is explicitly a two-arm interim +result rather than a completed four-arm ablation. Performance is omitted +because the immutable G11 and final G10 runs used different repository and +Transformer Engine revisions. + +The diagnostic (non-causal) stage-timing comparison can be regenerated from +the two sealed `run.log` files with +[`analyze_performance.py`](analyze_performance.py), following the commands in +[`REPRODUCTION.md`](REPRODUCTION.md#performance-analysis-commands). diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/REPRODUCTION.md b/examples/vime_qwen3_8b_tp4_cp2_200/REPRODUCTION.md new file mode 100644 index 00000000..0697fe4d --- /dev/null +++ b/examples/vime_qwen3_8b_tp4_cp2_200/REPRODUCTION.md @@ -0,0 +1,329 @@ +# Reproduce and profile the Qwen3-8B TP4/CP2 experiment + +This runbook records the commands used for the 200-step G10/G11 experiment and +the commands used to analyse its performance. Run every training replica under +a **new** run ID. The published G11 run +`g11-convergence-s1234-tp4-20260901e` is immutable and must never be submitted +again or overwritten. + +The historical pair is suitable for consistency validation and descriptive +performance diagnosis. It is not a causal performance ablation: G10 and G11 +used different RL-Kernel, VIME, and Transformer Engine revisions. A strict +operator-overhead claim requires a separately predeclared pair with identical +software, prompts, generated token lengths, and profiler settings. + +## Frozen configuration + +| Item | G10 | G11 | +|---|---|---| +| Run ID | `g10-convergence-s1234-tp4-20260901j` | `g11-convergence-s1234-tp4-20260901e` | +| RL-Kernel | `d2173e8d948e8cf062ac36be32bdf53bac75daa0` | `5403df6e3c5244343438916248ccfcc597dd96f6` | +| VIME | `1a113710e80aa7cfc271caa9bd90bcf348a7af08` | `a013293fb6dfdc5cd27152b54f64209ea2691d26` | +| Megatron-LM | `1dcf0dafa884ad52ffb243625717a3471643e087` | same | +| Transformer Engine | 2.18 local wheel | 2.11 local wheel | +| Attention backend | `fused` | `auto` | +| Attention / FFN / logp | P/P, P/P, native VIME | R/R, R/R, strict RL-Kernel | +| Framework logp | rollout logp reused | rollout logp reused | +| CUDA Graph | `FULL_DECODE_ONLY`, capture sizes 1 through 8 | same | + +Both runs used Python 3.11.15, PyTorch 2.9.1, vLLM 0.16.0, Ray 2.57.0, +Megatron TP4/CP2, two TP4 vLLM engines, BF16, one prompt and eight samples per +step, a 7168-token response limit, GRPO, and the `deepscaler` rule reward. A +reference checkpoint was supplied through `--ref-load`; this does not imply a +non-zero KL penalty unless that penalty is enabled by the training config. + +## Host paths and preflight + +The following values reproduce the original host layout. Change them together +when using another machine. + +```bash +set -euo pipefail + +export EXPERIMENT_ROOT=/home/ellm/ljj/vime_qwen3_8b_tp2_cp2_200_experiment +export DATA_ROOT=/data/ellm/vime_qwen3_8b_tp4_cp2_200_experiment +export RLK_ROOT=$EXPERIMENT_ROOT/RL-Kernel +export VIME_ROOT=$EXPERIMENT_ROOT/vime +export MEGATRON_ROOT=/home/ellm/ljj/Megatron-LM +export EXAMPLE_ROOT=$RLK_ROOT/examples/vime_qwen3_8b_tp4_cp2_200 +export RUNTIME_ROOT=/home/ellm/workspace/ljj/.conda/envs/rlk-attention-engines +export PYTHON=$RUNTIME_ROOT/bin/python3.11 +export RAY=$RUNTIME_ROOT/bin/ray +export CUDA_RUNTIME_ROOT=$RUNTIME_ROOT/lib/python3.11/site-packages/nvidia/cuda_runtime +export TE218_ROOT=/home/ellm/ljj/.te-2.18.0-site-packages +export TE211_ROOT=/home/ellm/ljj/.te-2.11.0-site-packages +export RUNTIME_SITE=/home/ellm/ljj/.runtime-site-packages +export CUDA_PYTHON_SITE=/home/ellm/ljj/.cuda-python-12.8-site-packages +export HF_MODEL_ROOT=/home/ellm/ljj/checkpoints/Qwen3-8B_vime_rlkernel_tp2_cp2 +export TORCH_DIST_ROOT=/home/ellm/ljj/checkpoints/Qwen3-8B_torch_dist +export PROMPT_DATA=$DATA_ROOT/datasets/dapo-math-17k.vime.jsonl +export RAY_API_SERVER_ADDRESS=http://127.0.0.1:8265 + +test "$(nvidia-smi --query-gpu=name --format=csv,noheader | wc -l)" -eq 8 +nvidia-smi --query-gpu=index,name,memory.total,driver_version --format=csv +df -h /dev/shm "$DATA_ROOT" +git -C "$RLK_ROOT" status --short +git -C "$VIME_ROOT" status --short +git -C "$MEGATRON_ROOT" status --short +"$PYTHON" - <<'PY' +import importlib.metadata as metadata +import platform + +print("python", platform.python_version()) +for package in ("torch", "vllm", "ray", "sympy", "pylatexenc"): + print(package, metadata.version(package)) +PY +``` + +Formal runs must start from clean repositories. Do not use `--allow-dirty` for +published evidence. + +## Checkout or clone the sources + +Clone once, then select the revisions listed above before launching each arm. +The commands below create the original directory layout without changing an +existing checkout. + +```bash +test -d "$RLK_ROOT/.git" || git clone https://github.com/RL-Align/RL-Kernel.git "$RLK_ROOT" +test -d "$VIME_ROOT/.git" || git clone https://github.com/RL-Align/vime.git "$VIME_ROOT" +test -d "$MEGATRON_ROOT/.git" || git clone https://github.com/NVIDIA/Megatron-LM.git "$MEGATRON_ROOT" + +git -C "$RLK_ROOT" fetch origin +git -C "$VIME_ROOT" fetch origin +git -C "$MEGATRON_ROOT" fetch origin +git -C "$MEGATRON_ROOT" checkout --detach 1dcf0dafa884ad52ffb243625717a3471643e087 +``` + +## Prepare the prompt data and checkpoint + +```bash +mkdir -p "$DATA_ROOT/datasets" +"$PYTHON" "$EXAMPLE_ROOT/prepare_dapo_data.py" \ + --download \ + --source "$DATA_ROOT/datasets/dapo-math-17k.parquet" \ + --output "$PROMPT_DATA" + +test "$(sha256sum "$PROMPT_DATA" | cut -d ' ' -f 1)" = \ + 73e2166517fd635e1157aff17202f86a5cced44ca1669e6f49d2d63a59bf509d +``` + +The actor and reference begin with the same Qwen3-8B weights. If the +`torch_dist` checkpoint is not already present, convert it with the model +arguments shipped by the selected VIME revision: + +```bash +if [ ! -d "$TORCH_DIST_ROOT" ]; then + cd "$VIME_ROOT" + source scripts/models/qwen3-8B.sh + PYTHONPATH="$MEGATRON_ROOT" "$PYTHON" tools/convert_hf_to_torch_dist.py \ + "${MODEL_ARGS[@]}" \ + --hf-checkpoint "$HF_MODEL_ROOT" \ + --save "$TORCH_DIST_ROOT" +fi +test -d "$HF_MODEL_ROOT" +test -d "$TORCH_DIST_ROOT" +``` + +## Start Ray + +The Ray start invocation was not sealed in either run manifest. The following +single-node command matches the observed dashboard, 8-GPU resource count, +200 GB object store, and temporary directory. Record the command used on a new +host because changing the object store or spilling path can affect timing. + +```bash +if ! "$RAY" status >/dev/null 2>&1; then + "$RAY" start --head \ + --include-dashboard=true \ + --dashboard-host=127.0.0.1 \ + --dashboard-port=8265 \ + --num-gpus=8 \ + --object-store-memory=200000000000 \ + --temp-dir=/tmp/rlk-mxs-perf +fi +"$RAY" status +``` + +## Submit one exact arm + +This function is the complete `run_arm.py` invocation used to construct the +Ray runtime environment and full `train.py` command. `run_arm.py` writes the +expanded `ray_command`, `train_command`, environment, revisions, hardware, +dataset hash, seeds, and topology to `manifest.json` before submission. + +```bash +submit_arm() { + local group=$1 + local run_id=$2 + local te_root=$3 + shift 3 + + env -u PYTHONPATH \ + CUDNN_FRONTEND_CUDART_LIB_NAME="$CUDA_RUNTIME_ROOT/lib/libcudart.so.12" \ + "$PYTHON" "$EXAMPLE_ROOT/run_arm.py" \ + --group "$group" \ + --num-rollout 200 \ + --seed 1234 \ + --rollout-seed 1234 \ + --run-id "$run_id" \ + --output-root "$DATA_ROOT/runs/convergence" \ + --rl-kernel-root "$RLK_ROOT" \ + --vime-root "$VIME_ROOT" \ + --megatron-root "$MEGATRON_ROOT" \ + --model-root "$HF_MODEL_ROOT" \ + --ref-load "$TORCH_DIST_ROOT" \ + --prompt-data "$PROMPT_DATA" \ + --python "$PYTHON" \ + --ray-bin "$RAY" \ + --extra-pythonpath "$RUNTIME_SITE" \ + --extra-pythonpath "$CUDA_PYTHON_SITE" \ + --extra-pythonpath "$te_root" \ + --ld-library-path "$CUDA_RUNTIME_ROOT/lib:$te_root/transformer_engine/wheel_lib" \ + "$@" +} +``` + +Use a fresh replica ID for G10. These checkouts reproduce the final G10 stack: + +```bash +git -C "$RLK_ROOT" checkout --detach d2173e8d948e8cf062ac36be32bdf53bac75daa0 +git -C "$VIME_ROOT" checkout --detach 1a113710e80aa7cfc271caa9bd90bcf348a7af08 +export G10_REPLICA_ID=g10-convergence-s1234-replica-$(date -u +%Y%m%d%H%M%S) +submit_arm G10 "$G10_REPLICA_ID" "$TE218_ROOT" +``` + +The following G11 command is intentionally `--dry-run`: it reconstructs and +records the historical command without launching a duplicate G11 job. Never +reuse the sealed ID. Removing `--dry-run` is only appropriate for a separately +approved replica with a new ID and output directory. + +```bash +git -C "$RLK_ROOT" checkout --detach 5403df6e3c5244343438916248ccfcc597dd96f6 +git -C "$VIME_ROOT" checkout --detach a013293fb6dfdc5cd27152b54f64209ea2691d26 +export G11_AUDIT_ID=g11-convergence-s1234-audit-$(date -u +%Y%m%d%H%M%S) +submit_arm G11 "$G11_AUDIT_ID" "$TE211_ROOT" --dry-run +``` + +To audit the fully expanded command without executing it: + +```bash +export AUDIT_MANIFEST=$DATA_ROOT/runs/convergence/$G11_AUDIT_ID/manifest.json +"$PYTHON" - "$AUDIT_MANIFEST" <<'PY' +import json +import shlex +import sys + +manifest = json.load(open(sys.argv[1], encoding="utf-8")) +print(shlex.join(manifest["ray_command"])) +print("\nTRAIN COMMAND\n", shlex.join(manifest["train_command"])) +PY +``` + +## Capture logs, status, and validation evidence + +For a submitted run, wait for a terminal Ray state, then save logs and validate. +The validator seals only a successful run that passes backend, native/provider, +CUDA Graph, mismatch, fallback, traceback, and step-count gates. + +```bash +export RUN_ID=$G10_REPLICA_ID +export SUBMISSION_ID=vime200-$RUN_ID +export RUN_DIR=$DATA_ROOT/runs/convergence/$RUN_ID + +while ! "$RAY" job status --address="$RAY_API_SERVER_ADDRESS" "$SUBMISSION_ID" \ + | grep -Eq 'SUCCEEDED|FAILED|STOPPED'; do + sleep 20 +done +"$RAY" job status --address="$RAY_API_SERVER_ADDRESS" "$SUBMISSION_ID" \ + > "$RUN_DIR/ray-status.txt" +"$RAY" job logs --address="$RAY_API_SERVER_ADDRESS" "$SUBMISSION_ID" \ + > "$RUN_DIR/run.log" + +grep -Ei 'actual_backend|linear-logp|native.*backend|cuda.graph|capture.size|mismatch|abs.diff|fallback|oom|spill|traceback' \ + "$RUN_DIR/run.log" > "$RUN_DIR/runtime-audit.txt" || true +"$PYTHON" "$EXAMPLE_ROOT/validate_run.py" --run-dir "$RUN_DIR" --seal +test -f "$RUN_DIR/COMPLETE" +``` + +## Collect learning and consistency results + +```bash +"$PYTHON" "$EXAMPLE_ROOT/collect_results.py" \ + --runs-root "$DATA_ROOT/runs" \ + --output-dir "$DATA_ROOT/results" + +"$PYTHON" "$EXAMPLE_ROOT/plot_results.py" \ + --rounds-csv "$DATA_ROOT/results/rounds.csv" \ + --phase convergence \ + --output-dir "$DATA_ROOT/results/figures" +``` + +## Performance analysis commands + +The performance analyser reads the emitted per-step timers instead of rounded +progress-bar durations. It separates rollout generation, weight update, +wake/offload residual, actor training, and train residual; controls for response +or total token length with OLS; reports token-normalized rollout and actor +throughput; and bootstraps mean timing gaps with 20,000 draws. + +Create a lightweight analysis environment if the training environment does not +contain NumPy and Matplotlib: + +```bash +export PERF_VENV=$EXPERIMENT_ROOT/.perf-analysis-venv +test -x "$PERF_VENV/bin/python" || python3 -m venv "$PERF_VENV" +"$PERF_VENV/bin/python" -m pip install --upgrade pip +"$PERF_VENV/bin/python" -m pip install numpy==2.4.1 matplotlib==3.10.8 +``` + +Run the sealed historical comparison without copying or editing either log: + +```bash +export G10_RUN=$DATA_ROOT/runs/convergence/g10-convergence-s1234-tp4-20260901j +export G11_RUN=$DATA_ROOT/runs/convergence/g11-convergence-s1234-tp4-20260901e +export PERF_OUT=$DATA_ROOT/results/performance_g10_g11 + +"$PERF_VENV/bin/python" "$EXAMPLE_ROOT/analyze_performance.py" \ + --g10-log "$G10_RUN/run.log" \ + --g11-log "$G11_RUN/run.log" \ + --output-dir "$PERF_OUT" + +column -s, -t < "$PERF_OUT/summary.csv" | less -S +"$PERF_VENV/bin/python" -m json.tool "$PERF_OUT/summary.json" \ + > "$PERF_OUT/summary.pretty.json" +``` + +Outputs are: + +- `step-metrics.csv`: all parsed and derived metrics for every step and arm; +- `summary.csv` and `summary.json`: descriptive statistics, stage gap shares, + length-control regressions, fully truncated and steady-step subsets, speedups, + and bootstrap intervals; +- `performance-decomposition.{png,pdf}`: mean step-time stack; +- `length-controlled-scaling.{png,pdf}`: time versus token length and OLS fits; +- `stage-time-series.{png,pdf}`: raw and 10-step moving-average stage times; +- `token-normalized-throughput.{png,pdf}`: rollout and actor throughput. + +For live diagnostics on a newly approved run, collect utilization and I/O in a +separate shell. These samplers do not alter the training command, but their +overhead and sampling interval must be identical across compared arms. + +```bash +mkdir -p "$RUN_DIR/perf-monitor" +nvidia-smi dmon -s pucvmet -d 1 -o DT > "$RUN_DIR/perf-monitor/nvidia-smi-dmon.log" & +export GPU_MONITOR_PID=$! +iostat -xz 1 > "$RUN_DIR/perf-monitor/iostat.log" & +export IO_MONITOR_PID=$! + +# After the Ray job reaches a terminal state: +kill "$GPU_MONITOR_PID" "$IO_MONITOR_PID" +wait "$GPU_MONITOR_PID" "$IO_MONITOR_PID" 2>/dev/null || true +"$RAY" memory --stats-only \ + > "$RUN_DIR/perf-monitor/ray-memory.txt" || true +``` + +Before interpreting a speedup, check response-length balance, truncation ratio, +CUDA Graph evidence for both engines, disk spilling, OOM/fallback/traceback +events, and the exact repository/TE revisions. Do not attribute the historical +G10/G11 timing gap to one operator because those controls are not matched. diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/aligned_python_entrypoint.sh b/examples/vime_qwen3_8b_tp4_cp2_200/aligned_python_entrypoint.sh new file mode 100755 index 00000000..2b611b6b --- /dev/null +++ b/examples/vime_qwen3_8b_tp4_cp2_200/aligned_python_entrypoint.sh @@ -0,0 +1,100 @@ +#!/usr/bin/env bash +set -euo pipefail + +REAL_PYTHON="${RL_KERNEL_REAL_PYTHON:?RL_KERNEL_REAL_PYTHON must name the real Python executable}" +RL_KERNEL_ROOT="${RL_KERNEL_ROOT:-$(cd -- "$(dirname -- "${BASH_SOURCE[0]}")/../.." && pwd)}" +export RL_KERNEL_ROOT +export PYTHONPATH="${RL_KERNEL_ROOT}:${PYTHONPATH:-}" + +if [[ "${1:-}" == "train.py" || "${1:-}" == */train.py ]]; then + # Ray job submission does not always forward the shell exports used by the + # outer launcher. A strict linear-logp job must still install the matching + # vLLM hooks; otherwise it silently falls back to native attention/FFN and + # loses the R/R performance path. Preserve explicit ablation selections. + strict_linear_logp=0 + rollout_batch_size="" + n_samples_per_prompt="1" + explicit_vllm_execution_config=0 + previous_arg="" + for current_arg in "$@"; do + if [[ "${previous_arg}" == "--linear-logp-provider-mode" && "${current_arg}" == "strict" ]]; then + strict_linear_logp=1 + elif [[ "${previous_arg}" == "--rollout-batch-size" ]]; then + rollout_batch_size="${current_arg}" + elif [[ "${previous_arg}" == "--n-samples-per-prompt" ]]; then + n_samples_per_prompt="${current_arg}" + fi + + case "${current_arg}" in + --linear-logp-provider-mode=strict) + strict_linear_logp=1 + ;; + --rollout-batch-size=*) + rollout_batch_size="${current_arg#*=}" + ;; + --n-samples-per-prompt=*) + n_samples_per_prompt="${current_arg#*=}" + ;; + --vllm-enforce-eager|--vllm-optimization-level|--vllm-optimization-level=*|--vllm-compilation-config|--vllm-compilation-config=*) + explicit_vllm_execution_config=1 + ;; + esac + previous_arg="${current_arg}" + done + + required_cudagraph_args=() + if [[ "${strict_linear_logp}" == "1" ]]; then + export RL_KERNEL_VLLM_INTEGRATION="${RL_KERNEL_VLLM_INTEGRATION:-1}" + export RL_KERNEL_CUDA_ONLY="${RL_KERNEL_CUDA_ONLY:-1}" + export VIME_RL_KERNEL_STRICT="${VIME_RL_KERNEL_STRICT:-1}" + export RL_KERNEL_ATTENTION_CASE="${RL_KERNEL_ATTENTION_CASE:-R/R}" + export RL_KERNEL_FFN_CASE="${RL_KERNEL_FFN_CASE:-R/R}" + export RL_KERNEL_LOGP_CASE="${RL_KERNEL_LOGP_CASE:-R/R}" + fi + + # CUDA Graph is a frozen matrix setting, independent of the linear-logp + # provider route. Capturing the complete decode graph removes per-layer + # host-launch gaps and keeps P/P and R/R performance comparisons aligned. + # Capture every exact batch size; explicit execution flags still win. + if [[ "${explicit_vllm_execution_config}" == "0" ]]; then + if [[ "${rollout_batch_size}" =~ ^[1-9][0-9]*$ && "${n_samples_per_prompt}" =~ ^[1-9][0-9]*$ ]]; then + max_capture_size=$((rollout_batch_size * n_samples_per_prompt)) + if [[ -n "${RL_KERNEL_VLLM_CUDAGRAPH_MAX_CAPTURE_SIZE:-}" ]]; then + max_capture_size="${RL_KERNEL_VLLM_CUDAGRAPH_MAX_CAPTURE_SIZE}" + fi + if ! [[ "${max_capture_size}" =~ ^[1-9][0-9]*$ ]]; then + echo "RL_KERNEL_VLLM_CUDAGRAPH_MAX_CAPTURE_SIZE must be a positive integer" >&2 + exit 2 + fi + + capture_sizes="[" + for ((batch_size = 1; batch_size <= max_capture_size; batch_size++)); do + if ((batch_size > 1)); then + capture_sizes+="," + fi + capture_sizes+="${batch_size}" + done + capture_sizes+="]" + compilation_config="{\"cudagraph_mode\":\"FULL_DECODE_ONLY\",\"cudagraph_capture_sizes\":${capture_sizes},\"max_cudagraph_capture_size\":${max_capture_size}}" + required_cudagraph_args=( + --vllm-optimization-level 0 + --vllm-compilation-config "${compilation_config}" + ) + echo "[RL-Kernel] required vLLM full-decode CUDA Graph capture sizes: ${capture_sizes}" >&2 + else + echo "[RL-Kernel] required CUDA Graph configuration lacks rollout batch sizes" >&2 + exit 2 + fi + fi + exec "${REAL_PYTHON}" "$@" \ + --seed "${RL_KERNEL_SEED:-1234}" \ + --rollout-seed "${RL_KERNEL_ROLLOUT_SEED:-42}" \ + --vllm-enable-deterministic-inference \ + --vllm-attention-backend flash_attn \ + --vllm-disable-custom-all-reduce \ + --deterministic-mode \ + --accumulate-allreduce-grads-in-fp32 \ + "${required_cudagraph_args[@]}" +fi + +exec "${REAL_PYTHON}" "$@" diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/analyze_performance.py b/examples/vime_qwen3_8b_tp4_cp2_200/analyze_performance.py new file mode 100644 index 00000000..7063ea28 --- /dev/null +++ b/examples/vime_qwen3_8b_tp4_cp2_200/analyze_performance.py @@ -0,0 +1,603 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors + +"""Compare stage timing and token-normalized throughput for sealed G10/G11 logs. + +The historical runs used different RL-Kernel, VIME, and Transformer Engine +revisions. The output is therefore descriptive diagnostic evidence, not a +causal estimate of the cost of the RL-Kernel consistency operators. +""" + +from __future__ import annotations + +import argparse +import ast +import csv +import json +import re +import statistics +from pathlib import Path +from typing import Any + +import matplotlib.pyplot as plt +import numpy as np + +RECORD_RE = re.compile(r"\b(rollout|step|perf)\s+(\d+):\s+(\{.*\})\s*$") +GENERATION_RE = re.compile(r"Rollout generation:.*100%.*\[(\d+):(\d+)<") +BLUE = "#2F67D8" +LIGHT_BLUE = "#93B4F4" +RED = "#E53935" +LIGHT_RED = "#F4A3A0" +GRID = "#D7DCE2" +TEXT = "#20242A" + + +def parse_log(path: Path) -> tuple[dict[str, dict[int, dict[str, Any]]], np.ndarray]: + records: dict[str, dict[int, dict[str, Any]]] = { + "rollout": {}, + "step": {}, + "rollout_perf": {}, + "train_perf": {}, + } + progress_seconds: list[float] = [] + for line in path.read_text(encoding="utf-8", errors="replace").splitlines(): + match = RECORD_RE.search(line) + if match: + try: + payload = ast.literal_eval(match.group(3)) + except (SyntaxError, ValueError): + payload = None + if isinstance(payload, dict): + kind = match.group(1) + step = int(match.group(2)) + if kind == "perf": + if "perf/rollout_time" in payload: + kind = "rollout_perf" + elif "perf/step_time" in payload: + kind = "train_perf" + else: + continue + records[kind][step] = payload + if "Rollout generation:" in line and "100%" in line: + match = GENERATION_RE.search(line) + if match: + progress_seconds.append(60 * int(match.group(1)) + int(match.group(2))) + + expected = list(range(200)) + for kind, values in records.items(): + if sorted(values) != expected: + raise RuntimeError( + f"{path}: {kind} has {len(values)} records; expected steps 0..199" + ) + if len(progress_seconds) != 400: + raise RuntimeError( + f"{path}: expected 400 duplicated rollout progress events, got " + f"{len(progress_seconds)}" + ) + pairs = np.asarray(progress_seconds, dtype=float).reshape(200, 2) + if not np.array_equal(pairs[:, 0], pairs[:, 1]): + raise RuntimeError(f"{path}: duplicated rollout progress events do not match") + return records, pairs[:, 0] + + +def array(records: dict[int, dict[str, Any]], key: str) -> np.ndarray: + return np.asarray([float(records[index][key]) for index in range(200)], dtype=float) + + +def rows_array(rows: list[dict[str, Any]], key: str) -> np.ndarray: + return np.asarray([float(row[key]) for row in rows], dtype=float) + + +def describe(values: np.ndarray) -> dict[str, float]: + return { + "mean": float(np.mean(values)), + "median": float(np.median(values)), + "p05": float(np.percentile(values, 5)), + "p95": float(np.percentile(values, 95)), + "min": float(np.min(values)), + "max": float(np.max(values)), + "sum": float(np.sum(values)), + "std": float(np.std(values)), + } + + +def moving_average(values: np.ndarray, window: int = 10) -> np.ndarray: + totals = np.convolve(values, np.ones(window, dtype=float), mode="full")[ + : len(values) + ] + counts = np.minimum(np.arange(1, len(values) + 1), window) + return totals / counts + + +def regression(x: np.ndarray, y: np.ndarray, common_x: float) -> dict[str, float]: + slope, intercept = np.polyfit(x, y, 1) + prediction = intercept + slope * common_x + return { + "slope": float(slope), + "intercept": float(intercept), + "correlation": float(np.corrcoef(x, y)[0, 1]), + "prediction_at_common_x": float(prediction), + "common_x": float(common_x), + } + + +def bootstrap_mean_difference( + left: np.ndarray, right: np.ndarray, *, draws: int = 20_000 +) -> dict[str, float]: + rng = np.random.default_rng(1234) + left_idx = rng.integers(0, len(left), size=(draws, len(left))) + right_idx = rng.integers(0, len(right), size=(draws, len(right))) + gaps = left[left_idx].mean(axis=1) - right[right_idx].mean(axis=1) + return { + "estimate": float(left.mean() - right.mean()), + "ci95_low": float(np.percentile(gaps, 2.5)), + "ci95_high": float(np.percentile(gaps, 97.5)), + "draws": draws, + } + + +def build_rows( + group: str, + records: dict[str, dict[int, dict[str, Any]]], + progress_seconds: np.ndarray, +) -> list[dict[str, float | int | str]]: + rollout = records["rollout"] + rollout_perf = records["rollout_perf"] + train_perf = records["train_perf"] + response_length = array(rollout, "rollout/response_lengths") + total_length = array(rollout, "rollout/total_lengths") + rollout_time = array(rollout_perf, "perf/rollout_time") + train_wait = array(train_perf, "perf/train_wait_time") + train_time = array(train_perf, "perf/train_time") + actor_train = array(train_perf, "perf/actor_train_time") + update_weights = array(train_perf, "perf/update_weights_time") + + rows: list[dict[str, float | int | str]] = [] + for step in range(200): + rows.append( + { + "group": group, + "step": step, + "response_length_mean": response_length[step], + "prompt_length_mean": total_length[step] - response_length[step], + "total_length_mean": total_length[step], + "response_tokens_total": 8 * response_length[step], + "rollout_time_s": rollout_time[step], + "rollout_progress_time_s_rounded": progress_seconds[step], + "rollout_tok_per_gpu_s": float( + rollout_perf[step]["perf/tokens_per_gpu_per_sec"] + ), + "rollout_aggregate_tok_s": 8 + * response_length[step] + / rollout_time[step], + "rollout_truncated_ratio": float( + rollout_perf[step]["rollout/truncated_ratio"] + ), + "update_weights_time_s": update_weights[step], + "wait_residual_time_s": ( + train_wait[step] - rollout_time[step] - update_weights[step] + ), + "train_wait_time_s": train_wait[step], + "actor_train_time_s": actor_train[step], + "train_residual_time_s": train_time[step] - actor_train[step], + "train_time_s": train_time[step], + "data_preprocess_time_s": float( + train_perf[step]["perf/data_preprocess_time"] + ), + "step_time_s": float(train_perf[step]["perf/step_time"]), + "actor_train_tok_s": float( + train_perf[step]["perf/actor_train_tok_per_s"] + ), + "actor_train_tflops": float( + train_perf[step]["perf/actor_train_tflops"] + ), + } + ) + return rows + + +def write_csv(path: Path, rows: list[dict[str, Any]]) -> None: + with path.open("w", encoding="utf-8", newline="") as handle: + writer = csv.DictWriter(handle, fieldnames=list(rows[0])) + writer.writeheader() + writer.writerows(rows) + + +def style_axis(axis: plt.Axes) -> None: + axis.set_facecolor("white") + axis.grid(True, color=GRID, linestyle="--", linewidth=0.8) + axis.set_axisbelow(True) + for spine in axis.spines.values(): + spine.set_color(TEXT) + spine.set_linewidth(1.0) + axis.tick_params(colors=TEXT, labelsize=10) + + +def save_figure(fig: plt.Figure, output_dir: Path, stem: str) -> None: + fig.savefig( + output_dir / f"{stem}.png", dpi=220, bbox_inches="tight", facecolor="white" + ) + fig.savefig(output_dir / f"{stem}.pdf", bbox_inches="tight", facecolor="white") + + +def plot_decomposition( + group_rows: dict[str, list[dict[str, Any]]], output_dir: Path +) -> None: + stages = [ + ("rollout_time_s", "Rollout generation", "#E15759"), + ("update_weights_time_s", "Weight update", "#F2B134"), + ("wait_residual_time_s", "Wake/offload residual", "#BAB0AC"), + ("actor_train_time_s", "Actor training", "#4E79A7"), + ("train_residual_time_s", "Train residual", "#76B7B2"), + ] + groups = ["G11", "G10"] + fig, axis = plt.subplots(figsize=(9.5, 6.5)) + style_axis(axis) + bottoms = np.zeros(2) + for key, label, color in stages: + values = np.asarray( + [ + statistics.fmean(float(row[key]) for row in group_rows[group]) + for group in groups + ] + ) + axis.bar(groups, values, bottom=bottoms, label=label, color=color, width=0.58) + for index, value in enumerate(values): + if value >= 2: + axis.text( + index, + bottoms[index] + value / 2, + f"{value:.1f}s", + ha="center", + va="center", + fontsize=9.5, + color=TEXT, + ) + bottoms += values + for index, total in enumerate(bottoms): + axis.text( + index, total + 2, f"{total:.1f}s / step", ha="center", fontweight="bold" + ) + axis.set_title( + "G11 vs G10 · Mean Step-Time Decomposition", + fontsize=16, + fontweight="bold", + pad=14, + ) + axis.set_ylabel("Seconds per training step") + axis.legend(ncol=2, loc="upper right", frameon=True) + axis.set_ylim(0, 1.16 * float(np.max(bottoms))) + fig.tight_layout() + save_figure(fig, output_dir, "performance-decomposition") + plt.close(fig) + + +def plot_scaling(group_rows: dict[str, list[dict[str, Any]]], output_dir: Path) -> None: + fig, axes = plt.subplots(1, 2, figsize=(14.2, 5.7)) + configs = {"G11": (BLUE, "o"), "G10": (RED, "s")} + panels = [ + ( + "response_length_mean", + "rollout_time_s", + "Rollout generation", + "Mean response length (tokens/sample)", + "Seconds", + ), + ( + "total_length_mean", + "actor_train_time_s", + "Actor training", + "Mean total length (tokens/sample)", + "Seconds", + ), + ] + for axis, (x_key, y_key, title, xlabel, ylabel) in zip(axes, panels, strict=True): + style_axis(axis) + for group, (color, marker) in configs.items(): + x = rows_array(group_rows[group], x_key) + y = rows_array(group_rows[group], y_key) + axis.scatter( + x, + y, + s=23, + alpha=0.30, + color=color, + marker=marker, + label=f"{group} steps", + ) + slope, intercept = np.polyfit(x, y, 1) + grid = np.linspace(float(np.min(x)), float(np.max(x)), 200) + axis.plot( + grid, + intercept + slope * grid, + color=color, + linewidth=2.3, + label=f"{group} OLS", + ) + axis.set_title(title, fontsize=12.5, pad=8) + axis.set_xlabel(xlabel) + axis.set_ylabel(ylabel) + axis.legend(ncol=2, frameon=True, fontsize=9.3) + fig.suptitle( + "G11 vs G10 · Length-Controlled Stage Scaling", + fontsize=17, + fontweight="bold", + y=0.99, + ) + fig.tight_layout(rect=(0, 0, 1, 0.95)) + save_figure(fig, output_dir, "length-controlled-scaling") + plt.close(fig) + + +def plot_time_series( + group_rows: dict[str, list[dict[str, Any]]], output_dir: Path +) -> None: + steps = np.arange(200) + fig, axes = plt.subplots(2, 1, figsize=(13.5, 7.8), sharex=True) + for group, raw_color, ma_color in ( + ("G11", LIGHT_BLUE, BLUE), + ("G10", LIGHT_RED, RED), + ): + for axis, key in zip( + axes, ("rollout_time_s", "actor_train_time_s"), strict=True + ): + values = rows_array(group_rows[group], key) + axis.plot(steps, values, color=raw_color, alpha=0.7, linewidth=0.9) + axis.plot( + steps, + moving_average(values), + color=ma_color, + linewidth=2.2, + label=f"{group} 10-step MA", + ) + for axis, title in zip(axes, ("Rollout generation", "Actor training"), strict=True): + style_axis(axis) + axis.set_title(title, fontsize=12.5, pad=8) + axis.set_ylabel("Seconds") + axis.legend(frameon=True) + axes[-1].set_xlabel("Training step") + axes[-1].set_xlim(0, 199) + axes[-1].set_xticks(np.arange(0, 200, 25)) + fig.suptitle( + "G11 vs G10 · Performance Across 200 Steps", + fontsize=17, + fontweight="bold", + y=0.99, + ) + fig.tight_layout(rect=(0, 0, 1, 0.96)) + save_figure(fig, output_dir, "stage-time-series") + plt.close(fig) + + +def plot_throughput( + group_rows: dict[str, list[dict[str, Any]]], output_dir: Path +) -> None: + fig, axes = plt.subplots(1, 2, figsize=(12.2, 5.3)) + panels = [ + ("rollout_tok_per_gpu_s", "Rollout throughput", "Tokens / GPU / s"), + ("actor_train_tok_s", "Actor training throughput", "Tokens / s"), + ] + for axis, (key, title, ylabel) in zip(axes, panels, strict=True): + style_axis(axis) + means = [rows_array(group_rows[group], key).mean() for group in ("G11", "G10")] + bars = axis.bar( + ["G11", "G10"], + means, + width=0.5, + color=["#A9C7ED", "#F2B566"], + edgecolor=[BLUE, "#D66A27"], + linewidth=1.2, + ) + for bar, value in zip(bars, means, strict=True): + axis.text( + bar.get_x() + bar.get_width() / 2, + value * 1.018, + f"{value:,.0f}", + ha="center", + color=TEXT, + ) + speedup = means[1] / means[0] + axis.text( + 0.5, + max(means) * 0.88, + f"G10 / G11 = {speedup:.2f}×", + ha="center", + fontweight="bold", + color=TEXT, + ) + axis.set_title(title, fontsize=12.5, pad=8) + axis.set_ylabel(ylabel) + axis.set_ylim(0, max(means) * 1.14) + fig.suptitle( + "G11 vs G10 · Token-Normalized Throughput", + fontsize=17, + fontweight="bold", + y=0.99, + ) + fig.tight_layout(rect=(0, 0, 1, 0.94)) + save_figure(fig, output_dir, "token-normalized-throughput") + plt.close(fig) + + +def summarize(group_rows: dict[str, list[dict[str, Any]]]) -> dict[str, Any]: + keys = [ + "response_length_mean", + "prompt_length_mean", + "total_length_mean", + "rollout_time_s", + "rollout_progress_time_s_rounded", + "rollout_tok_per_gpu_s", + "rollout_aggregate_tok_s", + "rollout_truncated_ratio", + "update_weights_time_s", + "wait_residual_time_s", + "train_wait_time_s", + "actor_train_time_s", + "train_residual_time_s", + "train_time_s", + "step_time_s", + "actor_train_tok_s", + "actor_train_tflops", + ] + all_rows = group_rows["G11"] + group_rows["G10"] + common_response_length = statistics.fmean( + float(row["response_length_mean"]) for row in all_rows + ) + common_total_length = statistics.fmean( + float(row["total_length_mean"]) for row in all_rows + ) + summary: dict[str, Any] = {"groups": {}} + for group, rows in group_rows.items(): + summary["groups"][group] = { + key: describe(rows_array(rows, key)) for key in keys + } + summary["groups"][group]["rollout_length_regression"] = regression( + rows_array(rows, "response_length_mean"), + rows_array(rows, "rollout_time_s"), + common_response_length, + ) + summary["groups"][group]["actor_length_regression"] = regression( + rows_array(rows, "total_length_mean"), + rows_array(rows, "actor_train_time_s"), + common_total_length, + ) + saturated = [ + row + for row in rows + if float(row["response_length_mean"]) == 7168.0 + and float(row["rollout_truncated_ratio"]) == 1.0 + ] + summary["groups"][group]["fully_truncated_subset"] = { + "count": len(saturated), + "rollout_time_s": describe(rows_array(saturated, "rollout_time_s")), + "actor_train_time_s": describe(rows_array(saturated, "actor_train_time_s")), + } + steady = rows[1:] + summary["groups"][group]["steady_steps_1_199"] = { + "step_time_s": describe(rows_array(steady, "step_time_s")), + "rollout_time_s": describe(rows_array(steady, "rollout_time_s")), + "actor_train_time_s": describe(rows_array(steady, "actor_train_time_s")), + } + + g11 = summary["groups"]["G11"] + g10 = summary["groups"]["G10"] + step_gap = g11["step_time_s"]["mean"] - g10["step_time_s"]["mean"] + stages = [ + "rollout_time_s", + "update_weights_time_s", + "wait_residual_time_s", + "actor_train_time_s", + "train_residual_time_s", + ] + summary["gap"] = { + "g11_minus_g10_step_time_s": step_gap, + "g11_minus_g10_total_hours": ( + g11["step_time_s"]["sum"] - g10["step_time_s"]["sum"] + ) + / 3600, + "g10_step_time_reduction_fraction": 1 + - g10["step_time_s"]["mean"] / g11["step_time_s"]["mean"], + "g10_end_to_end_speedup": g11["step_time_s"]["mean"] + / g10["step_time_s"]["mean"], + "g10_rollout_throughput_speedup": g10["rollout_tok_per_gpu_s"]["mean"] + / g11["rollout_tok_per_gpu_s"]["mean"], + "g10_actor_throughput_speedup": g10["actor_train_tok_s"]["mean"] + / g11["actor_train_tok_s"]["mean"], + "rollout_common_length_gap_s": g11["rollout_length_regression"][ + "prediction_at_common_x" + ] + - g10["rollout_length_regression"]["prediction_at_common_x"], + "actor_common_length_gap_s": g11["actor_length_regression"][ + "prediction_at_common_x" + ] + - g10["actor_length_regression"]["prediction_at_common_x"], + "stage_contributions": {}, + "bootstrap_mean_gap_ci95": {}, + } + for stage in stages: + gap = g11[stage]["mean"] - g10[stage]["mean"] + summary["gap"]["stage_contributions"][stage] = { + "g11_minus_g10_s": gap, + "share_of_step_gap": gap / step_gap, + } + for key in ("step_time_s", "rollout_time_s", "actor_train_time_s"): + summary["gap"]["bootstrap_mean_gap_ci95"][key] = bootstrap_mean_difference( + rows_array(group_rows["G11"], key), rows_array(group_rows["G10"], key) + ) + summary["method_notes"] = { + "rollout_time": "Exact perf/rollout_time emitted by RolloutManager; tqdm integer seconds are retained only as an audit cross-check.", + "wait_residual": "train_wait - exact rollout_time - update_weights_time; captures wake/offload/orchestration outside the two named timers.", + "throughput": "Uses emitted perf/tokens_per_gpu_per_sec and perf/actor_train_tok_per_s, avoiding response-length confounding.", + "common_length": "Separate OLS fits for each arm evaluated at the pooled mean length; descriptive rather than causal because the two runs use different operator stacks, TE versions, revisions and generated sequences.", + "bootstrap": "Independent non-parametric bootstrap of per-step mean gaps with fixed seed 1234 and 20,000 draws.", + } + return summary + + +def main() -> None: + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument( + "--data-dir", + type=Path, + help="directory containing g10.run.log and g11.run.log", + ) + parser.add_argument("--g10-log", type=Path, help="sealed G10 run.log") + parser.add_argument("--g11-log", type=Path, help="sealed G11 run.log") + parser.add_argument("--output-dir", type=Path, required=True) + args = parser.parse_args() + if args.data_dir is not None and ( + args.g10_log is not None or args.g11_log is not None + ): + parser.error("use either --data-dir or both --g10-log/--g11-log") + if args.data_dir is not None: + g10_log = args.data_dir / "g10.run.log" + g11_log = args.data_dir / "g11.run.log" + elif args.g10_log is not None and args.g11_log is not None: + g10_log = args.g10_log + g11_log = args.g11_log + else: + parser.error("provide --data-dir or both --g10-log and --g11-log") + args.output_dir.mkdir(parents=True, exist_ok=True) + + parsed = { + "G11": parse_log(g11_log), + "G10": parse_log(g10_log), + } + group_rows = { + group: build_rows(group, records, progress) + for group, (records, progress) in parsed.items() + } + all_rows = group_rows["G11"] + group_rows["G10"] + write_csv(args.output_dir / "step-metrics.csv", all_rows) + summary = summarize(group_rows) + (args.output_dir / "summary.json").write_text( + json.dumps(summary, indent=2, sort_keys=True) + "\n", encoding="utf-8" + ) + + comparison_rows: list[dict[str, Any]] = [] + for key, g11_value in summary["groups"]["G11"].items(): + g10_value = summary["groups"]["G10"].get(key) + if not isinstance(g11_value, dict) or "mean" not in g11_value: + continue + g11_mean = g11_value["mean"] + g10_mean = g10_value["mean"] + comparison_rows.append( + { + "metric": key, + "g11_mean": g11_mean, + "g10_mean": g10_mean, + "g11_minus_g10": g11_mean - g10_mean, + "g10_over_g11": g10_mean / g11_mean if g11_mean else None, + } + ) + write_csv(args.output_dir / "summary.csv", comparison_rows) + + plot_decomposition(group_rows, args.output_dir) + plot_scaling(group_rows, args.output_dir) + plot_time_series(group_rows, args.output_dir) + plot_throughput(group_rows, args.output_dir) + print(json.dumps(summary["gap"], indent=2, sort_keys=True)) + + +if __name__ == "__main__": + main() diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/collect_results.py b/examples/vime_qwen3_8b_tp4_cp2_200/collect_results.py new file mode 100644 index 00000000..70366da5 --- /dev/null +++ b/examples/vime_qwen3_8b_tp4_cp2_200/collect_results.py @@ -0,0 +1,250 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors + +"""Collect sealed VIME ablation runs into reproducible CSV and JSON tables.""" + +from __future__ import annotations + +import argparse +import ast +import csv +import json +import math +import re +from collections import defaultdict +from pathlib import Path +from statistics import mean +from typing import Any + + +RECORD_RE = re.compile(r"\b(rollout|step|perf)\s+(\d+):\s+(\{.*\})\s*$") + + +def _load(path: Path) -> dict[str, Any]: + value = json.loads(path.read_text(encoding="utf-8")) + if not isinstance(value, dict): + raise ValueError(f"expected a JSON object in {path}") + return value + + +def _records(path: Path) -> dict[str, dict[int, dict[str, Any]]]: + result: dict[str, dict[int, dict[str, Any]]] = { + "rollout": {}, + "step": {}, + "perf": {}, + } + for line in path.read_text(encoding="utf-8", errors="replace").splitlines(): + match = RECORD_RE.search(line) + if not match: + continue + try: + value = ast.literal_eval(match.group(3)) + except (SyntaxError, ValueError): + continue + if isinstance(value, dict): + result[match.group(1)].setdefault(int(match.group(2)), {}).update(value) + return result + + +def _value(mapping: dict[str, Any], name: str) -> float | None: + value = mapping.get(name) + if isinstance(value, (int, float)) and not isinstance(value, bool): + return float(value) + return None + + +def _run_rows(run_dir: Path) -> tuple[dict[str, Any], list[dict[str, Any]]]: + manifest = _load(run_dir / "manifest.json") + validation = _load(run_dir / "run-validation.json") + if not validation.get("passed"): + raise ValueError(f"sealed run has a failed validation: {run_dir}") + records = _records(run_dir / "run.log") + validation_rows = { + int(row["step"]): row for row in validation["train_rollout_logprob"]["rows"] + } + steps = sorted( + set(records["rollout"]) | set(records["step"]) | set(records["perf"]) + ) + rows = [] + for index in steps: + rollout = records["rollout"].get(index, {}) + step = records["step"].get(index, {}) + perf = records["perf"].get(index, {}) + exact = validation_rows.get(index, {}) + rows.append( + { + "phase": run_dir.parent.name, + "run_id": manifest["run_id"], + "group": manifest["arm"]["group"], + "seed": manifest["seed"], + "rollout_seed": manifest["rollout_seed"], + "step": index, + "framework_consistency": manifest["arm"][ + "framework_use_rollout_logprobs" + ], + "operator_case": manifest["arm"]["logp_case"], + "reward": _value(rollout, "rollout/rewards"), + "raw_reward": _value(rollout, "rollout/raw_reward"), + "response_length": _value(rollout, "rollout/response_lengths"), + "truncated_ratio": _value(rollout, "rollout/truncated"), + "rollout_kl": _value(rollout, "rollout/kl"), + "train_loss": _value(step, "train/loss"), + "pg_loss": _value(step, "train/pg_loss"), + "pg_clipfrac": _value(step, "train/pg_clipfrac"), + "entropy": _value(step, "train/entropy_loss"), + "ppo_kl": _value(step, "train/ppo_kl"), + "grad_norm": _value(step, "train/grad_norm"), + "mean_abs_dlogp": _value(step, "train/train_rollout_logprob_abs_diff"), + "max_abs_dlogp": _value( + step, "train/train_current_rollout_logprob_max_abs_diff" + ), + "mismatch_count": exact.get("bitwise_mismatch_count"), + "active_token_count": exact.get("active_token_count"), + "rollout_time": _value(perf, "perf/rollout_time"), + "train_time": _value(perf, "perf/train_time"), + "step_time": _value(perf, "perf/step_time"), + "update_weights_time": _value(perf, "perf/update_weights_time"), + "actor_tokens_per_second": _value(perf, "perf/actor_train_tok_per_s"), + "actor_train_tflops": _value(perf, "perf/actor_train_tflops"), + } + ) + run = { + "phase": run_dir.parent.name, + "run_id": manifest["run_id"], + "group": manifest["arm"]["group"], + "seed": manifest["seed"], + "rollout_seed": manifest["rollout_seed"], + "num_rollout": manifest["num_rollout"], + "rl_kernel_revision": manifest["revisions"]["rl_kernel"], + "vime_revision": manifest["revisions"]["vime"], + "prompt_data_sha256": manifest["prompt_data_sha256"], + "cudagraph_passed": validation["cudagraph"]["passed"], + "validation_passed": validation["passed"], + "offline_tensor_comparison": validation["offline_tensor_comparison"]["status"], + "rounds_observed": len(rows), + } + return run, rows + + +def _finite(rows: list[dict[str, Any]], name: str) -> list[float]: + values = [] + for row in rows: + value = row.get(name) + if isinstance(value, (int, float)) and math.isfinite(float(value)): + values.append(float(value)) + return values + + +def _summaries(rows: list[dict[str, Any]]) -> list[dict[str, Any]]: + grouped: dict[tuple[str, str], list[dict[str, Any]]] = defaultdict(list) + for row in rows: + grouped[(str(row["phase"]), str(row["group"]))].append(row) + summaries = [] + for (phase, group), items in sorted(grouped.items()): + mismatches = _finite(items, "mismatch_count") + tokens = _finite(items, "active_token_count") + mean_abs = _finite(items, "mean_abs_dlogp") + weighted_abs_numerator = sum( + float(item["mean_abs_dlogp"]) * float(item["active_token_count"]) + for item in items + if item.get("mean_abs_dlogp") is not None + and item.get("active_token_count") is not None + ) + token_total = sum(tokens) + summaries.append( + { + "phase": phase, + "group": group, + "run_count": len({str(item["run_id"]) for item in items}), + "round_count": len(items), + "active_token_exposure": token_total, + "bitwise_mismatch_count": sum(mismatches), + "bitwise_mismatch_rate": ( + sum(mismatches) / token_total if token_total else None + ), + "mean_abs_dlogp_token_weighted": ( + weighted_abs_numerator / token_total if token_total else None + ), + "max_abs_dlogp": max(_finite(items, "max_abs_dlogp"), default=None), + "reward_mean": ( + mean(_finite(items, "reward")) if _finite(items, "reward") else None + ), + "raw_reward_mean": ( + mean(_finite(items, "raw_reward")) + if _finite(items, "raw_reward") + else None + ), + "truncated_ratio_mean": ( + mean(_finite(items, "truncated_ratio")) + if _finite(items, "truncated_ratio") + else None + ), + "step_time_mean": ( + mean(_finite(items, "step_time")) + if _finite(items, "step_time") + else None + ), + "actor_tokens_per_second_mean": ( + mean(_finite(items, "actor_tokens_per_second")) + if _finite(items, "actor_tokens_per_second") + else None + ), + "unweighted_mean_abs_dlogp": mean(mean_abs) if mean_abs else None, + } + ) + return summaries + + +def _write_csv(path: Path, rows: list[dict[str, Any]]) -> None: + if not rows: + raise ValueError(f"no rows to write to {path}") + with path.open("w", encoding="utf-8", newline="") as handle: + writer = csv.DictWriter( + handle, + fieldnames=list(rows[0]), + lineterminator="\n", + ) + writer.writeheader() + writer.writerows(rows) + + +def main() -> int: + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--runs-root", type=Path, required=True) + parser.add_argument("--output-dir", type=Path, required=True) + args = parser.parse_args() + run_dirs = sorted( + path.parent + for path in args.runs_root.rglob("COMPLETE") + if (path.parent / "manifest.json").is_file() + and (path.parent / "run-validation.json").is_file() + and (path.parent / "run.log").is_file() + ) + if not run_dirs: + raise ValueError(f"no sealed runs found under {args.runs_root}") + runs: list[dict[str, Any]] = [] + rounds: list[dict[str, Any]] = [] + for run_dir in run_dirs: + run, run_rounds = _run_rows(run_dir) + runs.append(run) + rounds.extend(run_rounds) + summaries = _summaries(rounds) + args.output_dir.mkdir(parents=True, exist_ok=True) + _write_csv(args.output_dir / "runs.csv", runs) + _write_csv(args.output_dir / "rounds.csv", rounds) + _write_csv(args.output_dir / "summary.csv", summaries) + report = { + "schema_version": "rlkernel.vime_qwen3_8b_tp4_cp2_200.results.v1", + "sealed_run_count": len(runs), + "round_count": len(rounds), + "summaries": summaries, + } + (args.output_dir / "summary.json").write_text( + json.dumps(report, indent=2, sort_keys=True) + "\n", encoding="utf-8" + ) + print(json.dumps(report, indent=2, sort_keys=True)) + return 0 + + +if __name__ == "__main__": + raise SystemExit(main()) diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/experiment_matrix.json b/examples/vime_qwen3_8b_tp4_cp2_200/experiment_matrix.json new file mode 100644 index 00000000..56bd372f --- /dev/null +++ b/examples/vime_qwen3_8b_tp4_cp2_200/experiment_matrix.json @@ -0,0 +1,53 @@ +{ + "schema_version": "rlkernel.vime_qwen3_8b_tp4_cp2_200.matrix.v1", + "groups": { + "G00": { + "framework_use_rollout_logprobs": false, + "attention_case": "P/P", + "ffn_case": "P/P", + "logp_case": "P/P", + "description": "native VIME baseline" + }, + "G10": { + "framework_use_rollout_logprobs": true, + "attention_case": "P/P", + "ffn_case": "P/P", + "logp_case": "P/P", + "description": "VIME framework-level consistency only" + }, + "G01": { + "framework_use_rollout_logprobs": false, + "attention_case": "R/R", + "ffn_case": "R/R", + "logp_case": "R/R", + "description": "RL-Kernel operator-level consistency only" + }, + "G11": { + "framework_use_rollout_logprobs": true, + "attention_case": "R/R", + "ffn_case": "R/R", + "logp_case": "R/R", + "description": "framework-level and operator-level consistency" + } + }, + "phases": { + "short": {"num_rollout": 8, "seeds": [1234]}, + "precision": {"num_rollout": 8, "seeds": [1234, 2345, 3456]}, + "convergence": {"num_rollout": 200, "seeds": [1234]}, + "module_ablation": { + "num_rollout": 8, + "seeds": [1234], + "framework_use_rollout_logprobs": false, + "groups": { + "M000": {"attention_case": "P/P", "ffn_case": "P/P", "logp_case": "P/P"}, + "M100": {"attention_case": "R/R", "ffn_case": "P/P", "logp_case": "P/P"}, + "M010": {"attention_case": "P/P", "ffn_case": "R/R", "logp_case": "P/P"}, + "M001": {"attention_case": "P/P", "ffn_case": "P/P", "logp_case": "R/R"}, + "M110": {"attention_case": "R/R", "ffn_case": "R/R", "logp_case": "P/P"}, + "M101": {"attention_case": "R/R", "ffn_case": "P/P", "logp_case": "R/R"}, + "M011": {"attention_case": "P/P", "ffn_case": "R/R", "logp_case": "R/R"}, + "M111": {"attention_case": "R/R", "ffn_case": "R/R", "logp_case": "R/R"} + } + } + } +} diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/plot_results.py b/examples/vime_qwen3_8b_tp4_cp2_200/plot_results.py new file mode 100644 index 00000000..f2fa4fbd --- /dev/null +++ b/examples/vime_qwen3_8b_tp4_cp2_200/plot_results.py @@ -0,0 +1,216 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors + +"""Render consistency, learning, and throughput figures from rounds.csv.""" + +from __future__ import annotations + +import argparse +import csv +import math +from collections import defaultdict +from pathlib import Path +from statistics import mean, pstdev +from typing import Any + + +COLORS = {"G00": "#6b7280", "G10": "#2563eb", "G01": "#dc2626", "G11": "#059669"} +MARKERS = {"G00": "o", "G10": "x", "G01": "D", "G11": "+"} +LABELS = { + "G00": "G00: production", + "G10": "G10: framework", + "G01": "G01: operator", + "G11": "G11: framework + operator", +} + + +def _parse(value: str) -> Any: + if value == "": + return None + if value in {"True", "False"}: + return value == "True" + try: + return float(value) + except ValueError: + return value + + +def _load(path: Path, phase: str | None) -> list[dict[str, Any]]: + with path.open(encoding="utf-8", newline="") as handle: + rows = [ + {key: _parse(value) for key, value in row.items()} + for row in csv.DictReader(handle) + ] + if phase is not None: + rows = [row for row in rows if row["phase"] == phase] + if not rows: + raise ValueError("no rows match the requested phase") + return rows + + +def _series( + rows: list[dict[str, Any]], metric: str +) -> dict[str, tuple[list[int], list[float], list[float]]]: + grouped: dict[tuple[str, int], list[float]] = defaultdict(list) + for row in rows: + value = row.get(metric) + if isinstance(value, (int, float)) and math.isfinite(float(value)): + grouped[(str(row["group"]), int(row["step"]))].append(float(value)) + result = {} + for group in sorted({key[0] for key in grouped}): + steps = sorted(step for candidate, step in grouped if candidate == group) + centers = [mean(grouped[(group, step)]) for step in steps] + spreads = [pstdev(grouped[(group, step)]) for step in steps] + result[group] = (steps, centers, spreads) + return result + + +def _moving_average(values: list[float], window: int) -> list[float]: + return [ + mean(values[max(0, index - window + 1) : index + 1]) + for index in range(len(values)) + ] + + +def _plot_series(axis, rows, metric, title, *, moving_average=1, symlog=False): + for group, (steps, values, spreads) in _series(rows, metric).items(): + centers = _moving_average(values, moving_average) + color = COLORS.get(group) + axis.plot( + steps, + centers, + label=LABELS.get(group, group), + color=color, + linewidth=2, + marker=MARKERS.get(group, "o"), + markersize=5, + ) + if any(spreads): + lower = [ + center - spread for center, spread in zip(centers, spreads, strict=True) + ] + upper = [ + center + spread for center, spread in zip(centers, spreads, strict=True) + ] + axis.fill_between(steps, lower, upper, color=color, alpha=0.15) + if symlog: + axis.set_yscale("symlog", linthresh=1e-9) + axis.set_title(title) + axis.set_xlabel("training step") + axis.grid(alpha=0.25) + + +def _mismatch_rate_rows(rows: list[dict[str, Any]]) -> list[dict[str, Any]]: + result = [] + for row in rows: + updated = dict(row) + mismatch = row.get("mismatch_count") + tokens = row.get("active_token_count") + updated["mismatch_rate"] = ( + float(mismatch) / float(tokens) + if isinstance(mismatch, (int, float)) + and isinstance(tokens, (int, float)) + and tokens + else None + ) + result.append(updated) + return result + + +def _save_consistency(rows, output, dpi): + import matplotlib.pyplot as plt + + derived = _mismatch_rate_rows(rows) + figure, axes = plt.subplots(2, 2, figsize=(13, 8), constrained_layout=True) + _plot_series( + axes[0, 0], + derived, + "mean_abs_dlogp", + "Mean |train logp - rollout logp|", + symlog=True, + ) + _plot_series(axes[0, 1], derived, "max_abs_dlogp", "Maximum |Δlogp|", symlog=True) + _plot_series(axes[1, 0], derived, "mismatch_rate", "Bitwise mismatch rate") + _plot_series( + axes[1, 1], + derived, + "mismatch_count", + "Bitwise mismatch count per step", + symlog=True, + ) + handles, labels = axes[0, 0].get_legend_handles_labels() + figure.legend(handles, labels, loc="outside upper center", ncol=max(1, len(labels))) + figure.savefig(output / "consistency.png", dpi=dpi) + plt.close(figure) + + +def _save_learning(rows, output, dpi, window): + import matplotlib.pyplot as plt + + figure, axes = plt.subplots(2, 2, figsize=(13, 8), constrained_layout=True) + _plot_series( + axes[0, 0], + rows, + "raw_reward", + f"Raw reward ({window}-step MA)", + moving_average=window, + ) + _plot_series(axes[0, 1], rows, "ppo_kl", "Training PPO KL", symlog=True) + _plot_series(axes[1, 0], rows, "entropy", "Token entropy") + _plot_series(axes[1, 1], rows, "truncated_ratio", "Truncated response ratio") + handles, labels = axes[0, 0].get_legend_handles_labels() + figure.legend(handles, labels, loc="outside upper center", ncol=max(1, len(labels))) + figure.savefig(output / "learning.png", dpi=dpi) + plt.close(figure) + + +def _save_optimization(rows, output, dpi): + import matplotlib.pyplot as plt + + figure, axes = plt.subplots(2, 2, figsize=(13, 8), constrained_layout=True) + _plot_series( + axes[0, 0], rows, "pg_loss", "GRPO policy-gradient loss", symlog=True + ) + _plot_series(axes[0, 1], rows, "pg_clipfrac", "Policy ratio clipped fraction") + _plot_series(axes[1, 0], rows, "ppo_kl", "Training PPO KL", symlog=True) + _plot_series(axes[1, 1], rows, "grad_norm", "Gradient norm", symlog=True) + handles, labels = axes[0, 0].get_legend_handles_labels() + figure.legend(handles, labels, loc="outside upper center", ncol=max(1, len(labels))) + figure.savefig(output / "optimization.png", dpi=dpi) + plt.close(figure) + + +def _save_performance(rows, output, dpi): + import matplotlib.pyplot as plt + + figure, axes = plt.subplots(1, 3, figsize=(15, 4.5), constrained_layout=True) + _plot_series(axes[0], rows, "step_time", "End-to-end step time (s)") + _plot_series(axes[1], rows, "rollout_time", "Rollout time (s)") + _plot_series(axes[2], rows, "actor_tokens_per_second", "Actor tokens/s") + handles, labels = axes[0].get_legend_handles_labels() + figure.legend(handles, labels, loc="outside upper center", ncol=max(1, len(labels))) + figure.savefig(output / "performance.png", dpi=dpi) + plt.close(figure) + + +def main() -> int: + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--rounds-csv", type=Path, required=True) + parser.add_argument("--output-dir", type=Path, required=True) + parser.add_argument("--phase", default=None) + parser.add_argument("--moving-average", type=int, default=10) + parser.add_argument("--dpi", type=int, default=180) + args = parser.parse_args() + if args.moving_average <= 0: + raise ValueError("--moving-average must be positive") + rows = _load(args.rounds_csv, args.phase) + args.output_dir.mkdir(parents=True, exist_ok=True) + _save_consistency(rows, args.output_dir, args.dpi) + _save_learning(rows, args.output_dir, args.dpi, args.moving_average) + _save_optimization(rows, args.output_dir, args.dpi) + _save_performance(rows, args.output_dir, args.dpi) + return 0 + + +if __name__ == "__main__": + raise SystemExit(main()) diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/prepare_dapo_data.py b/examples/vime_qwen3_8b_tp4_cp2_200/prepare_dapo_data.py new file mode 100644 index 00000000..8027b182 --- /dev/null +++ b/examples/vime_qwen3_8b_tp4_cp2_200/prepare_dapo_data.py @@ -0,0 +1,126 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors + +"""Download and convert DAPO-Math-17k to VIME prompt/label JSONL.""" + +from __future__ import annotations + +import argparse +import hashlib +import json +import urllib.request +from pathlib import Path +from typing import Any + + +DEFAULT_URL = ( + "https://huggingface.co/datasets/BytedTsinghua-SIA/DAPO-Math-17k/" + "resolve/main/data/train-00000-of-00001.parquet" +) + + +def _sha256(path: Path) -> str: + digest = hashlib.sha256() + with path.open("rb") as handle: + for chunk in iter(lambda: handle.read(8 * 1024 * 1024), b""): + digest.update(chunk) + return digest.hexdigest() + + +def _download(url: str, destination: Path) -> None: + destination.parent.mkdir(parents=True, exist_ok=True) + partial = destination.with_suffix(destination.suffix + ".partial") + with urllib.request.urlopen(url) as response, partial.open("wb") as output: + while chunk := response.read(8 * 1024 * 1024): + output.write(chunk) + partial.replace(destination) + + +def _index(row: dict[str, Any]) -> str: + extra = row.get("extra_info") or {} + value = extra.get("index") if isinstance(extra, dict) else None + if not isinstance(value, str) or not value: + raise ValueError("DAPO row is missing extra_info.index") + return value + + +def convert(source: Path, output: Path) -> dict[str, Any]: + try: + import pyarrow.parquet as parquet + except ImportError as exc: + raise RuntimeError("prepare_dapo_data.py requires pyarrow") from exc + + output.parent.mkdir(parents=True, exist_ok=True) + partial = output.with_suffix(output.suffix + ".partial") + seen: set[str] = set() + rows_read = 0 + rows_written = 0 + with partial.open("w", encoding="utf-8") as destination: + parquet_file = parquet.ParquetFile(source) + for batch in parquet_file.iter_batches( + batch_size=4096, + columns=["data_source", "prompt", "ability", "reward_model", "extra_info"], + ): + for row in batch.to_pylist(): + rows_read += 1 + row_id = _index(row) + if row_id in seen: + continue + seen.add(row_id) + reward_model = row.get("reward_model") or {} + label = reward_model.get("ground_truth") + prompt = row.get("prompt") + if ( + not isinstance(prompt, list) + or not isinstance(label, str) + or not label + ): + raise ValueError(f"invalid prompt or ground truth for row {row_id}") + record = { + "prompt": prompt, + "label": label, + "data_source": row.get("data_source"), + "ability": row.get("ability"), + "source_index": row_id, + "reward_style": reward_model.get("style"), + } + destination.write( + json.dumps(record, ensure_ascii=False, separators=(",", ":")) + ) + destination.write("\n") + rows_written += 1 + partial.replace(output) + return { + "schema_version": "rlkernel.dapo_math_17k.v1", + "source": str(source.resolve()), + "source_sha256": _sha256(source), + "rows_read": rows_read, + "rows_written": rows_written, + "duplicates_removed": rows_read - rows_written, + "output": str(output.resolve()), + "output_sha256": _sha256(output), + } + + +def main() -> int: + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--source", type=Path, required=True) + parser.add_argument("--output", type=Path, required=True) + parser.add_argument("--download-url", default=DEFAULT_URL) + parser.add_argument("--download", action="store_true") + args = parser.parse_args() + if args.download and not args.source.exists(): + _download(args.download_url, args.source) + if not args.source.is_file(): + raise FileNotFoundError(args.source) + manifest = convert(args.source, args.output) + manifest_path = args.output.with_suffix(args.output.suffix + ".manifest.json") + manifest_path.write_text( + json.dumps(manifest, indent=2, sort_keys=True) + "\n", encoding="utf-8" + ) + print(json.dumps(manifest, indent=2, sort_keys=True)) + return 0 + + +if __name__ == "__main__": + raise SystemExit(main()) diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/qwen3_8b_multiround_math.jsonl b/examples/vime_qwen3_8b_tp4_cp2_200/qwen3_8b_multiround_math.jsonl new file mode 100644 index 00000000..64d5ae24 --- /dev/null +++ b/examples/vime_qwen3_8b_tp4_cp2_200/qwen3_8b_multiround_math.jsonl @@ -0,0 +1,32 @@ +{"prompt":[{"role":"user","content":"Solve 1+1. Give a concise solution and end with the final answer in \\boxed{}."}],"label":"2"} +{"prompt":[{"role":"user","content":"Solve 12*13. Give a concise solution and end with the final answer in \\boxed{}."}],"label":"156"} +{"prompt":[{"role":"user","content":"Compute 3/4 + 1/8. Give a concise solution and end with the final answer in \\boxed{}."}],"label":"7/8"} +{"prompt":[{"role":"user","content":"Solve 2x+3=11. Give a concise solution and end with the final answer in \\boxed{}."}],"label":"4"} +{"prompt":[{"role":"user","content":"What is 15 percent of 80? Give a concise solution and end with the final answer in \\boxed{}."}],"label":"12"} +{"prompt":[{"role":"user","content":"A rectangle has sides 7 and 4. What is its perimeter? Give a concise solution and end with the final answer in \\boxed{}."}],"label":"22"} +{"prompt":[{"role":"user","content":"What is the average of 6, 8, and 10? Give a concise solution and end with the final answer in \\boxed{}."}],"label":"8"} +{"prompt":[{"role":"user","content":"Compute 2^5+3. Give a concise solution and end with the final answer in \\boxed{}."}],"label":"35"} +{"prompt":[{"role":"user","content":"What is the positive square root of 144? Give a concise solution and end with the final answer in \\boxed{}."}],"label":"12"} +{"prompt":[{"role":"user","content":"If x^2=49 and x is positive, find x. Give a concise solution and end with the final answer in \\boxed{}."}],"label":"7"} +{"prompt":[{"role":"user","content":"Compute 17*19. Give a concise solution and end with the final answer in \\boxed{}."}],"label":"323"} +{"prompt":[{"role":"user","content":"Compute (2/3)/(4/5). Give a concise solution and end with the final answer in \\boxed{}."}],"label":"5/6"} +{"prompt":[{"role":"user","content":"An item costs 80 and is discounted by 25 percent. What is the sale price? Give a concise solution and end with the final answer in \\boxed{}."}],"label":"60"} +{"prompt":[{"role":"user","content":"What is the probability of getting heads twice in two fair coin flips? Give a concise solution and end with the final answer in \\boxed{}."}],"label":"1/4"} +{"prompt":[{"role":"user","content":"The sequence is 2,5,8,11,... What is its 10th term? Give a concise solution and end with the final answer in \\boxed{}."}],"label":"29"} +{"prompt":[{"role":"user","content":"Solve 3x+2=20. Give a concise solution and end with the final answer in \\boxed{}."}],"label":"6"} +{"prompt":[{"role":"user","content":"A triangle has base 10 and height 6. What is its area? Give a concise solution and end with the final answer in \\boxed{}."}],"label":"30"} +{"prompt":[{"role":"user","content":"Five workers finish a job in 12 days at the same rate. How many days would ten workers need? Give a concise solution and end with the final answer in \\boxed{}."}],"label":"6"} +{"prompt":[{"role":"user","content":"Write 0.125 as a reduced fraction. Give a concise solution and end with the final answer in \\boxed{}."}],"label":"1/8"} +{"prompt":[{"role":"user","content":"Solve x/3+2=7. Give a concise solution and end with the final answer in \\boxed{}."}],"label":"15"} +{"prompt":[{"role":"user","content":"What is the sum of the first 10 positive integers? Give a concise solution and end with the final answer in \\boxed{}."}],"label":"55"} +{"prompt":[{"role":"user","content":"What is 30 percent of 250? Give a concise solution and end with the final answer in \\boxed{}."}],"label":"75"} +{"prompt":[{"role":"user","content":"A right triangle has legs 3 and 4. What is its hypotenuse? Give a concise solution and end with the final answer in \\boxed{}."}],"label":"5"} +{"prompt":[{"role":"user","content":"Two numbers are in the ratio 2:3 and sum to 25. What is the larger number? Give a concise solution and end with the final answer in \\boxed{}."}],"label":"15"} +{"prompt":[{"role":"user","content":"Reduce 18/24 to lowest terms. Give a concise solution and end with the final answer in \\boxed{}."}],"label":"3/4"} +{"prompt":[{"role":"user","content":"Compute 4^3-2^3. Give a concise solution and end with the final answer in \\boxed{}."}],"label":"56"} +{"prompt":[{"role":"user","content":"A car travels at 40 km/h for 2.5 hours. How far does it travel? Give a concise solution and end with the final answer in \\boxed{}."}],"label":"100"} +{"prompt":[{"role":"user","content":"What is the probability that a fair six-sided die shows an even number? Give a concise solution and end with the final answer in \\boxed{}."}],"label":"1/2"} +{"prompt":[{"role":"user","content":"Compute 9^2-7^2. Give a concise solution and end with the final answer in \\boxed{}."}],"label":"32"} +{"prompt":[{"role":"user","content":"Solve 5x-4=31. Give a concise solution and end with the final answer in \\boxed{}."}],"label":"7"} +{"prompt":[{"role":"user","content":"What is the sum of the first 20 positive integers? Give a concise solution and end with the final answer in \\boxed{}."}],"label":"210"} +{"prompt":[{"role":"user","content":"If a number is increased by 7 to become 19, what was the original number? Give a concise solution and end with the final answer in \\boxed{}."}],"label":"12"} diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/qwen3_8b_tp4_cp2.json b/examples/vime_qwen3_8b_tp4_cp2_200/qwen3_8b_tp4_cp2.json new file mode 100644 index 00000000..b749c46e --- /dev/null +++ b/examples/vime_qwen3_8b_tp4_cp2_200/qwen3_8b_tp4_cp2.json @@ -0,0 +1,43 @@ +{ + "schema_version": "rlkernel.vime_qwen3_8b_tp4_cp2.v1", + "model": "Qwen/Qwen3-8B", + "training": { + "framework": "megatron", + "tensor_model_parallel_size": 4, + "context_parallel_size": 2, + "pipeline_model_parallel_size": 1, + "world_size": 8, + "dtype": "bf16" + }, + "rollout": { + "framework": "vllm", + "top_p": 1.0, + "logprobs_mode": "processed_logprobs" + }, + "linear_logp_provider": { + "path": "rl_engine.integrations.vime.linear_logp_provider.provider", + "mode": "strict", + "backend_id": "rlkernel.linear_logp.bitwise.v1", + "real_vocab_size": 151936, + "padded_vocab_size": 152064, + "num_vocab_tiles": 64 + }, + "operator_evidence": { + "logp": { + "training": "rl-kernel-provider", + "rollout": "vllm-native-processed-logprobs", + "required_runtime_marker": "linear_logp provider active" + }, + "attention": { + "training": "runtime-readback-required", + "rollout": "runtime-readback-required", + "status": "not_claimed_without_Megatron_and_vLLM_readback" + }, + "ffn": { + "training": "runtime-readback-required", + "rollout": "runtime-readback-required", + "status": "not_claimed_without_Megatron_and_vLLM_readback" + } + }, + "vime_script": "scripts/run-qwen3-8B-rlkernel-tp4-cp2.sh" +} diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/README.md b/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/README.md new file mode 100644 index 00000000..7bf6f5c9 --- /dev/null +++ b/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/README.md @@ -0,0 +1,92 @@ +# G10/G11 200-step convergence results + +This directory publishes the completed half of the Qwen3-8B TP4/CP2 +consistency matrix. It contains one sealed 200-step run for G10 and one for +G11. G00 and G01 were paused and are intentionally not represented here, so +these artifacts must not be interpreted as a completed four-arm ablation. + +## Result + +Both runs used one 8xH100 node, a TP4/CP2 Megatron actor, two colocated TP4 +vLLM engines, seed 1234, eight samples per prompt, and the same DAPO-Math-17k +data hash. Both validators passed the `FULL_DECODE_ONLY` CUDA Graph contract +with capture sizes 1 through 8 and found no fallback or runtime-route error. + +| Group | Runtime route | Active tokens | Bitwise mismatches | Mismatch rate | Token-weighted mean abs dlogp | Max abs dlogp | Mean raw reward | +|---|---|---:|---:|---:|---:|---:|---:| +| G11 | strict RL-Kernel R/R | 9,806,995 | 0 | 0 | 0 | 0 | 0.3950 | +| G10 | VIME-native P/P baseline | 9,927,045 | 5,764,529 | 58.0689% | 0.012567 | 1.191781 | 0.3794 | + +G11 therefore satisfies the strongest claim supported by this experiment: +zero runtime bitwise train/rollout log-probability mismatches over 9.81 million +active tokens. G10 is the native production comparison, not a failure gate; +its nonzero drift measures the train/rollout numerical gap of that route. +The result table reports a token-weighted mean absolute difference, while the +summary figure labels its separate unweighted mean across the 200 step means. + +The raw-reward curves are close because both runs start from the same model, +prompt order, sampling seed, reward function, and rollout topology. Reward is +an outcome-level, relatively coarse metric; it is much less sensitive than the +token-level log-probability comparison. With one seed, the small reward-mean +difference is descriptive rather than a statistical quality claim. + +## Figures + +![Per-step consistency trajectories](figures/consistency-trajectories.png) + +![Consistency summary](figures/consistency-summary.png) + +![Reward and optimization dynamics](figures/training-dynamics.png) + +The trajectory figure compares per-step mismatch counts directly. The summary +figure reports the cumulative comparison: 0 mismatches for G11 strict +RL-Kernel versus 5,764,529 for the G10 VIME-native baseline. + +The near-zero G11 scalar training loss and PPO KL do not mean that gradients +were absent. With rollout log-probability reuse, the pre-update policy ratio is +exactly one; GRPO group-centers advantages, so positive and negative terms can +cancel in the reported scalar while their parameter derivatives remain +nonzero. The plotted PPO KL is the old/rollout-policy diagnostic, not KL to a +reference model. + +Neither run activated a reference model: `kl_coef=0`, `use_kl_loss=false`, and +`kl_loss_coef=0`. A `ref_load` path was recorded, but VIME does not load the +reference checkpoint under those settings. This is intentional for the +train/rollout consistency study and keeps reference regularization from +changing its objective. + +## Provenance and limitations + +| Group | Run ID | RL-Kernel | VIME | Megatron | Transformer Engine | +|---|---|---|---|---|---| +| G11 | `g11-convergence-s1234-tp4-20260901e` | `5403df6` | `a013293` | `1dcf0da` | 2.11 | +| G10 | `g10-convergence-s1234-tp4-20260901j` | `d2173e8` | `1a113710` | `1dcf0da` | 2.18 | + +The G11 run is the immutable sealed run selected by the experiment protocol. +The later G10 run includes production-route verification and CUDA Graph/provider +decoupling fixes. Because the repository revisions and Transformer Engine +versions differ, these two artifacts establish route-specific consistency but +are not used for a performance comparison. No performance figure is published. + +Stopped G10h and G10i attempts are audit-only and excluded. Offline tensor +comparison is unavailable because the VIME debug dump did not capture training +`log_probs`; the accepted consistency evidence is VIME's runtime `torch.ne`, +maximum, and mean absolute-difference instrumentation recorded by the sealed +validator. + +## Artifacts + +- `runs.csv`: run identities, revisions, hashes, and validation status. +- `rounds.csv`: 400 per-step records (200 each for G10 and G11). +- `summary.csv` and `summary.json`: machine-readable aggregate results. +- `plot_report.py`: reproduces the three PNG/PDF figures from `rounds.csv`. +- `figures/`: publication-ready PNG and PDF outputs. + +Regenerate the figures from this directory with: + +```bash +python3 plot_report.py --rounds-csv rounds.csv --output-dir figures +``` + +The plotting command requires Matplotlib; the CSV/JSON artifacts themselves +use only standard text formats. diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/figures/consistency-summary.pdf b/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/figures/consistency-summary.pdf new file mode 100644 index 00000000..0769d50b Binary files /dev/null and b/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/figures/consistency-summary.pdf differ diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/figures/consistency-summary.png b/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/figures/consistency-summary.png new file mode 100644 index 00000000..2a4c340e Binary files /dev/null and b/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/figures/consistency-summary.png differ diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/figures/consistency-trajectories.pdf b/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/figures/consistency-trajectories.pdf new file mode 100644 index 00000000..29bcb6cb Binary files /dev/null and b/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/figures/consistency-trajectories.pdf differ diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/figures/consistency-trajectories.png b/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/figures/consistency-trajectories.png new file mode 100644 index 00000000..6e2dfda5 Binary files /dev/null and b/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/figures/consistency-trajectories.png differ diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/figures/training-dynamics.pdf b/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/figures/training-dynamics.pdf new file mode 100644 index 00000000..1248c942 Binary files /dev/null and b/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/figures/training-dynamics.pdf differ diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/figures/training-dynamics.png b/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/figures/training-dynamics.png new file mode 100644 index 00000000..1805c5d1 Binary files /dev/null and b/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/figures/training-dynamics.png differ diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/plot_report.py b/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/plot_report.py new file mode 100644 index 00000000..4007a508 --- /dev/null +++ b/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/plot_report.py @@ -0,0 +1,347 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors + +"""Plot the sealed G10/G11 consistency report from collected round records.""" + +from __future__ import annotations + +import argparse +import csv +import math +from collections import defaultdict +from pathlib import Path +from statistics import mean +from typing import Any, Callable + + +BLUE = "#2F67D8" +LIGHT_BLUE = "#93B4F4" +RED = "#E53935" +LIGHT_RED = "#F4A3A0" +BAR_BLUE = "#A9C7ED" +BAR_BLUE_EDGE = "#356AC3" +BAR_ORANGE = "#F2B566" +BAR_ORANGE_EDGE = "#D66A27" +GRID = "#D7DCE2" +TEXT = "#20242A" +GROUP_ORDER = ("G11", "G10") +LABELS = {"G11": "G11 strict RL-Kernel", "G10": "G10 VIME-native baseline"} + + +def _parse(value: str) -> Any: + if value == "": + return None + if value in {"True", "False"}: + return value == "True" + try: + return float(value) + except ValueError: + return value + + +def load_rounds(path: Path) -> dict[str, list[dict[str, Any]]]: + grouped: dict[str, list[dict[str, Any]]] = defaultdict(list) + with path.open(encoding="utf-8", newline="") as handle: + for raw in csv.DictReader(handle): + row = {key: _parse(value) for key, value in raw.items()} + group = str(row["group"]) + if group in GROUP_ORDER: + grouped[group].append(row) + for group in GROUP_ORDER: + grouped[group].sort(key=lambda row: int(row["step"])) + steps = [int(row["step"]) for row in grouped[group]] + if steps != list(range(200)): + raise ValueError(f"{group}: expected steps 0..199, found {steps[:2]}..{steps[-2:]}") + return grouped + + +def series(rows: list[dict[str, Any]], metric: str) -> list[float]: + values = [row.get(metric) for row in rows] + if any(not isinstance(value, (int, float)) for value in values): + raise ValueError(f"metric {metric!r} is incomplete") + return [float(value) for value in values] + + +def trailing_ma(values: list[float], window: int = 10) -> list[float]: + return [mean(values[max(0, index - window + 1) : index + 1]) for index in range(len(values))] + + +def percentile(values: list[float], quantile: float) -> float: + ordered = sorted(values) + location = (len(ordered) - 1) * quantile + lower = math.floor(location) + upper = math.ceil(location) + if lower == upper: + return ordered[lower] + return ordered[lower] + (ordered[upper] - ordered[lower]) * (location - lower) + + +def style_axis(axis) -> None: + axis.set_facecolor("white") + axis.grid(True, color=GRID, linestyle="--", linewidth=0.8) + axis.set_axisbelow(True) + for spine in axis.spines.values(): + spine.set_color(TEXT) + spine.set_linewidth(1.0) + axis.tick_params(colors=TEXT, labelsize=10) + + +def save(figure, output_dir: Path, stem: str) -> None: + output_dir.mkdir(parents=True, exist_ok=True) + figure.savefig(output_dir / f"{stem}.png", dpi=220, bbox_inches="tight", facecolor="white") + figure.savefig(output_dir / f"{stem}.pdf", bbox_inches="tight", facecolor="white") + + +def add_report_header(figure, title: str, subtitle: str) -> None: + """Place a non-overlapping title/subtitle pair at the top of a report figure.""" + figure.suptitle(title, fontsize=19, fontweight="bold", y=0.99) + figure.text( + 0.5, + 0.945, + subtitle, + ha="center", + va="top", + fontsize=11.5, + color="#4A515A", + ) + + +def plot_pair(axis, steps, strict_values, native_values) -> None: + axis.plot(steps, strict_values, color=LIGHT_BLUE, linewidth=1.0, alpha=0.78) + axis.plot(steps, trailing_ma(strict_values), color=BLUE, linewidth=2.35) + axis.plot(steps, native_values, color=LIGHT_RED, linewidth=1.0, alpha=0.72) + axis.plot(steps, trailing_ma(native_values), color=RED, linewidth=2.35) + + +def consistency_trajectories(rows, output_dir: Path) -> None: + import matplotlib.pyplot as plt + from matplotlib.lines import Line2D + + steps = list(range(200)) + fig, axes = plt.subplots(3, 1, figsize=(14.5, 10.8), sharex=True) + add_report_header( + fig, + "VIME · Qwen3-8B · 200-step GRPO (8×H100, TP4/CP2)", + "Consistency dynamics: G11 strict RL-Kernel vs G10 VIME-native baseline", + ) + panels = ( + ("mean_abs_dlogp", r"Mean absolute $\Delta\log p$", "Mean |Δ logp|"), + ("max_abs_dlogp", r"Maximum absolute $\Delta\log p$", "Max |Δ logp|"), + ) + for axis, (metric, title, ylabel) in zip(axes[:2], panels, strict=True): + style_axis(axis) + strict_values = series(rows["G11"], metric) + native_values = series(rows["G10"], metric) + plot_pair(axis, steps, strict_values, native_values) + axis.set_title(title, fontsize=12.5, pad=8) + axis.set_ylabel(ylabel, fontsize=11.5) + upper = max(native_values) + axis.set_ylim(-0.025 * upper, 1.08 * upper) + mismatch = {group: series(rows[group], "mismatch_count") for group in GROUP_ORDER} + style_axis(axes[2]) + plot_pair(axes[2], steps, mismatch["G11"], mismatch["G10"]) + axes[2].set_title("Bitwise log-probability mismatch count per step", fontsize=12.5, pad=8) + axes[2].set_ylabel("Mismatch count", fontsize=11.5) + axes[2].set_xlabel("Training Step", fontsize=12) + mismatch_upper = max(mismatch["G10"]) + axes[2].set_ylim(-0.025 * mismatch_upper, 1.08 * mismatch_upper) + axes[2].set_xlim(0, 199) + axes[2].set_xticks(range(0, 200, 25)) + fig.legend( + handles=[ + Line2D([0], [0], color=LIGHT_RED, label="G10 per-step"), + Line2D([0], [0], color=RED, linewidth=2.4, label="G10 10-step MA"), + Line2D([0], [0], color=BLUE, linewidth=1.9, label="G11 strict (exact zero)"), + ], + loc="upper left", + bbox_to_anchor=(0.095, 0.905), + ncol=3, + frameon=True, + facecolor="white", + edgecolor="#C9CDD2", + fontsize=10.5, + ) + fig.tight_layout(rect=(0.025, 0.03, 0.99, 0.845), h_pad=1.15) + save(fig, output_dir, "consistency-trajectories") + plt.close(fig) + + +def consistency_summary(rows, output_dir: Path) -> None: + import matplotlib.pyplot as plt + from matplotlib.patches import Patch + + summaries = {} + for group in GROUP_ORDER: + mismatches = series(rows[group], "mismatch_count") + tokens = series(rows[group], "active_token_count") + summaries[group] = { + "agreement": 100.0 * (1.0 - sum(mismatches) / sum(tokens)), + "mismatch_count": sum(mismatches), + "mean_abs": mean(series(rows[group], "mean_abs_dlogp")), + "p95_max": percentile(series(rows[group], "max_abs_dlogp"), 0.95), + } + panels: tuple[tuple[str, str, str, Callable[[float], str]], ...] = ( + ("agreement", "Bitwise agreement", "%", lambda value: f"{value:.2f}%"), + ( + "mean_abs", + r"Unweighted mean per-step $|\Delta\log p|$", + "Mean per-step |Δ logp|", + lambda value: "0" if value == 0 else f"{value:.5f}", + ), + ( + "p95_max", + r"P95 maximum $|\Delta\log p|$", + "P95 max |Δ logp|", + lambda value: "0" if value == 0 else f"{value:.4f}", + ), + ( + "mismatch_count", + "Total bitwise mismatches", + "Mismatch count", + lambda value: f"{value:,.0f}", + ), + ) + fig, axes = plt.subplots(2, 2, figsize=(13.5, 8.7)) + add_report_header( + fig, + "VIME Consistency Summary · Qwen3-8B · 200-step GRPO", + "G11 strict RL-Kernel vs G10 VIME-native baseline", + ) + for axis, (key, title, ylabel, formatter) in zip(axes.flat, panels, strict=True): + style_axis(axis) + values = [summaries[group][key] for group in GROUP_ORDER] + bars = axis.bar( + [0, 1], + values, + width=0.46, + color=[BAR_BLUE, BAR_ORANGE], + edgecolor=[BAR_BLUE_EDGE, BAR_ORANGE_EDGE], + linewidth=1.4, + ) + axis.set_title(title, fontsize=12.5, pad=8) + axis.set_ylabel(ylabel, fontsize=11) + axis.set_xticks([0, 1], ["G11 strict", "G10 native"]) + axis.grid(False, axis="x") + if key == "mismatch_count": + from matplotlib.ticker import StrMethodFormatter + + axis.yaxis.set_major_formatter(StrMethodFormatter("{x:,.0f}")) + upper = max(values) + axis.set_ylim(0, 1.18 * upper if upper else 1) + for bar, value, edge in zip(bars, values, [BAR_BLUE_EDGE, BAR_ORANGE_EDGE], strict=True): + axis.text( + bar.get_x() + bar.get_width() / 2, + value + 0.025 * (upper or 1), + formatter(value), + ha="center", + fontsize=10.5, + color=edge, + fontweight="bold", + ) + fig.legend( + handles=[ + Patch(facecolor=BAR_BLUE, edgecolor=BAR_BLUE_EDGE, label=LABELS["G11"]), + Patch(facecolor=BAR_ORANGE, edgecolor=BAR_ORANGE_EDGE, label=LABELS["G10"]), + ], + loc="upper center", + bbox_to_anchor=(0.5, 0.895), + ncol=2, + frameon=True, + facecolor="white", + edgecolor="#C9CDD2", + fontsize=10.5, + ) + fig.text( + 0.5, + 0.022, + "Agreement = 1 − total bitwise mismatches / total active tokens; " + "mismatch-count bars span all 200 training steps.", + ha="center", + fontsize=9.5, + color="#606872", + ) + fig.tight_layout(rect=(0.03, 0.06, 0.99, 0.825), h_pad=2.1, w_pad=2.0) + save(fig, output_dir, "consistency-summary") + plt.close(fig) + + +def training_dynamics(rows, output_dir: Path) -> None: + import matplotlib.pyplot as plt + from matplotlib.lines import Line2D + from matplotlib.ticker import ScalarFormatter + + steps = list(range(200)) + panels = ( + ("raw_reward", "Rollout raw reward", "Raw reward", False), + ("train_loss", "Training loss", "Loss", True), + ("ppo_kl", "Train PPO KL", "PPO KL", True), + ("pg_clipfrac", "PPO policy-ratio clip fraction", "Clip fraction", True), + ) + fig, axes = plt.subplots(4, 1, figsize=(14.5, 14.0), sharex=True) + add_report_header( + fig, + "VIME · Qwen3-8B · 200-step GRPO (8×H100, TP4/CP2)", + "Training dynamics: G11 strict RL-Kernel vs G10 VIME-native baseline", + ) + for axis, (metric, title, ylabel, zero_line) in zip(axes, panels, strict=True): + style_axis(axis) + strict_values = series(rows["G11"], metric) + native_values = series(rows["G10"], metric) + plot_pair(axis, steps, strict_values, native_values) + if zero_line: + axis.axhline(0, color="#8B939D", linewidth=0.9, linestyle="--") + axis.set_title(title, fontsize=12.5, pad=8) + axis.set_ylabel(ylabel, fontsize=11.5) + lower = min(min(strict_values), min(native_values), 0.0 if zero_line else math.inf) + upper = max(max(strict_values), max(native_values), 0.0 if zero_line else -math.inf) + span = upper - lower or max(abs(upper), 1e-8) + axis.set_ylim(lower - 0.07 * span, upper + 0.07 * span) + if metric in {"train_loss", "ppo_kl"}: + formatter = ScalarFormatter(useMathText=True) + formatter.set_powerlimits((-3, 3)) + axis.yaxis.set_major_formatter(formatter) + axes[-1].set_xlabel("Training Step", fontsize=12) + axes[-1].set_xlim(0, 199) + axes[-1].set_xticks(range(0, 200, 25)) + fig.legend( + handles=[ + Line2D([0], [0], color=LIGHT_BLUE, label="G11 per-step"), + Line2D([0], [0], color=BLUE, linewidth=2.5, label="G11 10-step MA"), + Line2D([0], [0], color=LIGHT_RED, label="G10 per-step"), + Line2D([0], [0], color=RED, linewidth=2.5, label="G10 10-step MA"), + ], + loc="upper left", + bbox_to_anchor=(0.095, 0.905), + ncol=4, + frameon=True, + facecolor="white", + edgecolor="#C9CDD2", + fontsize=10.5, + ) + fig.text( + 0.986, 0.015, "Ratio metric: train/pg_clipfrac", ha="right", fontsize=9.3, color="#68717C" + ) + fig.tight_layout(rect=(0.025, 0.028, 0.99, 0.85), h_pad=1.1) + save(fig, output_dir, "training-dynamics") + plt.close(fig) + + +def main() -> int: + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--rounds-csv", type=Path, default=Path("rounds.csv")) + parser.add_argument("--output-dir", type=Path, default=Path("figures")) + args = parser.parse_args() + import matplotlib.pyplot as plt + + plt.rcParams.update( + {"font.family": "DejaVu Sans", "font.size": 10.5, "figure.facecolor": "white"} + ) + rows = load_rounds(args.rounds_csv) + consistency_trajectories(rows, args.output_dir) + consistency_summary(rows, args.output_dir) + training_dynamics(rows, args.output_dir) + return 0 + + +if __name__ == "__main__": + raise SystemExit(main()) diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/rounds.csv b/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/rounds.csv new file mode 100644 index 00000000..b167a3b7 --- /dev/null +++ b/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/rounds.csv @@ -0,0 +1,401 @@ +phase,run_id,group,seed,rollout_seed,step,framework_consistency,operator_case,reward,raw_reward,response_length,truncated_ratio,rollout_kl,train_loss,pg_loss,pg_clipfrac,entropy,ppo_kl,grad_norm,mean_abs_dlogp,max_abs_dlogp,mismatch_count,active_token_count,rollout_time,train_time,step_time,update_weights_time,actor_tokens_per_second,actor_train_tflops +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,0,True,P/P,-3.725290298461914e-09,0.875,6224.5,0.0,0.0,4.0920451283454895e-05,4.0920451283454895e-05,0.00479112658649683,0.2703757882118225,0.0009074220433831215,0.5684469883039317,0.012120857834815979,0.6449768543243408,28793.0,49796.0,66.38983988761902,14.01167893409729,101.14814448356628,15.651453495025635,3664.0339111627586,23.40111855574183 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,1,True,P/P,0.0,0.75,1872.875,0.0,0.0,0.0003606565296649933,0.0003606565296649933,0.004792533814907074,0.2821452021598816,0.0005411133752204478,1.1886790595964172,0.01283409632742405,0.39875030517578125,9735.0,14983.0,18.011947631835938,5.7568440437316895,43.52386736869812,14.786272048950195,2807.3989566640353,16.566681741745942 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,2,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.29148590564727783,0.000903440872207284,0.0,0.012715650722384453,0.6922649145126343,33486.0,57344.0,67.59457468986511,6.333472967147827,93.80023503303528,14.855926275253296,9409.472325643894,61.05110863254142 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,3,True,P/P,0.0,0.5,6640.375,0.375,0.0,0.0001495853066444397,0.0001495853066444397,0.0034133412409573793,0.2758287191390991,0.0006878783460706472,0.5141960051846164,0.012148238718509674,0.7331101894378662,31039.0,53123.0,66.24389886856079,6.377356052398682,92.41868710517883,14.845603227615356,8624.261727074225,55.466679835480456 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,4,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.33314695954322815,0.0008298142347484827,0.0,0.01418837159872055,0.7521541118621826,40089.0,57344.0,67.38078737258911,6.478285074234009,93.4860508441925,14.828738689422607,9100.38983198467,58.96921216057516 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,5,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3323231339454651,0.0007460463675670326,0.0,0.014367220923304558,0.8157639503479004,38351.0,57344.0,67.46892142295837,6.4851953983306885,93.88024520874023,14.826070547103882,9147.095727236167,59.3224393987542 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,6,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3847472071647644,0.0011771873105317354,0.0,0.015973418951034546,0.7159717082977295,45156.0,57344.0,67.39624357223511,6.404048681259155,93.97422695159912,15.121649742126465,9222.347001907696,59.77069610880832 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,7,True,P/P,0.0,0.0,4738.5,0.0,0.0,0.0,0.0,0.0,0.4053499698638916,0.0009968606755137444,0.0,0.014486962929368019,0.7170383930206299,26428.0,37908.0,61.20672917366028,5.4058005809783936,86.41077661514282,14.969678401947021,7317.505746853442,45.65465070301993 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,8,True,P/P,-3.725290298461914e-09,0.875,4922.5,0.0,0.0,-0.0001212824136018753,-0.0001212824136018753,0.004630363080650568,0.23015181720256805,0.00084553228225559,0.5951896080068688,0.010534742847084999,0.6438437700271606,22203.0,39380.0,50.92466378211975,5.510761260986328,76.49048089981079,15.153024435043335,7382.185963814407,46.07908735354079 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,9,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.43740057945251465,0.0007440592162311077,0.0,0.016963258385658264,0.7328830361366272,46917.0,57344.0,67.99629330635071,6.38693904876709,94.40991854667664,15.097957134246826,9413.05614001392,61.14723196346941 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,10,True,P/P,1.4901161193847656e-08,0.375,6353.125,0.5,0.0,0.00014613568782806396,0.00014613568782806396,0.003131676698103547,0.3399895429611206,0.000670980429276824,0.6249132096111497,0.013060002587735653,0.7455546855926514,31557.0,50825.0,65.89491367340088,6.322525978088379,92.1033399105072,15.02694821357727,8294.050973737403,53.11322272563074 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,11,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.4585250914096832,0.0006513426778838038,0.0,0.01645077019929886,0.6775834560394287,43997.0,57344.0,67.82830953598022,6.348127365112305,94.12309336662292,14.981138229370117,9389.679960044341,60.92995957088571 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,12,True,P/P,0.0,0.5,4174.375,0.0,0.0,-3.910437226295471e-05,-3.910437226295471e-05,0.00358419306576252,0.2151423692703247,0.0006257632048800588,0.7482235055063424,0.010273435153067112,0.5051577091217041,17010.0,33395.0,45.05213737487793,4.5881593227386475,69.39601612091064,14.856159210205078,7672.172282628071,47.32012025889719 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,13,True,P/P,-3.725290298461914e-09,0.875,4290.875,0.0,0.0,0.0002124495804309845,0.0002124495804309845,0.0051225051283836365,0.27130717039108276,0.0005801066290587187,0.8109315993254208,0.012389028444886208,0.5364299416542053,21671.0,34327.0,43.18650507926941,4.742377758026123,67.75360441207886,15.117964267730713,7578.966127534194,46.77146280603078 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,14,True,P/P,-1.4901161193847656e-08,0.625,4237.875,0.0,0.0,0.0003873556852340698,0.0003873556852340698,0.004298387095332146,0.23583880066871643,0.00043751904740929604,0.6870980821888364,0.011153671890497208,0.6332810521125793,19296.0,33903.0,43.96708559989929,5.316413640975952,69.02921891212463,14.891580581665039,6603.241275616848,40.69907780084455 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,15,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.34094175696372986,0.0007889511180110276,0.0,0.014646660536527634,0.7661538124084473,42189.0,57344.0,67.71523976325989,6.343425512313843,93.75518155097961,14.821417093276978,9327.67241061209,60.4718858665828 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,16,True,P/P,0.0,0.75,4350.875,0.125,0.0,1.979619264602661e-05,1.979619264602661e-05,0.004776099696755409,0.27704596519470215,0.0009653666638769209,0.6813714006728129,0.012729418464004993,0.5915727615356445,21198.0,34807.0,62.018779277801514,4.737968683242798,86.58805251121521,14.951040029525757,7835.480823613509,48.69035376539358 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,17,True,P/P,0.0,0.5,5744.125,0.125,0.0,0.00023774057626724243,0.00023774057626724243,0.0036657596938312054,0.2870829105377197,0.000623806961812079,0.5967342020907382,0.012528423219919205,0.7349861860275269,29655.0,45953.0,64.46116399765015,6.235980749130249,90.48583698272705,14.88133692741394,7643.224441381263,48.427231765386196 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,18,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.24985149502754211,0.0005341017968021333,0.0,0.010709828697144985,0.6611655950546265,31045.0,57344.0,67.51829314231873,6.3801703453063965,94.03099393844604,15.157917022705078,9290.346630981201,60.24428462597861 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,19,True,P/P,0.0,0.75,4527.25,0.0,0.0,0.0003070533275604248,0.0003070533275604248,0.004181845113635063,0.2931738495826721,0.0006898391293361783,0.6365560311993935,0.012632078491151333,0.61952805519104,22955.0,36218.0,50.636481523513794,5.478410005569458,76.11529660224915,15.098262071609497,6933.442597294526,43.05049417582007 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,20,True,P/P,1.4901161193847656e-08,0.375,6711.5,0.625,0.0,0.0003246329724788666,0.0003246329724788666,0.003028714098036289,0.2523394525051117,0.0009607656393200159,0.5312549366441127,0.011415356770157814,0.6908265352249146,29414.0,53692.0,66.54918098449707,6.355381011962891,92.98513317108154,15.227828025817871,8745.308141322097,56.334857406725845 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,21,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.42058873176574707,0.000759697169996798,0.0,0.014996391721069813,0.7838053107261658,42896.0,57344.0,67.45322680473328,6.360701322555542,93.75775361061096,15.068852424621582,9306.348909448629,60.33158585428821 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,22,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.24330207705497742,0.0005154722603037953,0.0,0.010491483844816685,0.6533262729644775,34671.0,57344.0,67.3405556678772,6.425260305404663,93.64012289047241,14.979021072387695,9175.486146803476,59.451766496438914 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,23,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3639666438102722,0.000785493350122124,0.0,0.014399819076061249,0.7272769808769226,39032.0,57344.0,67.44907855987549,6.367808103561401,93.60742592811584,14.909992933273315,9301.364940821186,60.306476170997875 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,24,True,P/P,-3.725290298461914e-09,0.875,5936.125,0.125,0.0,0.000288182869553566,0.000288182869553566,0.004802369512617588,0.25187551975250244,0.0010903594084084034,0.5948213648911591,0.011529816314578056,0.7873067855834961,27313.0,47489.0,65.37698674201965,6.298056602478027,91.37925720214844,14.841277599334717,8136.341018017501,51.97587203565573 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,25,True,P/P,-3.725290298461914e-09,0.875,5395.25,0.0,0.0,0.0003207121044397354,0.0003207121044397354,0.004798966459929943,0.2643960118293762,0.0009967524092644453,0.625912572294836,0.011647357605397701,0.6367483139038086,26014.0,43162.0,56.07310223579407,6.219217300415039,82.3099935054779,15.129562139511108,7183.511299445129,45.19522047820603 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,26,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.284249871969223,0.0007087260019034147,0.0,0.01306111179292202,0.6667622327804565,35369.0,57344.0,67.3579409122467,6.395200252532959,93.74777007102966,15.06612753868103,9215.632416656272,59.718005820502235 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,27,True,P/P,1.4901161193847656e-08,0.375,6398.875,0.25,0.0,7.593631744384766e-05,7.593631744384766e-05,0.003754957113415003,0.3934386968612671,0.0007342100143432617,0.6641373311683457,0.015292856842279434,0.8349349498748779,35827.0,51191.0,65.6477906703949,6.3069281578063965,91.9380235671997,15.044588327407837,8386.89947716577,53.7040163707761 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,28,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.22138556838035583,0.0006103095947764814,0.0,0.010613763704895973,0.7080219388008118,30482.0,57344.0,67.78959560394287,6.37563681602478,93.94752788543701,14.877451181411743,9490.94775857754,61.705697221906554 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,29,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3695732355117798,0.0012376437662169337,0.0,0.017199477180838585,1.1279423236846924,43859.0,57344.0,67.48557734489441,6.526968240737915,94.05903053283691,15.170441150665283,9073.346014940178,58.83210283115573 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,30,True,P/P,0.0,1.0,2549.5,0.0,0.0,0.0,0.0,0.0,0.20721665024757385,0.000709145562723279,0.0,0.009911423549056053,0.3120119869709015,9779.0,20396.0,25.343219757080078,2.4734652042388916,47.90480422973633,15.267510890960693,8949.106287818653,53.48815905468905 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,31,True,P/P,0.0,1.0,5311.5,0.0,0.0,0.0,0.0,0.0,0.22483859956264496,0.0007269021589308977,0.0,0.010751388967037201,0.6687332987785339,22506.0,42492.0,54.75827193260193,6.23038125038147,80.86785864830017,15.21037220954895,7068.439997759417,44.40446181490209 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,32,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.270622193813324,0.0006680283695459366,0.0,0.012416140176355839,0.7035144567489624,33382.0,57344.0,67.88064384460449,6.3983473777771,94.15135025978088,15.011314630508423,9515.958611658387,61.916715963915046 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,33,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.24268731474876404,0.0008039348758757114,0.0,0.010907072573900223,0.7081736922264099,31473.0,57344.0,67.73456716537476,6.369226455688477,93.97797203063965,15.146235466003418,9329.540826109478,60.51598387150924 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,34,True,P/P,0.0,0.5,6628.75,0.5,0.0,4.2691826820373535e-06,4.2691826820373535e-06,0.0033109639771282673,0.24494783580303192,0.0005348019185476005,0.5515517938887827,0.011215727776288986,0.6709189414978027,28017.0,53030.0,66.72232055664062,6.353013277053833,92.72533822059631,14.973822832107544,8884.52621571104,57.31683983301582 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,35,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.30807313323020935,0.0007465713424608111,0.0,0.013447504490613937,0.6765857338905334,39265.0,57344.0,67.68034791946411,6.356398344039917,94.03023052215576,15.070174932479858,9294.985496478512,60.24866708864331 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,36,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.27146345376968384,0.0007877630414441228,0.0,0.012432310730218887,0.7376691102981567,32232.0,57344.0,67.34695601463318,6.336257696151733,93.47525548934937,14.998578071594238,9309.136895588663,60.32804008939835 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,37,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2860282063484192,0.0006167554529383779,0.0,0.012087345123291016,0.6726268529891968,38384.0,57344.0,67.42448329925537,6.36902928352356,93.96417570114136,15.188588380813599,9325.42207869322,60.486173701666694 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,38,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2782611846923828,0.000695729220751673,0.0,0.012096730060875416,0.7536394596099854,34590.0,57344.0,68.10038185119629,6.365238904953003,94.32513403892517,15.004334688186646,9579.810861709631,62.34383217449996 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,39,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.29391223192214966,0.0007897160248830914,0.0,0.012727180495858192,0.9134600162506104,35866.0,57344.0,67.41058945655823,6.4006102085113525,93.75015377998352,15.068745613098145,9245.643982957565,59.94213540854445 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,40,True,P/P,1.4901161193847656e-08,0.375,6819.875,0.5,0.0,0.0004830881953239441,0.0004830881953239441,0.003308583050966263,0.25014039874076843,0.0005372173618525267,0.5299652496677567,0.012198736891150475,0.87153160572052,31648.0,54559.0,66.93494820594788,6.3802266120910645,93.26677227020264,14.984029293060303,8914.66793939741,57.549137951648774 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,41,True,P/P,3.725290298461914e-09,0.125,6868.0,0.875,0.0,0.0003119371831417084,0.0003119371831417084,0.0029573081992566586,0.25413885712623596,0.0008641704334877431,0.5759921812033462,0.011810272932052612,0.8075380325317383,33113.0,54944.0,66.97346782684326,6.3576812744140625,93.18131494522095,14.973583221435547,8945.905789720715,57.80574746442397 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,42,True,P/P,0.0,0.0,7165.75,0.875,0.0,0.0,0.0,0.0,0.3768514394760132,0.0006598061881959438,0.0,0.014488877728581429,0.7221826314926147,39612.0,57326.0,67.3595335483551,6.3213050365448,94.62271332740784,14.997324705123901,9425.998263977033,61.162247404827674 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,43,True,P/P,3.725290298461914e-09,0.125,7052.0,0.875,0.0,3.729574382305145e-05,3.729574382305145e-05,0.002978515811264515,0.3057391941547394,0.0006695918855257332,0.5148845326503656,0.01412786915898323,0.7158371210098267,34723.0,56416.0,67.14346051216125,6.485588550567627,93.62628316879272,15.074557304382324,8980.454105249224,58.123624461271014 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,44,True,P/P,0.0,1.0,4179.5,0.0,0.0,0.0,0.0,0.0,0.2068227082490921,0.0008308816468343139,0.0,0.010345418006181717,0.5491979122161865,16073.0,33436.0,45.593576192855835,4.588877439498901,70.05687832832336,15.082818508148193,7624.88939700639,46.980004339261015 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,45,True,P/P,0.0,1.0,5467.125,0.0,0.0,0.0,0.0,0.0,0.19777163863182068,0.0008107761386781931,0.0,0.009436438791453838,0.7277637124061584,21563.0,43737.0,58.08453059196472,6.27031135559082,84.90361189842224,15.736234188079834,7257.745257757198,45.76329974769817 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,46,True,P/P,1.4901161193847656e-08,0.375,6648.25,0.625,0.0,0.00020582973957061768,0.00020582973957061768,0.003618717659264803,0.314630925655365,0.0009237479534931481,0.6064977208587643,0.013668986968696117,0.7044045925140381,35492.0,53186.0,66.36095762252808,6.325867176055908,92.73200869560242,15.137802362442017,8686.330379661544,55.92948853286966 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,47,True,P/P,3.725290298461914e-09,0.125,7126.25,0.875,0.0,0.00013535842299461365,0.00013535842299461365,0.0027693971060216427,0.42261332273483276,0.0007535129552707076,0.5918395527473749,0.01701831817626953,0.7283152341842651,42981.0,57010.0,67.75252676010132,6.369811296463013,94.2602710723877,15.044878482818604,9346.259622985497,60.641910378338565 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,48,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.47952818870544434,0.0006161052733659744,0.0,0.015877891331911087,0.6626149415969849,46377.0,57344.0,67.5626118183136,6.357044219970703,93.77983283996582,14.993746995925903,9362.21846040489,60.74140708074354 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,49,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.29638099670410156,0.000947827473282814,0.0,0.013257095590233803,0.6850106120109558,35343.0,57344.0,67.37736177444458,6.392014980316162,93.6281807422638,15.063201904296875,9223.710591444718,59.77137308213693 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,50,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2594350576400757,0.0007713806116953492,0.0,0.01235903799533844,0.7498743534088135,31943.0,57344.0,67.62307953834534,6.333682060241699,93.93626046180725,14.989206075668335,9332.352731252433,60.49500409427063 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,51,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2508096694946289,0.0009767418960109353,0.0,0.011743410490453243,0.7691302299499512,30754.0,57344.0,67.28418636322021,6.369722127914429,93.82317733764648,15.303053855895996,9230.259772907475,59.7933971475037 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,52,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.26666396856307983,0.0008693314157426357,0.0,0.01264298614114523,0.7149931192398071,33782.0,57344.0,67.71640610694885,6.371352434158325,93.83361625671387,14.990099906921387,9269.330195777808,60.0813482596744 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,53,True,P/P,0.0,1.0,5233.5,0.0,0.0,0.0,0.0,0.0,0.23909378051757812,0.0008377513149753213,0.0,0.011449085548520088,0.7537249326705933,21770.0,41868.0,52.469178915023804,6.257869720458984,78.8018491268158,15.135361433029175,6905.171236126931,43.302127417388206 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,54,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.28538092970848083,0.0008234899141825736,0.0,0.013129748404026031,0.7198138236999512,38231.0,57344.0,67.50475215911865,6.443543195724487,93.90998101234436,15.054537773132324,9203.433604169111,59.685776580657645 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,55,True,P/P,0.0,0.75,6677.75,0.25,0.0,-0.00026410818099975586,-0.00026410818099975586,0.005386143922805786,0.2663571238517761,0.0009096003486774862,0.521946248884377,0.012554297223687172,0.7343463897705078,28903.0,53422.0,65.9813277721405,6.370711326599121,92.42027711868286,15.168467283248901,8618.754613021312,55.395723324670335 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,56,True,P/P,0.0,0.75,5243.125,0.25,0.0,0.000331193208694458,0.000331193208694458,0.0037555256858468056,0.18771588802337646,0.00043148433906026185,0.5189496276699032,0.009228058159351349,0.71006178855896,18917.0,41945.0,63.95794081687927,5.428218841552734,89.30646991729736,15.010161638259888,8047.684730199778,50.808059036030244 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,57,True,P/P,0.0,1.0,3335.75,0.0,0.0,0.0,0.0,0.0,0.24421292543411255,0.0006768715684302151,0.0,0.012128313072025776,0.4239920377731323,15431.0,26686.0,32.396788358688354,3.2927286624908447,55.692744731903076,15.113965272903442,8657.349169927984,52.49058081852197 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,58,True,P/P,0.0,1.0,5181.625,0.0,0.0,0.0,0.0,0.0,0.2310042679309845,0.0008368862909264863,0.0,0.010792052373290062,0.5970793962478638,21315.0,41453.0,56.80596470832825,5.50947904586792,82.09441423416138,15.05582308769226,7861.800115470691,49.40792635743914 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,59,True,P/P,0.0,1.0,5421.875,0.0,0.0,0.0,0.0,0.0,0.21457432210445404,0.0009814996737986803,0.0,0.010522632859647274,0.7265615463256836,22194.0,43375.0,56.95116090774536,6.271579027175903,83.14483404159546,15.092703580856323,7230.100430063068,45.553877405776426 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,60,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2874848544597626,0.0008291081758216023,0.0,0.013705053366720676,0.9079194664955139,36965.0,57344.0,67.52107119560242,6.44205379486084,93.89828491210938,15.134571075439453,9239.47325673072,59.94606674666258 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,61,True,P/P,3.725290298461914e-09,0.125,7130.5,0.875,0.0,0.00015364587306976318,0.00015364587306976318,0.0028858319856226444,0.322539746761322,0.0006234465981833637,0.5457902504271389,0.0141767468303442,0.7929098010063171,38669.0,57044.0,67.48453283309937,6.356516361236572,94.40430998802185,15.311083793640137,9340.478188637007,60.58160343609727 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,62,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3531469702720642,0.0007891095010563731,0.0,0.015959134325385094,0.755255937576294,39392.0,57344.0,67.7796242237091,6.358715534210205,94.46851587295532,15.388328790664673,9512.605417068153,61.84124508371341 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,63,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.34104862809181213,0.001190597889944911,0.0,0.016013741493225098,0.742705225944519,41928.0,57344.0,67.47012829780579,6.372115850448608,93.99813866615295,15.13271713256836,9296.272028548108,60.27756805806605 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,64,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3318078815937042,0.0011185111943632364,0.0,0.015093645080924034,0.6923420429229736,41591.0,57344.0,67.63527059555054,6.376891136169434,94.16206693649292,15.214966297149658,9239.556117490467,59.87098959726776 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,65,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.31942218542099,0.000817720138002187,0.0,0.014478446915745735,0.7410677671432495,38154.0,57344.0,67.45664286613464,6.340852975845337,93.7758355140686,15.091269254684448,9364.131274775244,60.73414092213436 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,66,True,P/P,0.0,1.0,5184.25,0.0,0.0,0.0,0.0,0.0,0.2977432608604431,0.0010779425501823425,0.0,0.014307957142591476,0.7934126853942871,26329.0,41474.0,56.170511960983276,6.304281234741211,82.49020171165466,15.16488265991211,6917.932369192023,43.44208886189577 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,67,True,P/P,1.4901161193847656e-08,0.375,6906.125,0.625,0.0,-1.7605721950531006e-05,-1.7605721950531006e-05,0.003012950997799635,0.20642954111099243,0.0005201281164772809,0.505730748853425,0.009977955371141434,0.8390500545501709,27317.0,55249.0,66.8147840499878,6.344203472137451,93.36304473876953,15.265349388122559,9037.267023738159,58.38128806050065 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,68,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.40800905227661133,0.0009083639015443623,0.0,0.015391799621284008,0.8063873648643494,39677.0,57344.0,67.70261430740356,6.374029636383057,94.3294529914856,15.257869005203247,9407.49045459393,61.08610782315811 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,69,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3426581621170044,0.0007669199840165675,0.0,0.01561337523162365,0.8409221172332764,40966.0,57344.0,67.73448967933655,6.368502855300903,94.23263359069824,15.148190021514893,9290.456176465175,60.22958323389484 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,70,True,P/P,0.0,1.0,3451.5,0.0,0.0,0.0,0.0,0.0,0.22270476818084717,0.0003573326685000211,0.0,0.010563607327640057,0.38571321964263916,13941.0,27612.0,36.543121099472046,3.2321152687072754,59.970720052719116,15.195897102355957,9077.115424428282,55.23636442884131 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,71,True,P/P,0.0,0.25,6849.75,0.75,0.0,0.00013498961925506592,0.00013498961925506592,0.0038289539515972137,0.30832812190055847,0.0012091277167201042,0.6333301687370676,0.014399094507098198,0.7864307165145874,36712.0,54798.0,66.84653806686401,6.351085424423218,93.21266031265259,15.115710735321045,8915.396384112533,57.5451130522199 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,72,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.24538514018058777,0.0007454284350387752,0.0,0.011631857603788376,0.7179243564605713,32049.0,57344.0,67.89812135696411,6.35034441947937,94.44485402107239,15.026439189910889,9427.782840969976,61.20848961016192 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,73,True,P/P,-1.4901161193847656e-08,0.625,6320.0,0.375,0.0,0.00024543702602386475,0.00024543702602386475,0.00435059517621994,0.30234283208847046,0.0007018998730927706,0.5793996111468537,0.013330023735761642,0.7841492891311646,32889.0,50560.0,65.85671305656433,6.281113386154175,92.14855670928955,15.059760570526123,8402.240008054876,53.85327575990681 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,74,True,P/P,-3.725290298461914e-09,0.875,4209.0,0.0,0.0,-9.070895612239838e-05,-9.070895612239838e-05,0.004914691671729088,0.22938336431980133,0.0009615623275749385,0.7480114960832521,0.011085063219070435,0.642704963684082,17099.0,33672.0,57.405741453170776,3.9923715591430664,81.42496538162231,15.105658054351807,8934.295046695432,55.26404855721005 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,75,True,P/P,0.0,0.25,6717.375,0.375,0.0,0.00027301907539367676,0.00027301907539367676,0.003536336123943329,0.29691582918167114,0.0009919428266584873,0.6048217515977958,0.013966390863060951,0.9174409508705139,33047.0,53739.0,67.52310752868652,6.301604509353638,93.81405210494995,15.10212779045105,9359.421313553998,60.68632712076373 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,76,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.28536003828048706,0.0010829069651663303,0.0,0.014203024096786976,0.7892655730247498,36181.0,57344.0,67.58572125434875,6.306264162063599,93.90244245529175,15.071896076202393,9470.896061795922,61.47268475763268 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,77,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2768442928791046,0.0008173760725185275,0.0,0.012953069992363453,0.8070613145828247,37758.0,57344.0,67.77657151222229,6.313199758529663,94.1372721195221,15.038735628128052,9434.212207099263,61.212670779231296 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,78,True,P/P,1.4901161193847656e-08,0.375,6907.875,0.625,0.0,4.710257053375244e-05,4.710257053375244e-05,0.00345963379368186,0.4023745656013489,0.0008983981679193676,0.6182430753185115,0.016065331175923347,0.6866913437843323,39257.0,55263.0,66.95761704444885,6.293524265289307,93.4561288356781,15.171188592910767,9157.797593029287,59.206008653750295 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,79,True,P/P,0.0,1.0,5177.0,0.0,0.0,0.0,0.0,0.0,0.2582058906555176,0.000961288926191628,0.0,0.012080110609531403,0.7023210525512695,24337.0,41416.0,62.360743284225464,6.189118146896362,88.55841255187988,15.11869192123413,6906.81262205094,43.33693299438297 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,80,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.27567991614341736,0.0006376994424499571,0.0,0.013767771422863007,0.7610621452331543,34702.0,57344.0,67.7992639541626,6.317007064819336,94.24509692192078,15.112595319747925,9399.697099385892,60.96481472844181 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,81,True,P/P,0.0,1.0,3255.625,0.0,0.0,0.0,0.0,0.0,0.20142816007137299,0.0008714735740795732,0.0,0.00992856826633215,0.355027437210083,13286.0,26045.0,30.744930505752563,3.2266063690185547,54.27278137207031,15.229177713394165,8593.600617886697,52.01685037529812 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,82,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3487071692943573,0.0010963281383737922,0.0,0.016134493052959442,0.7537202835083008,40329.0,57344.0,67.49972295761108,6.347830772399902,94.03335332870483,15.157233953475952,9315.455241823478,60.3885603349791 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,83,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.4032915234565735,0.001041030278429389,0.0,0.015653368085622787,0.8212316036224365,41023.0,57344.0,67.78195571899414,6.405786514282227,94.27942323684692,15.172584772109985,9285.850097761542,60.23258434935497 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,84,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.5525314211845398,0.000841391971334815,0.0,0.018820416182279587,0.7246725559234619,51533.0,57344.0,67.52990913391113,6.367157220840454,93.97012305259705,15.087298154830933,9347.769370289576,60.64662862387802 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,85,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.32111823558807373,0.0008562636794522405,0.0,0.014897716231644154,1.1041171550750732,37946.0,57344.0,67.629967212677,6.333343744277954,93.9117419719696,15.01301097869873,9460.640424939087,61.43018279714903 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,86,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.289934366941452,0.0007786848582327366,0.0,0.014075987972319126,0.7617355585098267,34927.0,57344.0,67.6239824295044,6.329829931259155,94.09118723869324,15.05236268043518,9296.541507799227,60.22996546750563 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,87,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.32439881563186646,0.0010151541791856289,0.0,0.01591780036687851,0.8044065237045288,39168.0,57344.0,67.56782031059265,6.456043004989624,93.99975967407227,15.061241388320923,9238.87835011852,59.957533159626486 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,88,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.22830970585346222,0.0005531225469894707,0.0,0.011596232652664185,0.8259167671203613,31588.0,57344.0,67.85029101371765,6.373074293136597,94.26994061470032,15.085735321044922,9392.487340235473,60.98037751126591 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,89,True,P/P,-3.725290298461914e-09,0.875,6284.75,0.25,0.0,0.00022034719586372375,0.00022034719586372375,0.007445703260600567,0.34967494010925293,0.001077240682207048,0.6236411103969806,0.01632785052061081,0.768391489982605,34120.0,50278.0,65.53394746780396,6.3353705406188965,91.96562623977661,15.057074069976807,8182.328637147139,52.29276347426315 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,90,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.24793218076229095,0.0008368950802832842,0.0,0.011843780055642128,0.9270305633544922,32491.0,57344.0,67.5321147441864,6.359710693359375,93.86002683639526,15.07808804512024,9352.915803157042,60.675880314248495 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,91,True,P/P,0.0,0.5,6485.5,0.25,0.0,0.0002458207309246063,0.0002458207309246063,0.003156268037855625,0.45404618978500366,0.0007038084440864623,0.6737294442716021,0.015728330239653587,0.6842485070228577,42019.0,51884.0,65.98905062675476,6.336275577545166,92.34740161895752,15.068203687667847,8527.996623432944,54.7612074184069 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,92,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3222174644470215,0.0007807333022356033,0.0,0.01512668002396822,0.8013651371002197,40439.0,57344.0,67.52010416984558,6.418318510055542,94.03331065177917,15.12215781211853,9232.706519840453,59.8684687961251 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,93,True,P/P,-3.725290298461914e-09,0.875,3809.125,0.125,0.0,0.0002673100680112839,0.0002673100680112839,0.00449429452419281,0.20175130665302277,0.0005862729740329087,0.630310292048599,0.010278457775712013,0.5403860807418823,15341.0,30473.0,60.92590284347534,4.01291298866272,85.12385177612305,15.196085453033447,7996.14734302801,49.35822604085398 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,94,True,P/P,1.4901161193847656e-08,0.375,7073.5,0.625,0.0,0.00010643899440765381,0.00010643899440765381,0.0029509635642170906,0.20436625182628632,0.0006970911053940654,0.47853547658419515,0.009630871005356312,0.9365644454956055,27989.0,56588.0,67.21600699424744,6.381656169891357,93.82421565055847,15.305514335632324,9102.918854973526,58.88087864747597 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,95,True,P/P,3.725290298461914e-09,0.125,7157.75,0.875,0.0,0.00011398456990718842,0.00011398456990718842,0.00274492590688169,0.22009414434432983,0.00035898119676858187,0.47746923829398313,0.010730587877333164,0.775492787361145,31519.0,57262.0,67.24652791023254,6.37655234336853,93.88536214828491,15.288172721862793,9246.616695763058,59.91249860326271 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,96,True,P/P,0.0,1.0,5302.625,0.0,0.0,0.0,0.0,0.0,0.3616517186164856,0.0010625217109918594,0.0,0.015830131247639656,0.9405617117881775,29110.0,42421.0,53.95995593070984,6.334820985794067,80.3621711730957,15.137881517410278,6979.458244479836,43.87306495291692 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,97,True,P/P,0.0,0.25,6841.25,0.625,0.0,6.068870425224304e-05,6.068870425224304e-05,0.003941177856177092,0.3091956675052643,0.0008720500627532601,0.599975474174233,0.014346808195114136,0.7740021347999573,34134.0,54730.0,66.60762047767639,6.390402317047119,93.18908071517944,15.058161497116089,8837.694836735485,57.00664937199882 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,98,True,P/P,0.0,0.5,6676.0,0.5,0.0,0.000473756343126297,0.000473756343126297,0.0033199626486748457,0.1819540560245514,0.0008638116996735334,0.5889159731164741,0.00939169991761446,0.8200641870498657,23144.0,53408.0,66.35380601882935,6.359963417053223,92.71371531486511,14.873374938964844,8675.044479292574,55.82637257947092 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,99,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.25524458289146423,0.0007820262690074742,0.0,0.012559186667203903,0.9112148284912109,31755.0,57344.0,67.75846457481384,6.343771696090698,93.99711084365845,15.032795906066895,9363.941569894354,60.737052837089465 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,100,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.21459826827049255,0.0005396216874942183,0.0,0.01118271891027689,0.7919857501983643,29023.0,57344.0,67.38548970222473,6.364102602005005,93.77971887588501,15.03586745262146,9299.426136055086,60.289791937350884 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,101,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3231019973754883,0.0008840380469337106,0.0,0.015511088073253632,0.8265482187271118,37437.0,57344.0,67.53284502029419,6.410966157913208,94.14853549003601,15.202080726623535,9305.71540023192,60.39334372574686 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,102,True,P/P,0.0,0.0,6704.75,0.5,0.0,0.0,0.0,0.0,0.15616990625858307,0.0005610703956335783,0.0,0.008077564649283886,0.745896577835083,21511.0,53638.0,66.41035842895508,6.318365573883057,92.99974870681763,15.219489812850952,8784.936212151397,56.59112039225815 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,103,True,P/P,0.0,0.25,7044.75,0.75,0.0,3.750622272491455e-05,3.750622272491455e-05,0.0036741215735673904,0.22662468254566193,0.0006427011685445905,1.2974886533942986,0.011584131047129631,0.8680052161216736,31310.0,56358.0,67.41820883750916,6.352008581161499,93.86368584632874,15.158912181854248,9380.84729682628,60.87674688828737 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,104,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2581730782985687,0.0008885071729309857,0.0,0.012260248884558678,0.716458797454834,31777.0,57344.0,67.9188084602356,6.328139781951904,94.42959809303284,15.201119184494019,9387.820814126702,60.891940061430205 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,105,True,P/P,0.0,0.75,6231.625,0.25,0.0,0.00018107891082763672,0.00018107891082763672,0.006084096617996693,0.26297760009765625,0.0009297077194787562,0.7220760181171657,0.013475578278303146,0.8152446150779724,30230.0,49853.0,65.60352110862732,6.345219850540161,92.21387934684753,15.295037984848022,8215.192837603761,52.545168176605 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,106,True,P/P,3.725290298461914e-09,0.125,7119.5,0.875,0.0,0.00024863891303539276,0.00024863891303539276,0.003051809500902891,0.3428218960762024,0.0009720748639665544,0.5868321772624211,0.01460409164428711,0.7211999893188477,39823.0,56956.0,67.5756254196167,6.3554160594940186,94.30917835235596,15.128238201141357,9254.628141536003,59.947951586349596 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,107,True,P/P,0.0,1.0,4053.75,0.0,0.0,0.0,0.0,0.0,0.18050113320350647,0.0009462847374379635,0.0,0.009583141654729843,0.6617510914802551,15023.0,32430.0,42.88018274307251,4.004113674163818,66.78155517578125,15.027718782424927,8432.684251894107,51.784800014998936 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,108,True,P/P,0.0,1.0,2943.625,0.0,0.0,0.0,0.0,0.0,0.12977071106433868,0.0005963249132037163,0.0,0.006788601167500019,0.35369646549224854,8585.0,23549.0,30.756223917007446,3.1332974433898926,53.8142511844635,15.030449628829956,7962.532227259259,47.95946219278491 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,109,True,P/P,1.4901161193847656e-08,0.375,6727.625,0.625,0.0,0.00047989189624786377,0.00047989189624786377,0.004445475526154041,0.2655128240585327,0.0009601892670616508,0.6786154041704476,0.012863274663686752,0.9146977066993713,29432.0,53821.0,66.6486144065857,6.389866590499878,93.18601536750793,15.336591482162476,8744.384434875823,56.37552647138486 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,110,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.22396065294742584,0.0005007424624636769,0.0,0.011492710560560226,0.8535133004188538,29744.0,57344.0,67.79088234901428,6.389566659927368,94.32488322257996,15.270131349563599,9320.33496498645,60.47173145220219 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,111,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3294116258621216,0.0009693795000202954,0.0,0.01510331779718399,0.7086734771728516,37289.0,57344.0,67.65007615089417,6.417985916137695,94.37420177459717,15.279042720794678,9358.778743306617,60.782225810593296 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,112,True,P/P,0.0,1.0,3824.625,0.0,0.0,0.0,0.0,0.0,0.16740870475769043,0.0006907394854351878,0.0,0.008181029930710793,0.5028310418128967,14014.0,30597.0,38.695157051086426,3.8263845443725586,62.64647150039673,15.135695219039917,8424.66643842618,51.53686315828813 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,113,True,P/P,1.4901161193847656e-08,0.375,7022.625,0.625,0.0,4.380941390991211e-05,4.380941390991211e-05,0.004547429271042347,0.24836452305316925,0.0011233146069571376,0.6007962734690347,0.012330872938036919,0.7677915692329407,29958.0,56181.0,66.79024291038513,6.401947021484375,93.38442349433899,15.237576246261597,9002.729715408916,58.179181615721355 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,114,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2899550199508667,0.0008754355367273092,0.0,0.013584412634372711,0.838972270488739,34625.0,57344.0,67.53426337242126,6.391494274139404,94.10552191734314,15.186551332473755,9341.159998668036,60.62131020074859 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,115,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3727809190750122,0.0009608711116015911,0.0,0.01832321099936962,0.8680381178855896,44823.0,57344.0,67.8729989528656,6.396614074707031,94.66840243339539,15.176841020584106,9324.121886320794,60.503519778931775 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,116,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2120274007320404,0.0007585033308714628,0.0,0.011088858358561993,0.9034280776977539,26632.0,57344.0,67.5569269657135,6.382016897201538,94.16959118843079,15.221373796463013,9363.599825259678,60.77418667994555 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,117,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.26867133378982544,0.0008386123226955533,0.0,0.012352518737316132,0.7651790976524353,31529.0,57344.0,67.49231600761414,6.358332395553589,93.95627737045288,15.119542121887207,9365.037978581928,60.75555712089617 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,118,True,P/P,-1.4901161193847656e-08,0.625,6473.375,0.375,0.0,9.848177433013916e-05,9.848177433013916e-05,0.0037386524491012096,0.18658193945884705,0.0007126008276827633,0.5406107024137897,0.00965781882405281,0.806789755821228,24270.0,51787.0,65.89793968200684,6.412827730178833,92.288747549057,15.112858295440674,8355.86437487944,53.58953390881044 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,119,True,P/P,0.0,1.0,3175.625,0.0,0.0,0.0,0.0,0.0,0.15957479178905487,0.0008452131878584623,0.0,0.00831315666437149,0.4673055410385132,10883.0,25405.0,34.21227788925171,3.2686314582824707,57.39112305641174,14.98010778427124,8411.48361176403,50.89627663260414 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,120,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.19870617985725403,0.0008531083585694432,0.0,0.010987568646669388,0.9920612573623657,27750.0,57344.0,67.35731172561646,6.444689512252808,93.73879480361938,15.02696704864502,9142.49833735049,59.23499158531836 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,121,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2698844373226166,0.0008366795955225825,0.0,0.01390913873910904,0.7911263704299927,33319.0,57344.0,67.63070464134216,6.393565654754639,94.00140118598938,15.00385570526123,9217.104504000712,59.7244870454866 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,122,True,P/P,0.0,1.0,3401.0,0.0,0.0,0.0,0.0,0.0,0.1408957988023758,0.0006220638751983643,0.0,0.007766854949295521,0.43279772996902466,10426.0,27208.0,45.714789628982544,3.227693796157837,68.98838329315186,15.232011318206787,9002.738171117528,55.005750314176346 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,123,True,P/P,-1.4901161193847656e-08,0.625,6797.0,0.375,0.0,-0.00020376592874526978,-0.00020376592874526978,0.004046706482768059,0.2049463987350464,0.0006638794438913465,0.5639818065123509,0.010217580944299698,0.7826426029205322,27920.0,54376.0,66.32177448272705,6.3992133140563965,92.87898397445679,15.191540479660034,8761.464086220309,56.44788420988028 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,124,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.22444286942481995,0.0007461743662133813,0.0,0.011409621685743332,0.8394088745117188,27260.0,57344.0,67.76019549369812,6.370972394943237,94.23970603942871,15.095263242721558,9317.202939319426,60.428741421456344 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,125,True,P/P,0.0,1.0,4684.125,0.0,0.0,0.0,0.0,0.0,0.22589600086212158,0.0008215614943765104,0.0,0.010807998478412628,0.6057238578796387,19761.0,37473.0,60.92041540145874,4.757734060287476,85.68621802330017,15.111179828643799,8318.131379850107,51.91282743122392 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,126,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2420257180929184,0.0011073609348386526,0.0,0.012395400553941727,0.964026927947998,30410.0,57344.0,67.46127128601074,6.364480972290039,93.91958808898926,15.101147174835205,9305.605757257988,60.33603024631874 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,127,True,P/P,0.0,0.0,6329.875,0.375,0.0,0.0,0.0,0.0,0.3864273130893707,0.0009531343821436167,0.0,0.016605578362941742,0.8147666454315186,36655.0,50639.0,65.87735247612,6.334499359130859,92.6009681224823,15.194900751113892,8315.470576462092,53.27223431938662 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,128,True,P/P,-3.725290298461914e-09,0.875,5847.0,0.125,0.0,0.00028762221336364746,0.00028762221336364746,0.0068326652981340885,0.31820106506347656,0.0007392630213871598,0.671833037019647,0.014878787100315094,0.8907882571220398,29980.0,46776.0,64.74574065208435,6.30497407913208,90.95399260520935,15.12839651107788,7672.117146946558,48.763508690451445 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,129,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.24428725242614746,0.0007687550969421864,0.0,0.012878907844424248,0.961449146270752,32656.0,57344.0,67.41790127754211,6.351522207260132,93.7916522026062,15.119333982467651,9305.98257072852,60.32406510227149 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,130,True,P/P,0.0,1.0,3718.125,0.0,0.0,0.0,0.0,0.0,0.16459180414676666,0.0004213511710986495,0.0,0.008301196619868279,0.4871922433376312,12929.0,29745.0,41.61287879943848,3.8084099292755127,65.74057269096375,15.315964698791504,8248.647612989891,50.41117617994279 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,131,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3769303560256958,0.0007942329393699765,0.0,0.01626121625304222,0.8060085773468018,39872.0,57344.0,67.53471183776855,6.46666145324707,93.98153233528137,15.046279191970825,9171.1456883963,59.47739840596354 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,132,True,P/P,0.0,1.0,4948.0,0.0,0.0,0.0,0.0,0.0,0.19264428317546844,0.000975073198787868,0.0,0.01040636096149683,0.8829958438873291,19475.0,39584.0,49.6487979888916,6.1610496044158936,75.71030807495117,15.062144994735718,6773.987994463669,42.37753221989919 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,133,True,P/P,3.725290298461914e-09,0.125,7103.75,0.875,0.0,-0.00016242451965808868,-0.00016242451965808868,0.0027353153564035892,0.15642967820167542,0.0007390099344775081,0.4206542614512282,0.008199939504265785,0.8248308897018433,23763.0,56830.0,67.60750389099121,6.390092849731445,94.1050636768341,15.106611251831055,9121.35103567051,59.03015744429538 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,134,True,P/P,0.0,1.0,3821.375,0.0,0.0,0.0,0.0,0.0,0.19853805005550385,0.0009581922786310315,0.0,0.010553708299994469,0.6697672605514526,15703.0,30571.0,41.505096435546875,4.614266872406006,66.87284421920776,15.829049348831177,6972.69413607747,42.71086548691335 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,135,True,P/P,-1.4901161193847656e-08,0.625,5290.875,0.375,0.0,0.00030531734228134155,0.00030531734228134155,0.0038873597513884306,0.20316621661186218,0.0007267269538715482,0.6150346219776832,0.01011617761105299,0.5982325077056885,20172.0,42327.0,64.10828304290771,4.821778059005737,88.84078192710876,15.071986198425293,9100.311873761275,57.51723656365943 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,136,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2335154116153717,0.0005321192438714206,0.0,0.011660189367830753,0.9391834735870361,30861.0,57344.0,67.47000670433044,6.362977981567383,94.06638312339783,15.305525541305542,9291.297285613797,60.23092593169597 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,137,True,P/P,0.0,1.0,4715.75,0.0,0.0,0.0,0.0,0.0,0.14842967689037323,0.000646679662168026,0.0,0.007632039487361908,0.7413687705993652,15508.0,37726.0,52.297096252441406,5.499350070953369,78.29396486282349,15.307559728622437,7162.710235517878,44.57509958090757 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,138,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.30734947323799133,0.0010758552234619856,0.0,0.013788452371954918,0.950189471244812,34222.0,57344.0,67.50150656700134,6.46800684928894,94.04072046279907,15.170518159866333,9159.082612444405,59.39004975422685 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,139,True,P/P,0.0,1.0,5772.5,0.0,0.0,0.0,0.0,0.0,0.1854669153690338,0.0006992141716182232,0.0,0.009977208450436592,0.9331016540527344,22507.0,46180.0,60.51004076004028,6.312273263931274,87.04251146316528,14.95940899848938,7634.124408456202,48.406042089338385 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,140,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2495022416114807,0.0007222970016300678,0.0,0.012976430356502533,0.7987363338470459,32104.0,57344.0,67.34665822982788,6.409116268157959,93.82383418083191,15.097327470779419,9196.429805937132,59.59458847984681 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,141,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.23605415225028992,0.0008368485723622143,0.0,0.013072646223008633,0.9113425016403198,30996.0,57344.0,67.97559118270874,6.388979911804199,94.62856340408325,15.128662586212158,9577.613232088323,62.35812896254111 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,142,True,P/P,0.0,1.0,4240.625,0.0,0.0,0.0,0.0,0.0,0.19680729508399963,0.0009119964670389891,0.0,0.010004622861742973,0.5930615663528442,16143.0,33925.0,43.94130277633667,4.562709093093872,68.55502605438232,15.097142219543457,7819.760139265238,48.217747261311665 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,143,True,P/P,0.0,1.0,5477.0,0.0,0.0,0.0,0.0,0.0,0.29717138409614563,0.0008487887680530548,0.0,0.014442061074078083,0.8446353077888489,27295.0,43816.0,59.54244256019592,6.243229150772095,85.74082636833191,15.023627042770386,7356.896750602015,46.39727431276039 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,144,True,P/P,0.0,1.0,4114.75,0.0,0.0,0.0,0.0,0.0,0.2036965787410736,0.0008262950577773154,0.0,0.010516749694943428,0.6002041101455688,15712.0,32918.0,42.91950178146362,4.737855672836304,67.6206681728363,14.990589618682861,7537.227150021851,46.49729933988225 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,145,True,P/P,0.0,0.75,6292.125,0.25,0.0,0.00017589330673217773,0.00017589330673217773,0.005394375883042812,0.24921990931034088,0.0008468691376037896,0.6882064397643071,0.012495797127485275,0.9255133271217346,27914.0,50337.0,65.51732969284058,6.3486857414245605,91.79705572128296,14.9741792678833,8178.159319688835,52.25986943591566 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,146,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2891666889190674,0.0011333490256220102,0.0,0.01507475133985281,1.0332951545715332,35714.0,57344.0,67.69168281555176,6.402607440948486,94.2870762348175,15.250696659088135,9245.125818038103,59.934686250411445 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,147,True,P/P,-1.4901161193847656e-08,0.625,6358.875,0.375,0.0,0.00041356682777404785,0.00041356682777404785,0.004172147251665592,0.20559647679328918,0.0005689900135621428,0.6636702557162895,0.010723788291215897,0.8848750591278076,24425.0,50871.0,66.29557609558105,6.331425428390503,92.81145811080933,15.213338375091553,8623.528929371716,55.44128285703759 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,148,True,P/P,0.0,1.0,4304.5,0.0,0.0,0.0,0.0,0.0,0.16409261524677277,0.0005355008761398494,0.0,0.008349365554749966,0.5875869989395142,14014.0,34436.0,47.57465410232544,4.725721597671509,72.59109568595886,15.123056888580322,7595.680844433473,46.92448734529482 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,149,True,P/P,0.0,0.25,6891.375,0.75,0.0,0.0002533271908760071,0.0002533271908760071,0.0037943022325634956,0.21072345972061157,0.0006141605554148555,0.5821516949244386,0.011038634926080704,0.8056501746177673,26802.0,55131.0,66.89948320388794,6.340243101119995,93.6099865436554,15.158060550689697,9017.353635315481,58.247305933880064 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,150,True,P/P,0.0,1.0,2047.5,0.0,0.0,0.0,0.0,0.0,0.2602422833442688,0.0013120456133037806,0.0,0.01269424706697464,0.2950223684310913,9649.0,16380.0,26.685163259506226,5.675842046737671,52.45035243034363,15.149111986160278,3075.6202767379104,18.230240653307728 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,151,True,P/P,0.0,0.75,6120.375,0.25,0.0,0.00012321770191192627,0.00012321770191192627,0.004093406721949577,0.15295477211475372,0.0007919669151306152,0.5557343182120928,0.008738866075873375,0.9508228302001953,20396.0,48963.0,65.32458329200745,6.214555025100708,91.31304979324341,15.067192077636719,8124.17422809154,51.8873386829652 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,152,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2728036046028137,0.0007697996916249394,0.0,0.014116967096924782,0.9532368183135986,32830.0,57344.0,67.61130046844482,6.443777561187744,94.1735589504242,15.169871807098389,9149.567362542131,59.290911065000216 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,153,True,P/P,-3.725290298461914e-09,0.875,6037.25,0.125,0.0,-5.689449608325958e-05,-5.689449608325958e-05,0.004953292198479176,0.20750746130943298,0.0009269799338653684,0.6236529616133345,0.010864203795790672,0.7915712594985962,24327.0,48298.0,65.09736490249634,6.411322593688965,91.76114988327026,15.339327573776245,7822.775828307237,49.86603321668947 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,154,True,P/P,0.0,1.0,5065.5,0.0,0.0,0.0,0.0,0.0,0.19748957455158234,0.00040615227771922946,0.0,0.010148395784199238,0.9533203840255737,18652.0,40524.0,54.850642919540405,6.267246961593628,81.39117956161499,15.312497854232788,6725.774852137331,42.09835807020983 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,155,True,P/P,0.0,1.0,4628.5,0.0,0.0,0.0,0.0,0.0,0.13087503612041473,0.0004619663523044437,0.0,0.00703656580299139,0.6616048812866211,13695.0,37028.0,56.03143858909607,5.49026083946228,81.67119550704956,15.287111520767212,7027.384340359694,43.68028328964508 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,156,True,P/P,0.0,1.0,5501.75,0.0,0.0,0.0,0.0,0.0,0.21596330404281616,0.0005705939838662744,0.0,0.010894479230046272,0.8228352069854736,24039.0,44014.0,58.97835922241211,6.398430109024048,85.63122415542603,15.294768810272217,7203.8155665755985,45.52058530417774 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,157,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.16358080506324768,0.0006987003143876791,0.0,0.008753277361392975,0.8698636293411255,22313.0,57344.0,67.38585567474365,6.364393472671509,93.80694365501404,15.196990489959717,9279.92047736058,60.149991665670704 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,158,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.16927310824394226,0.00026272900868207216,0.0,0.00899159349501133,0.768256664276123,22556.0,57344.0,67.6019675731659,6.347429275512695,94.10388469696045,15.151665687561035,9281.726341042342,60.14219358781933 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,159,True,P/P,1.4901161193847656e-08,0.375,2245.875,0.0,0.0,5.7227909564971924e-05,5.7227909564971924e-05,0.0032173278741538525,0.15500381588935852,0.00014810706488788128,1.1751806810129763,0.008441154845058918,0.3328143358230591,8089.0,17967.0,21.768691539764404,2.4849560260772705,44.237343549728394,15.047176837921143,7939.356659739811,47.16764685484971 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,160,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2860425114631653,0.0012022587470710278,0.0,0.014927936717867851,1.1796411275863647,34541.0,57344.0,67.52543377876282,6.347728729248047,93.86641454696655,14.963432788848877,9364.178103285725,60.74272946925774 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,161,True,P/P,0.0,1.0,4696.625,0.0,0.0,0.0,0.0,0.0,0.18663722276687622,0.00041961483657360077,0.0,0.009219438768923283,0.7830262184143066,17936.0,37573.0,50.27669978141785,5.3662168979644775,75.5648078918457,15.012318134307861,7280.847164868487,45.28232419253379 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,162,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.20692777633666992,0.0009613584261387587,0.0,0.011188654229044914,0.9793136119842529,26677.0,57344.0,68.13257265090942,6.374447584152222,94.64094758033752,15.098812818527222,9518.835535265933,61.908064625828814 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,163,True,P/P,0.0,0.5,6531.75,0.5,0.0,0.0004509352147579193,0.0004509352147579193,0.004052438773214817,0.2205243855714798,0.0008835654007270932,0.595391789359994,0.011234045960009098,0.8363525867462158,28554.0,52254.0,66.25490951538086,6.3563008308410645,92.49207353591919,14.94364309310913,8564.047072516369,55.05520775039848 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,164,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.21983692049980164,0.0008145634201355278,0.0,0.01136244647204876,0.8669316172599792,28566.0,57344.0,67.38202738761902,6.3707075119018555,94.91982913017273,16.15557622909546,9274.759316031357,60.117564096315114 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,165,True,P/P,0.0,1.0,5717.25,0.0,0.0,0.0,0.0,0.0,0.19667577743530273,0.0008433894254267216,0.0,0.009959597140550613,0.8910927772521973,23100.0,45738.0,58.59423542022705,6.323958873748779,85.03307962417603,15.10973048210144,7475.184153264588,47.273119743430634 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,166,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2859199047088623,0.0008461159304715693,0.0,0.01465662568807602,1.1917808055877686,36026.0,57344.0,67.39150381088257,6.375938177108765,93.71526980400085,14.978984594345093,9240.929590934164,59.87988950687532 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,167,True,P/P,0.0,1.0,4335.75,0.0,0.0,0.0,0.0,0.0,0.18341320753097534,0.0006446861079894006,0.0,0.009684107266366482,0.8891302347183228,16779.0,34686.0,43.16801953315735,5.332619905471802,68.75280785560608,15.034950017929077,6829.594114049937,42.185139803873575 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,168,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.28637540340423584,0.000846132286824286,0.0,0.014090701937675476,0.8165465593338013,32614.0,57344.0,67.40119886398315,6.355722427368164,93.78769111633301,15.001725196838379,9298.940711147165,60.27019062829377 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,169,True,P/P,-3.725290298461914e-09,0.875,6109.625,0.125,0.0,0.00032423436641693115,0.00032423436641693115,0.004752073436975479,0.19525915384292603,0.000645607418846339,0.7374463482107769,0.010096126236021519,0.9483441710472107,22210.0,48877.0,65.30311989784241,6.296417713165283,91.71633815765381,15.135236263275146,8121.565703325586,51.878292456517535 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,170,True,P/P,0.0,1.0,4239.375,0.0,0.0,0.0,0.0,0.0,0.23934924602508545,0.000816589395981282,0.0,0.012065509334206581,0.609393298625946,17434.0,33915.0,41.42259883880615,5.447986602783203,66.94624543190002,15.123604536056519,6529.809137150889,40.242079149334415 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,171,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2552044987678528,0.0007539574871771038,0.0,0.012957851402461529,0.9256450533866882,30061.0,57344.0,67.40864515304565,6.411708354949951,93.7338399887085,15.03031325340271,9215.985489869168,59.73863961884915 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,172,True,P/P,3.725290298461914e-09,0.125,7077.875,0.875,0.0,0.0002167224884033203,0.0002167224884033203,0.0029397797770798206,0.20122964680194855,0.0006979977479204535,0.5678713036731126,0.010546840727329254,0.9061625003814697,27890.0,56623.0,67.26449012756348,6.354372978210449,93.82501912117004,15.14084506034851,9280.46751782287,60.14706623983376 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,173,True,P/P,0.0,0.75,5838.5,0.25,0.0,0.00023398548364639282,0.00023398548364639282,0.0054329088889062405,0.24299390614032745,0.0009631355060264468,0.6828829093962452,0.012573221698403358,0.7721552848815918,26791.0,46708.0,65.03313517570496,5.584020137786865,90.48845648765564,14.99523377418518,8697.18721366084,55.52279553051693 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,174,True,P/P,1.4901161193847656e-08,0.375,6983.25,0.625,0.0,0.00025425106287002563,0.00025425106287002563,0.00342313339933753,0.20679974555969238,0.0005789826391264796,0.5137594460224592,0.009869875386357307,0.7231879234313965,28242.0,55866.0,67.17415571212769,6.33830189704895,93.71159434318542,15.196976900100708,9137.93010127306,59.09653582878099 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,175,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.25974133610725403,0.0012134959688410163,0.0,0.013561592437326908,0.846505880355835,32993.0,57344.0,67.65587949752808,6.417689800262451,94.1069815158844,15.054551839828491,9367.976470046644,60.85128626527749 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,176,True,P/P,0.0,1.0,1252.25,0.0,0.0,0.0,0.0,0.0,0.1305636465549469,0.0009179672342725098,0.0,0.008086515590548515,0.22500760853290558,4010.0,10018.0,11.509917259216309,1.5398945808410645,33.09617209434509,15.089199542999268,7364.167967723222,42.89479418151605 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,177,True,P/P,1.4901161193847656e-08,0.375,6949.75,0.625,0.0,0.0004509836435317993,0.0004509836435317993,0.0045364368706941605,0.2722417712211609,0.0011573819210752845,0.6449212483819498,0.013816691935062408,0.9089870452880859,32162.0,55598.0,66.86636567115784,6.389747381210327,93.30280351638794,15.0562424659729,8981.768159360427,58.02786198053718 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,178,True,P/P,-3.725290298461914e-09,0.875,5255.5,0.125,0.0,4.1114166378974915e-05,4.1114166378974915e-05,0.004084807820618153,0.1544552445411682,0.0004317459824960679,0.6373427907792019,0.00821889378130436,0.8623244166374207,17460.0,42044.0,63.416361808776855,6.276356220245361,89.91302061080933,15.247311353683472,6909.448261183894,43.40726174571626 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,179,True,P/P,0.0,0.5,6909.25,0.5,0.0,-5.879253149032593e-05,-5.879253149032593e-05,0.004462003707885742,0.19104966521263123,0.0009416970424354076,0.6148268545851585,0.010440506041049957,0.8600207567214966,24553.0,55274.0,66.57276940345764,6.341283082962036,92.86210107803345,15.166271209716797,8964.768904470911,57.84796858039182 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,180,True,P/P,3.725290298461914e-09,0.125,7037.25,0.875,0.0,0.00016324780881404877,0.00016324780881404877,0.004414601717144251,0.3274317979812622,0.001246619038283825,0.8820591640870187,0.016325250267982483,0.8919916152954102,38493.0,56298.0,67.10969758033752,6.457458972930908,93.63968777656555,15.071508884429932,8988.063796236187,58.14888074629406 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,181,True,P/P,0.0,0.0,4575.625,0.0,0.0,0.0,0.0,0.0,0.19366517663002014,0.0007857180899009109,0.0,0.010247091762721539,0.7131105661392212,16755.0,36605.0,53.91102457046509,4.756424188613892,78.66994094848633,15.04392695426941,8052.004451179911,50.024106491281486 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,182,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.21798762679100037,0.001187777379527688,0.0,0.012423906475305557,0.9762205481529236,30353.0,57344.0,67.74124574661255,6.34259033203125,94.11235070228577,15.022711753845215,9359.415301790017,60.70251885800432 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,183,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.14147895574569702,0.0007206299342215061,0.0,0.008028434589505196,0.8915302753448486,21705.0,57344.0,67.43572235107422,6.3658576011657715,94.05638575553894,15.179531574249268,9341.142433576546,60.5964028736278 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,184,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.24620607495307922,0.0005266744992695749,0.0,0.011814305558800697,0.844855785369873,28747.0,57344.0,67.4690613746643,6.345963716506958,94.14406156539917,15.330825567245483,9345.747903122656,60.60664102356627 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,185,True,P/P,0.0,1.0,5382.25,0.0,0.0,0.0,0.0,0.0,0.18546247482299805,0.0007315294351428747,0.0,0.00967741385102272,0.9320632219314575,21787.0,43058.0,55.72701811790466,6.280710458755493,82.03851652145386,15.106676578521729,7193.522733830432,45.284342241065 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,186,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2745491862297058,0.0007207000744529068,0.0,0.014043432660400867,0.9837619066238403,33468.0,57344.0,67.36679768562317,6.386097192764282,93.94292640686035,15.264317750930786,9251.260767335461,59.96525011605322 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,187,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.21635808050632477,0.000841704080812633,0.0,0.011873781681060791,0.9511567950248718,27834.0,57344.0,67.54174494743347,6.3643717765808105,93.96286296844482,15.151348352432251,9352.359683785462,60.6774440396661 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,188,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.27580082416534424,0.0008800469804555178,0.0,0.013793015852570534,0.7458294630050659,34833.0,57344.0,67.61964130401611,6.388808012008667,94.0644359588623,15.057947874069214,9219.15011615689,59.7346833979277 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,189,True,P/P,0.0,0.5,6832.75,0.5,0.0,3.7495046854019165e-05,3.7495046854019165e-05,0.0062036155723035336,0.3026096224784851,0.0011508762836456299,0.6599964196657112,0.014925300143659115,0.8683499097824097,35634.0,54662.0,66.48062181472778,6.390932083129883,92.9448173046112,15.12624740600586,8825.180962428167,56.89573155053837 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,190,True,P/P,0.0,1.0,4533.125,0.0,0.0,0.0,0.0,0.0,0.14411865174770355,0.00047171468031592667,0.0,0.008243563584983349,0.7398809194564819,15540.0,36265.0,49.725037813186646,5.532226800918579,75.28740453720093,15.095679759979248,6833.501650285773,42.383301755044215 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,191,True,P/P,0.0,1.0,5052.375,0.0,0.0,0.0,0.0,0.0,0.14498800039291382,0.000302958651445806,0.0,0.007821770384907722,0.8040862083435059,16049.0,40419.0,53.57848525047302,6.215216636657715,79.91000509262085,15.174911737442017,6761.98245886905,42.304219143733704 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,192,True,P/P,-3.725290298461914e-09,0.875,5780.75,0.125,0.0,5.202367901802063e-05,5.202367901802063e-05,0.004775315523147583,0.1864776909351349,0.0005962266586720943,0.5353243487238736,0.009686596691608429,0.7664518356323242,19929.0,46246.0,64.85405325889587,6.292330741882324,91.43981695175171,15.347677230834961,7746.311820037525,49.21714567652037 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,193,True,P/P,3.725290298461914e-09,0.125,7141.5,0.875,0.0,0.0001041218638420105,0.0001041218638420105,0.004757222253829241,0.2716138958930969,0.0009540076134726405,0.7792197359001206,0.014989292249083519,0.8919297456741333,36059.0,57132.0,67.39301705360413,6.448682069778442,94.02946376800537,15.200502634048462,9191.752715171217,59.60613314284012 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,194,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.25145411491394043,0.0009266398847103119,0.0,0.013535451143980026,1.0708682537078857,33060.0,57344.0,67.32725286483765,6.441653728485107,93.91038370132446,15.220194816589355,9149.255608045822,59.280796161634186 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,195,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.529685378074646,0.0007142938557080925,0.0,0.017430029809474945,0.6006650924682617,47837.0,57344.0,67.5083909034729,6.414337635040283,94.17481541633606,15.269554376602173,9268.145460488066,60.12184330758554 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,196,True,P/P,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.17592614889144897,0.00040048049413599074,0.0,0.009420156478881836,1.0189785957336426,23020.0,57344.0,67.67054271697998,6.424857139587402,94.41756987571716,15.337631464004517,9216.024801518202,59.75418271883274 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,197,True,P/P,0.0,1.0,4851.625,0.0,0.0,0.0,0.0,0.0,0.17630070447921753,0.00047824339708313346,0.0,0.009460071101784706,0.9492120146751404,18535.0,38813.0,55.72050333023071,6.203697443008423,82.28286814689636,15.311254501342773,6478.92010921221,40.37982987598357 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,198,True,P/P,0.0,1.0,3726.875,0.0,0.0,0.0,0.0,0.0,0.26832443475723267,0.0015311014140024781,0.0,0.01334262266755104,0.5232681035995483,16927.0,29815.0,41.66998744010925,3.8468334674835205,65.65469741821289,15.163749694824219,8285.263775855268,50.682203685164396 +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,199,True,P/P,3.725290298461914e-09,0.125,6207.875,0.5,0.0,0.00010925158858299255,0.00010925158858299255,0.0035345242358744144,0.26132121682167053,0.0007576454663649201,0.9326679315792483,0.012930657714605331,0.9914264678955078,28506.0,49663.0,65.8700532913208,6.1737165451049805,92.20786714553833,15.261356115341187,8385.257777202136,53.748698893673684 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,0,True,R/R,-3.725290298461914e-09,0.875,6117.5,0.125,0.0,-1.4901161193847656e-08,-1.4901161193847656e-08,0.0,0.27224838733673096,0.0,0.8293514167290124,0.0,0.0,0.0,48940.0,98.81510138511658,13.872129201889038,132.95428729057312,15.586602210998535,3638.375339035335,23.188270327705833 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,1,True,R/R,1.4901161193847656e-08,0.375,2443.25,0.0,0.0,1.1175870895385742e-08,1.1175870895385742e-08,0.0,0.2995075285434723,0.0,1.8035462308901338,0.0,0.0,0.0,19546.0,42.12234663963318,4.593104124069214,66.38132929801941,14.868260860443115,4543.584330697696,27.165475877807562 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,2,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2952701449394226,0.0,0.0,0.0,0.0,0.0,57344.0,106.4632658958435,9.851000547409058,135.93639945983887,14.823027849197388,6016.990269896805,39.03980094689073 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,3,True,R/R,0.0,0.25,6857.5,0.75,0.0,0.0,0.0,0.0,0.28477537631988525,0.0,0.8737660290035038,0.0,0.0,0.0,54860.0,103.59118270874023,9.776686429977417,133.44563913345337,15.096342086791992,5776.9824754721285,37.29026238995406 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,4,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3236057162284851,0.0,0.0,0.0,0.0,0.0,57344.0,106.20321273803711,9.858642578125,135.84337162971497,14.916146516799927,5951.010630133317,38.56168965254508 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,5,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3471332788467407,0.0,0.0,0.0,0.0,0.0,57344.0,106.58334636688232,9.849411725997925,136.21419596672058,14.957462072372437,5996.165146018466,38.88744079068263 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,6,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.38764721155166626,0.0,0.0,0.0,0.0,0.0,57344.0,106.20599412918091,9.87514328956604,136.16213536262512,14.95702075958252,5948.514515466185,38.55275162919132 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,7,True,R/R,0.0,0.25,4486.0,0.0,0.0,-7.450580596923828e-09,-7.450580596923828e-09,0.0,0.4042288064956665,0.0,1.0160940407687022,0.0,0.0,0.0,35888.0,70.99342083930969,8.506583452224731,99.17592096328735,14.921639919281006,4383.07796209097,27.13571388741932 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,8,True,R/R,0.0,1.0,5226.625,0.0,0.0,0.0,0.0,0.0,0.21955132484436035,0.0,0.0,0.0,0.0,0.0,41813.0,81.5780291557312,8.575538873672485,110.04261326789856,15.050976514816284,5001.001415488422,31.377001681584908 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,9,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.43446874618530273,0.0,0.0,0.0,0.0,0.0,57344.0,106.62837481498718,9.81551456451416,136.2713418006897,14.880455017089844,6095.870311441934,39.59878602746921 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,10,True,R/R,0.0,0.5,6705.125,0.375,0.0,-3.725290298461914e-09,-3.725290298461914e-09,0.0,0.3331888020038605,0.0,0.8515880125931149,0.0,0.0,0.0,53641.0,102.01668453216553,10.841449499130249,132.59974479675293,15.03275179862976,5067.550968146655,32.59920157449957 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,11,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.41952162981033325,0.0,0.0,0.0,0.0,0.0,57344.0,106.75976800918579,9.822849750518799,136.2189061641693,14.87039589881897,6040.227552058683,39.195246495296686 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,12,True,R/R,-1.4901161193847656e-08,0.625,4056.0,0.0,0.0,2.2351741790771484e-08,2.2351741790771484e-08,0.0,0.20587757229804993,0.0,1.1861186777469594,0.0,0.0,0.0,32448.0,59.82288956642151,7.0359296798706055,86.56057786941528,14.984462976455688,4834.059470993585,29.689394285427575 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,13,True,R/R,-1.4901161193847656e-08,0.625,4301.25,0.0,0.0,2.9802322387695312e-08,2.9802322387695312e-08,0.0,0.2728572189807892,0.0,1.1366043295411832,0.0,0.0,0.0,34410.0,67.91412258148193,7.2561540603637695,94.9500617980957,15.013824224472046,4933.758642102313,30.46598259295703 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,14,True,R/R,0.0,0.75,4660.0,0.0,0.0,-1.1175870895385742e-08,-1.1175870895385742e-08,0.0,0.23002943396568298,0.0,1.0333441654273359,0.0,0.0,0.0,37280.0,78.74428701400757,7.409151554107666,105.73263192176819,14.909169435501099,5178.125269877145,32.203017312834 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,15,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.33019769191741943,0.0,0.0,0.0,0.0,0.0,57344.0,106.21630501747131,9.864837408065796,135.7345154285431,14.903329133987427,5971.6242021335165,38.71441462488549 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,16,True,R/R,0.0,1.0,4689.875,0.0,0.0,0.0,0.0,0.0,0.276747465133667,0.0,0.0,0.0,0.0,0.0,37519.0,80.39818501472473,8.431106567382812,108.31882071495056,14.893224954605103,4689.210162405293,29.2184809711265 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,17,True,R/R,1.4901161193847656e-08,0.375,6066.875,0.25,0.0,-2.60770320892334e-08,-2.60770320892334e-08,0.0,0.28465336561203003,0.0,0.8873158263796865,0.0,0.0,0.0,48535.0,97.19358205795288,9.30036997795105,126.12310886383057,14.866850852966309,5381.465821192442,34.29128877268018 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,18,True,R/R,3.725290298461914e-09,0.125,7132.5,0.75,0.0,7.450580596923828e-09,7.450580596923828e-09,0.0,0.25973057746887207,0.0,0.7212044644319678,0.0,0.0,0.0,57060.0,105.796058177948,9.898650169372559,135.1678385734558,14.826037645339966,5931.204229946745,38.43860829296844 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,19,True,R/R,-3.725290298461914e-09,0.875,4652.375,0.0,0.0,-1.4901161193847656e-08,-1.4901161193847656e-08,0.0,0.29449933767318726,0.0,0.9413672791325095,0.0,0.0,0.0,37219.0,76.11164999008179,8.281607627868652,104.26063537597656,15.113470554351807,4681.206579602738,29.099304288535333 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,20,True,R/R,1.4901161193847656e-08,0.375,6816.875,0.625,0.0,-3.725290298461914e-08,-3.725290298461914e-08,0.0,0.24625220894813538,0.0,0.8152251392493883,0.0,0.0,0.0,54535.0,103.47640776634216,9.686241149902344,132.92202854156494,15.027012586593628,5796.830105195905,37.40130199533276 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,21,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.41630423069000244,0.0,0.0,0.0,0.0,0.0,57344.0,106.26401209831238,9.841235160827637,135.62841892242432,14.919235706329346,5982.2227976844615,38.78180281329721 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,22,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2508235573768616,0.0,0.0,0.0,0.0,0.0,57344.0,106.1765387058258,9.85792851448059,135.67486262321472,14.9394690990448,5946.71487341644,38.531223132893125 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,23,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.36989474296569824,0.0,0.0,0.0,0.0,0.0,57344.0,106.6138846874237,9.927589893341064,136.03686952590942,14.929622173309326,5935.141033373024,38.48117384683777 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,24,True,R/R,-3.725290298461914e-09,0.875,5981.125,0.125,0.0,-5.587935447692871e-09,-5.587935447692871e-09,0.0,0.2737797498703003,0.0,0.838666294543068,0.0,0.0,0.0,47849.0,96.8875789642334,9.3179612159729,125.96221828460693,14.972325086593628,5512.711307602965,35.21945175095536 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,25,True,R/R,0.0,0.75,4911.875,0.0,0.0,2.9802322387695312e-08,2.9802322387695312e-08,0.0,0.25254160165786743,0.0,0.9503059283615437,0.0,0.0,0.0,39295.0,76.27589797973633,8.336561918258667,104.34955954551697,15.04804515838623,4871.781731055424,30.391672453229415 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,26,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2708841562271118,0.0,0.0,0.0,0.0,0.0,57344.0,106.2290210723877,9.865432500839233,135.7734396457672,14.896524906158447,5946.760697074138,38.53546602826195 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,27,True,R/R,3.725290298461914e-09,0.125,6500.5,0.5,0.0,3.725290298461914e-08,3.725290298461914e-08,0.0,0.38082581758499146,0.0,1.0635068274360884,0.0,0.0,0.0,52004.0,100.56073665618896,9.54113507270813,129.76586151123047,14.899926900863647,5605.325232218392,36.00769321343525 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,28,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2273862659931183,0.0,0.0,0.0,0.0,0.0,57344.0,107.02325057983398,9.903135538101196,136.6555371284485,14.947538137435913,6082.129774363316,39.5431592152662 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,29,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3708828091621399,0.0,0.0,0.0,0.0,0.0,57344.0,106.32073712348938,9.81855297088623,135.91527724266052,14.975807428359985,6004.367014262157,38.93266464625175 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,30,True,R/R,0.0,1.0,3108.5,0.0,0.0,0.0,0.0,0.0,0.20954370498657227,0.0,0.0,0.0,0.0,0.0,24868.0,60.98711681365967,5.526556968688965,86.19332838058472,14.876710891723633,4748.48139736718,28.84679357334061 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,31,True,R/R,0.0,1.0,5337.875,0.0,0.0,0.0,0.0,0.0,0.22268211841583252,0.0,0.0,0.0,0.0,0.0,42703.0,87.66281962394714,8.869309663772583,116.26382637023926,14.943670988082886,4969.989521897735,31.289083527934324 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,32,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2883322238922119,0.0,0.0,0.0,0.0,0.0,57344.0,107.22105598449707,9.8440682888031,136.98107171058655,15.096868991851807,6156.570289452184,40.05845648139664 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,33,True,R/R,1.4901161193847656e-08,0.375,6945.625,0.625,0.0,-2.2351741790771484e-08,-2.2351741790771484e-08,0.0,0.2602505087852478,0.0,0.7806399443076306,0.0,0.0,0.0,55565.0,104.282883644104,9.866642951965332,133.9430251121521,14.915541887283325,5810.937330354535,37.56417413616166 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,34,True,R/R,0.0,0.5,6497.25,0.5,0.0,-1.1175870895385742e-08,-1.1175870895385742e-08,0.0,0.2375715970993042,0.0,0.7908053271644351,0.0,0.0,0.0,51978.0,101.46691846847534,9.53995943069458,130.77171444892883,14.974579572677612,5777.917730592018,37.238323488052 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,35,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.32077518105506897,0.0,0.0,0.0,0.0,0.0,57344.0,106.19772243499756,9.952434301376343,135.98732161521912,15.066859483718872,5906.555419452072,38.28538422593681 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,36,True,R/R,3.725290298461914e-09,0.125,7068.0,0.875,0.0,1.862645149230957e-08,1.862645149230957e-08,0.0,0.26625216007232666,0.0,0.8614848506833118,0.0,0.0,0.0,56544.0,105.29426527023315,9.810298442840576,134.96621346473694,15.152300357818604,5900.946675132422,38.18230598038146 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,37,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2816567122936249,0.0,0.0,0.0,0.0,0.0,57344.0,106.56787848472595,9.82964825630188,136.12128615379333,14.950417757034302,6013.8226928252125,39.006612348448655 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,38,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.29822665452957153,0.0,0.0,0.0,0.0,0.0,57344.0,106.97879648208618,9.84903621673584,136.60949420928955,14.850255727767944,6162.899802982872,40.107137455201766 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,39,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.30728551745414734,0.0,0.0,0.0,0.0,0.0,57344.0,106.5908796787262,9.838685512542725,136.20074677467346,15.003084897994995,5987.9298831157885,38.82144981271271 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,40,True,R/R,1.4901161193847656e-08,0.375,6855.625,0.625,0.0,-1.4901161193847656e-08,-1.4901161193847656e-08,0.0,0.2551245093345642,0.0,0.9143306958200721,0.0,0.0,0.0,54845.0,103.71215915679932,9.697691679000854,133.26171445846558,15.107828140258789,5867.454027559696,37.893810311611404 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,41,True,R/R,0.0,0.25,7050.625,0.75,0.0,4.842877388000488e-08,4.842877388000488e-08,0.0,0.28912225365638733,0.0,0.7836617437435477,0.0,0.0,0.0,56405.0,105.20541858673096,9.899551630020142,134.90533542633057,15.015484094619751,5866.315344762962,37.97075420766998 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,42,True,R/R,0.0,0.25,6968.25,0.625,0.0,1.4901161193847656e-08,1.4901161193847656e-08,0.0,0.36653923988342285,0.0,0.9532017911400337,0.0,0.0,0.0,55746.0,104.5614914894104,9.760751008987427,133.96329832077026,14.927266597747803,5912.034577244047,38.24201852978717 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,43,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.32147958874702454,0.0,0.0,0.0,0.0,0.0,57344.0,106.29735445976257,9.897690534591675,135.9020278453827,14.952976703643799,5952.980509683057,38.59683815630621 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,44,True,R/R,0.0,1.0,4243.75,0.0,0.0,0.0,0.0,0.0,0.21138995885849,0.0,0.0,0.0,0.0,0.0,33950.0,66.31960105895996,7.066495895385742,93.350013256073,14.974580764770508,4992.346842812979,30.778111474692942 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,45,True,R/R,0.0,1.0,4874.0,0.0,0.0,0.0,0.0,0.0,0.20010879635810852,0.0,0.0,0.0,0.0,0.0,38992.0,72.55945420265198,8.476034879684448,100.88292741775513,15.098487377166748,4784.557821646986,29.820402257482584 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,46,True,R/R,0.0,0.75,6037.125,0.25,0.0,0.0,0.0,0.0,0.3164231777191162,0.0,0.8480147908738217,0.0,0.0,0.0,48297.0,96.58438777923584,9.305993795394897,125.73517966270447,15.065187931060791,5353.792606594454,34.11683074795467 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,47,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.47959819436073303,0.0,0.0,0.0,0.0,0.0,57344.0,106.82682824134827,9.812131643295288,136.32354831695557,14.990311861038208,6073.427954826934,39.43352192914919 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,48,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.4745352864265442,0.0,0.0,0.0,0.0,0.0,57344.0,106.49479532241821,9.810868978500366,136.1824493408203,14.934295177459717,6038.778052974463,39.17916223994184 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,49,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.30813440680503845,0.0,0.0,0.0,0.0,0.0,57344.0,106.24778628349304,9.804491519927979,135.6864812374115,14.865354299545288,5985.280277714886,38.78573779323426 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,50,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.25222480297088623,0.0,0.0,0.0,0.0,0.0,57344.0,106.66037225723267,9.81129264831543,136.2628858089447,15.021666765213013,5996.166543108947,38.8688822659432 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,51,True,R/R,1.4901161193847656e-08,0.375,6949.25,0.5,0.0,-2.9802322387695312e-08,-2.9802322387695312e-08,0.0,0.2668120861053467,0.0,0.7702904843169147,0.0,0.0,0.0,55594.0,104.00471186637878,9.981399774551392,133.74109363555908,14.951888084411621,5684.76766754759,36.69337198810463 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,52,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.30809175968170166,0.0,0.0,0.0,0.0,0.0,57344.0,106.72320866584778,9.898323059082031,136.53400659561157,15.013869285583496,5942.1887171893295,38.51569662549951 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,53,True,R/R,0.0,1.0,5540.875,0.0,0.0,0.0,0.0,0.0,0.24656343460083008,0.0,0.0,0.0,0.0,0.0,44327.0,88.91462755203247,8.945593357086182,118.57406449317932,15.92919111251831,5088.491727795756,32.10356184865426 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,54,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2938944697380066,0.0,0.0,0.0,0.0,0.0,57344.0,106.40048933029175,9.902449131011963,137.03448700904846,15.912866830825806,5964.016726558475,38.677626761323395 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,55,True,R/R,-3.725290298461914e-09,0.875,6562.375,0.125,0.0,-3.725290298461914e-09,-3.725290298461914e-09,0.0,0.2775299549102783,0.0,0.780283406257209,0.0,0.0,0.0,52499.0,100.62243008613586,9.710201978683472,130.02934336662292,15.001430034637451,5537.3042058111,35.51775898214651 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,56,True,R/R,-1.4901161193847656e-08,0.625,6170.25,0.375,0.0,4.470348358154297e-08,4.470348358154297e-08,0.0,0.2046162635087967,0.0,0.7861293626792512,0.0,0.0,0.0,49362.0,97.864675283432,9.190547466278076,126.71147108078003,14.999943971633911,5538.948791285636,35.43850656176271 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,57,True,R/R,-3.725290298461914e-09,0.875,2978.75,0.0,0.0,-2.60770320892334e-08,-2.60770320892334e-08,0.0,0.2358812391757965,0.0,1.4702926001990881,0.0,0.0,0.0,23830.0,49.6390118598938,5.527188777923584,74.83760261535645,14.985628366470337,4581.695489050815,27.64256719480009 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,58,True,R/R,0.0,1.0,4510.875,0.0,0.0,0.0,0.0,0.0,0.24237072467803955,0.0,0.0,0.0,0.0,0.0,36087.0,69.9599142074585,7.91239857673645,97.57237434387207,14.941673040390015,4768.262988647216,29.550461733600347 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,59,True,R/R,0.0,1.0,5017.625,0.0,0.0,0.0,0.0,0.0,0.2102598249912262,0.0,0.0,0.0,0.0,0.0,40141.0,78.26121830940247,8.109006404876709,106.02419710159302,14.863475322723389,5170.745895265587,32.376801662276 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,60,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.28096139430999756,0.0,0.0,0.0,0.0,0.0,57344.0,106.51224207878113,9.877191543579102,136.1457769870758,14.98118782043457,5997.908731405037,38.914668310909704 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,61,True,R/R,1.4901161193847656e-08,0.375,7036.25,0.625,0.0,-2.2351741790771484e-08,-2.2351741790771484e-08,0.0,0.3298438787460327,0.0,0.9002060773474415,0.0,0.0,0.0,56290.0,105.42333936691284,9.90343189239502,135.16806197166443,14.97621750831604,5888.386674553906,38.13348647118055 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,62,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.36733803153038025,0.0,0.0,0.0,0.0,0.0,57344.0,106.75248694419861,9.894654512405396,136.4301495552063,14.97618055343628,6083.057013407893,39.54582400205364 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,63,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3548415005207062,0.0,0.0,0.0,0.0,0.0,57344.0,106.32678389549255,9.813729286193848,135.99433970451355,15.049269199371338,6007.40541231056,38.95236579576498 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,64,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.37114471197128296,0.0,0.0,0.0,0.0,0.0,57344.0,106.22202038764954,9.825403213500977,135.77802300453186,15.010293960571289,5969.125482489209,38.679071280316414 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,65,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3503168225288391,0.0,0.0,0.0,0.0,0.0,57344.0,106.46703314781189,9.836357593536377,136.16922092437744,14.923014640808105,6006.590600323505,38.957711001373305 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,66,True,R/R,0.0,1.0,4224.0,0.0,0.0,0.0,0.0,0.0,0.31599417328834534,0.0,0.0,0.0,0.0,0.0,33792.0,65.19593548774719,7.601423978805542,92.6731686592102,15.138614416122437,4703.056736681896,29.018996544982876 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,67,True,R/R,1.4901161193847656e-08,0.375,7027.0,0.625,0.0,-4.470348358154297e-08,-4.470348358154297e-08,0.0,0.22928835451602936,0.0,0.7965533037279625,0.0,0.0,0.0,56216.0,105.34095072746277,9.878845930099487,134.99712133407593,15.019336462020874,5869.5127786628855,37.982721137592016 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,68,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.41317498683929443,0.0,0.0,0.0,0.0,0.0,57344.0,106.55006647109985,9.864076137542725,136.19565725326538,14.96815824508667,6048.294656413811,39.27368104290816 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,69,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.38546353578567505,0.0,0.0,0.0,0.0,0.0,57344.0,106.59827208518982,9.854967832565308,136.3621141910553,15.0770423412323,5976.796378861196,38.747285185475846 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,70,True,R/R,0.0,1.0,3331.625,0.0,0.0,0.0,0.0,0.0,0.23300614953041077,0.0,0.0,0.0,0.0,0.0,26653.0,53.46547722816467,5.598073244094849,78.79816770553589,15.010534524917603,4998.08255547503,30.328894297252425 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,71,True,R/R,3.725290298461914e-09,0.125,7121.625,0.875,0.0,3.725290298461914e-08,3.725290298461914e-08,0.0,0.3257186710834503,0.0,0.8702095916010776,0.0,0.0,0.0,56973.0,105.86921048164368,9.850673913955688,135.4983196258545,15.00736927986145,5945.6365445202955,38.520093608425555 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,72,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.23375561833381653,0.0,0.0,0.0,0.0,0.0,57344.0,106.54213905334473,9.843961715698242,136.27898144721985,15.143894910812378,6052.797785396634,39.29689690664787 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,73,True,R/R,0.0,1.0,5918.375,0.0,0.0,0.0,0.0,0.0,0.30153530836105347,0.0,0.0,0.0,0.0,0.0,47347.0,93.54915070533752,9.397582054138184,122.95200419425964,15.188085079193115,5246.38654360158,33.385740661337515 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,74,True,R/R,0.0,0.5,4155.75,0.0,0.0,-7.450580596923828e-09,-7.450580596923828e-09,0.0,0.23623663187026978,0.0,1.2216588055894206,0.0,0.0,0.0,33246.0,69.1568500995636,7.360716342926025,96.44385313987732,15.163073062896729,4741.963367901251,29.264940081630925 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,75,True,R/R,-1.4901161193847656e-08,0.625,6052.75,0.125,0.0,7.450580596923828e-09,7.450580596923828e-09,0.0,0.30924585461616516,0.0,1.0787967855622715,0.0,0.0,0.0,48422.0,97.9136700630188,9.674626588821411,127.4682092666626,14.976666688919067,5513.771884856736,35.37368359499907 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,76,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3060561716556549,0.0,0.0,0.0,0.0,0.0,57344.0,106.52511405944824,9.903577327728271,136.1592755317688,14.949386835098267,6001.928443375349,38.95667872713606 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,77,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.33469098806381226,0.0,0.0,0.0,0.0,0.0,57344.0,106.81695032119751,9.847840070724487,137.4838376045227,15.9849271774292,6019.4064857897365,39.05614368349134 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,78,True,R/R,0.0,0.5,6752.875,0.5,0.0,-1.862645149230957e-08,-1.862645149230957e-08,0.0,0.41742831468582153,0.0,1.0214754472564866,0.0,0.0,0.0,54023.0,102.81700229644775,9.736615657806396,132.2736575603485,14.944515943527222,5764.129191040407,37.17656550376393 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,79,True,R/R,0.0,1.0,5067.625,0.0,0.0,0.0,0.0,0.0,0.2915680706501007,0.0,0.0,0.0,0.0,0.0,40541.0,77.52631855010986,8.729440689086914,106.09501242637634,15.120133876800537,4775.482209090981,29.86245640545893 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,80,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.29485151171684265,0.0,0.0,0.0,0.0,0.0,57344.0,106.7823543548584,9.843806982040405,136.4395546913147,15.002877473831177,6003.200113845156,38.93572089064885 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,81,True,R/R,0.0,1.0,3349.625,0.0,0.0,0.0,0.0,0.0,0.2397172451019287,0.0,0.0,0.0,0.0,0.0,26797.0,55.97676181793213,5.720515251159668,81.57199048995972,15.074735641479492,4924.813717378232,29.893305679884858 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,82,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.39020875096321106,0.0,0.0,0.0,0.0,0.0,57344.0,106.35240912437439,9.853389263153076,135.8438367843628,14.898651361465454,5972.2950537274155,38.716121845669164 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,83,True,R/R,3.725290298461914e-09,0.125,7129.875,0.875,0.0,3.5390257835388184e-08,3.5390257835388184e-08,0.0,0.44043827056884766,0.0,0.8855494044418792,0.0,0.0,0.0,57039.0,106.36530947685242,9.858709812164307,142.44967007637024,14.922964096069336,5969.183547900011,38.69480955938866 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,84,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.5726093053817749,0.0,0.0,0.0,0.0,0.0,57344.0,106.47894048690796,9.87036395072937,136.27907395362854,15.09342646598816,6000.813125980792,38.9321848535696 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,85,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3491520285606384,0.0,0.0,0.0,0.0,0.0,57344.0,107.067143201828,9.824211120605469,136.6291332244873,15.085466861724854,6069.717994967984,39.41211896978199 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,86,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3279462456703186,0.0,0.0,0.0,0.0,0.0,57344.0,106.30723237991333,9.843793630599976,136.12620329856873,14.889881372451782,5949.784084838917,38.54716183091577 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,87,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3608044981956482,0.0,0.0,0.0,0.0,0.0,57344.0,106.57853746414185,9.828668117523193,136.37852001190186,15.226176738739014,6039.920446653738,39.19726147894369 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,88,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.251244455575943,0.0,0.0,0.0,0.0,0.0,57344.0,107.01968455314636,9.884571552276611,136.7266767024994,15.078819990158081,6026.666582580033,39.127910427510635 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,89,True,R/R,0.0,1.0,5813.875,0.0,0.0,0.0,0.0,0.0,0.393779993057251,0.0,0.0,0.0,0.0,0.0,46511.0,89.95168161392212,9.089494943618774,118.86382627487183,15.018502712249756,5260.085745247838,33.313098963505404 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,90,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.28112536668777466,0.0,0.0,0.0,0.0,0.0,57344.0,106.75732016563416,9.86957597732544,136.56151676177979,15.126234769821167,5998.259047507745,38.912961018807444 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,91,True,R/R,1.4901161193847656e-08,0.375,6975.25,0.625,0.0,-3.725290298461914e-08,-3.725290298461914e-08,0.0,0.5207419395446777,0.0,0.9746730684249759,0.0,0.0,0.0,55802.0,104.65662789344788,9.841892719268799,134.41095662117004,15.10407567024231,5864.999530210155,37.93936006623262 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,92,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3850135803222656,0.0,0.0,0.0,0.0,0.0,57344.0,106.39062929153442,9.838010549545288,136.21383929252625,15.19633936882019,5994.91052093602,38.873337161422334 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,93,True,R/R,0.0,1.0,3256.75,0.0,0.0,0.0,0.0,0.0,0.22374367713928223,0.0,0.0,0.0,0.0,0.0,26054.0,52.55140733718872,5.567930459976196,77.95188665390015,15.052792310714722,4923.183400068513,29.858651962374598 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,94,True,R/R,1.4901161193847656e-08,0.375,6752.5,0.625,0.0,0.0,0.0,0.0,0.22581343352794647,0.0,0.7457431149038612,0.0,0.0,0.0,54020.0,102.60830640792847,9.712428569793701,131.75494360923767,14.90739369392395,5682.151644089104,36.57832168312189 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,95,True,R/R,3.725290298461914e-09,0.125,7148.125,0.875,0.0,-7.450580596923828e-09,-7.450580596923828e-09,0.0,0.2239258736371994,0.0,0.7658441580241572,0.0,0.0,0.0,57185.0,105.98967838287354,9.991149663925171,135.9380567073822,15.13006854057312,5861.9912745915235,37.976109180348 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,96,True,R/R,-3.725290298461914e-09,0.875,4928.875,0.0,0.0,-2.0489096641540527e-08,-2.0489096641540527e-08,0.0,0.4052520990371704,0.0,1.2423779803408328,0.0,0.0,0.0,39431.0,77.97472834587097,8.468602418899536,106.30384039878845,15.071055889129639,4846.443823544827,30.277527453630057 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,97,True,R/R,3.725290298461914e-09,0.125,6745.5,0.75,0.0,1.1175870895385742e-08,1.1175870895385742e-08,0.0,0.3589279055595398,0.0,1.0690678508661433,0.0,0.0,0.0,53964.0,102.65692257881165,9.622698545455933,132.00204038619995,14.914534091949463,5762.504404158269,37.137722049879024 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,98,True,R/R,0.0,0.75,6144.625,0.25,0.0,-1.1175870895385742e-08,-1.1175870895385742e-08,0.0,0.20730414986610413,0.0,0.7951150420461451,0.0,0.0,0.0,49157.0,97.55708241462708,9.419418334960938,126.7019693851471,15.065852165222168,5375.583731119635,34.29645375773765 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,99,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2882508337497711,0.0,0.0,0.0,0.0,0.0,57344.0,106.68945717811584,9.834395170211792,136.37773275375366,15.062776565551758,6011.66905708709,38.99330836635217 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,100,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2735469937324524,0.0,0.0,0.0,0.0,0.0,57344.0,106.28754448890686,9.806392192840576,135.9364790916443,15.042081594467163,6005.374243866311,38.9338824107502 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,101,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.361169695854187,0.0,0.0,0.0,0.0,0.0,57344.0,106.8209319114685,9.901817798614502,136.48715496063232,14.942768335342407,5997.253930072009,38.92169515525892 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,102,True,R/R,0.0,0.0,6724.0,0.5,0.0,0.0,0.0,0.0,0.18880769610404968,0.0,0.0,0.0,0.0,0.0,53792.0,102.39194560050964,9.553484678268433,131.79536747932434,15.029273271560669,5800.723553909788,37.35912684574977 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,103,True,R/R,3.725290298461914e-09,0.125,7016.875,0.875,0.0,2.60770320892334e-08,2.60770320892334e-08,0.0,0.2530791163444519,0.0,0.8328052556222023,0.0,0.0,0.0,56135.0,105.5742506980896,9.827818870544434,135.1044671535492,15.00511360168457,6011.749909847417,39.00494428804341 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,104,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.29566729068756104,0.0,0.0,0.0,0.0,0.0,57344.0,106.69199013710022,9.851261854171753,136.30255961418152,15.153182029724121,6001.256569681742,38.92577016216928 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,105,True,R/R,0.0,0.75,5590.75,0.25,0.0,-2.2351741790771484e-08,-2.2351741790771484e-08,0.0,0.32656675577163696,0.0,0.9455473742042074,0.0,0.0,0.0,44726.0,93.33641409873962,8.94580626487732,122.42829370498657,15.104126453399658,5224.581578467926,33.129176828294774 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,106,True,R/R,3.725290298461914e-09,0.125,7049.375,0.875,0.0,2.9802322387695312e-08,2.9802322387695312e-08,0.0,0.3635643720626831,0.0,0.9229689427561741,0.0,0.0,0.0,56395.0,105.63518595695496,9.83190655708313,135.2051763534546,14.914693832397461,5891.9273261128,38.12742905737707 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,107,True,R/R,0.0,1.0,3624.75,0.0,0.0,0.0,0.0,0.0,0.21269957721233368,0.0,0.0,0.0,0.0,0.0,28998.0,56.18070149421692,6.38499641418457,82.35863637924194,14.968626022338867,4704.07131306218,28.66321240318907 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,108,True,R/R,0.0,1.0,3041.875,0.0,0.0,0.0,0.0,0.0,0.15064120292663574,0.0,0.0,0.0,0.0,0.0,24335.0,52.651655197143555,5.404544830322266,77.93315744400024,15.059983253479004,4708.672321417239,28.46702605389734 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,109,True,R/R,0.0,0.25,6751.375,0.75,0.0,0.0,0.0,0.0,0.27480289340019226,0.0,1.1846440756931296,0.0,0.0,0.0,54011.0,102.89007329940796,9.654811143875122,132.3534152507782,15.029286861419678,5775.618409426517,37.270949923295845 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,110,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.24997618794441223,0.0,0.0,0.0,0.0,0.0,57344.0,106.71621870994568,9.79280424118042,136.17845392227173,14.963071346282959,6051.600350358378,39.26369091000331 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,111,True,R/R,0.0,0.0,6863.5,0.625,0.0,0.0,0.0,0.0,0.3493388593196869,0.0,0.0,0.0,0.0,0.0,54908.0,103.96071672439575,9.721823930740356,133.46536540985107,14.94528579711914,5892.6128200437815,38.08725674699258 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,112,True,R/R,0.0,1.0,4166.0,0.0,0.0,0.0,0.0,0.0,0.18109069764614105,0.0,0.0,0.0,0.0,0.0,33328.0,70.31276106834412,7.003165006637573,96.95837330818176,14.908791065216064,4950.507115892953,30.537230313934376 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,113,True,R/R,-3.725290298461914e-09,0.875,6562.375,0.125,0.0,-1.30385160446167e-08,-1.30385160446167e-08,0.0,0.3049176037311554,0.0,0.8726368054142074,0.0,0.0,0.0,52499.0,100.60189032554626,9.73605489730835,130.11544919013977,14.971629619598389,5504.721831374881,35.30069018986287 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,114,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.34199514985084534,0.0,0.0,0.0,0.0,0.0,57344.0,106.5103063583374,9.830886363983154,136.22577285766602,15.204817295074463,6039.328355522597,39.193418986134105 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,115,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.4919268488883972,0.0,0.0,0.0,0.0,0.0,57344.0,106.46011233329773,9.80409049987793,136.1487057209015,15.043242692947388,6050.799644977336,39.26318000361286 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,116,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.26069390773773193,0.0,0.0,0.0,0.0,0.0,57344.0,106.49633073806763,9.82240080833435,136.1509108543396,15.061467170715332,6050.456554512545,39.27032156488468 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,117,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.28970998525619507,0.0,0.0,0.0,0.0,0.0,57344.0,106.69473791122437,9.824509859085083,136.35557746887207,15.041386365890503,6026.6201529596965,39.09761666599564 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,118,True,R/R,-1.4901161193847656e-08,0.625,6242.25,0.375,0.0,4.470348358154297e-08,4.470348358154297e-08,0.0,0.21079449355602264,0.0,0.8558641329989549,0.0,0.0,0.0,49938.0,98.41873836517334,9.40349531173706,127.6920440196991,15.154156923294067,5471.657773900555,34.97376839176818 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,119,True,R/R,0.0,1.0,2841.75,0.0,0.0,0.0,0.0,0.0,0.17748993635177612,0.0,0.0,0.0,0.0,0.0,22734.0,40.90266513824463,5.349882364273071,66.22668194770813,15.083173274993896,4576.318820390376,27.504851040464022 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,120,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2349124401807785,0.0,0.0,0.0,0.0,0.0,57344.0,106.51370573043823,9.858698844909668,136.20049285888672,15.042034387588501,5942.688822932927,38.50316504652622 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,121,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.320705771446228,0.0,0.0,0.0,0.0,0.0,57344.0,106.25722432136536,9.844151735305786,135.83053636550903,14.951059579849243,5960.026702399443,38.61945336777992 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,122,True,R/R,0.0,1.0,2942.125,0.0,0.0,0.0,0.0,0.0,0.1814902275800705,0.0,0.0,0.0,0.0,0.0,23537.0,46.6624436378479,5.30816388130188,71.7007896900177,14.92196249961853,4716.41507517789,28.420477099550258 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,123,True,R/R,0.0,0.5,6723.25,0.5,0.0,-1.1175870895385742e-08,-1.1175870895385742e-08,0.0,0.23358435928821564,0.0,0.7345190109924371,0.0,0.0,0.0,53786.0,102.27165150642395,9.670225381851196,131.6536180973053,14.887176752090454,5706.9307407977785,36.737154771879865 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,124,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.26362764835357666,0.0,0.0,0.0,0.0,0.0,57344.0,106.37381839752197,9.834612607955933,136.29474258422852,14.961465120315552,6007.93605269899,38.96577294383925 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,125,True,R/R,0.0,1.0,4858.5,0.0,0.0,0.0,0.0,0.0,0.275801420211792,0.0,0.0,0.0,0.0,0.0,38868.0,79.54970455169678,8.505369186401367,107.8731255531311,14.983816146850586,4779.379014277609,29.844267734431252 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,126,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2718536853790283,0.0,0.0,0.0,0.0,0.0,57344.0,106.65323376655579,9.859356880187988,136.22447156906128,14.932157039642334,5980.302397366537,38.775305524652076 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,127,True,R/R,0.0,0.0,6372.75,0.375,0.0,0.0,0.0,0.0,0.42641231417655945,0.0,0.0,0.0,0.0,0.0,50982.0,99.5882842540741,9.38172197341919,128.86347913742065,15.061702251434326,5627.754692144418,36.08764338232775 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,128,True,R/R,0.0,1.0,5203.5,0.0,0.0,0.0,0.0,0.0,0.3665175139904022,0.0,0.0,0.0,0.0,0.0,41628.0,81.31767416000366,8.66981291770935,109.8952419757843,15.10841155052185,4959.071390306749,31.099986499639126 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,129,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.29115769267082214,0.0,0.0,0.0,0.0,0.0,57344.0,106.30830907821655,9.835144758224487,136.03457236289978,15.058851480484009,5981.96530109182,38.77682571614169 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,130,True,R/R,0.0,1.0,3412.875,0.0,0.0,0.0,0.0,0.0,0.2000093311071396,0.0,0.0,0.0,0.0,0.0,27303.0,53.072378158569336,5.607032299041748,78.32843971252441,14.878668069839478,5127.23637919981,31.146842906051766 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,131,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.4319975972175598,0.0,0.0,0.0,0.0,0.0,57344.0,106.41915535926819,9.913719177246094,136.12615847587585,14.964315176010132,5956.874260242546,38.63196547830976 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,132,True,R/R,0.0,1.0,4393.25,0.0,0.0,0.0,0.0,0.0,0.23911124467849731,0.0,0.0,0.0,0.0,0.0,35146.0,69.23645186424255,7.768442153930664,97.78591728210449,15.974226474761963,4777.430558370154,29.57121464575847 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,133,True,R/R,0.0,0.25,6742.75,0.75,0.0,0.0,0.0,0.0,0.19753292202949524,0.0,0.7621552219414436,0.0,0.0,0.0,53942.0,102.59901881217957,9.685630321502686,132.27884340286255,15.205614805221558,5689.731437322571,36.67403979749687 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,134,True,R/R,0.0,1.0,3884.75,0.0,0.0,0.0,0.0,0.0,0.24136915802955627,0.0,0.0,0.0,0.0,0.0,31078.0,61.74530005455017,6.307037591934204,87.9226496219635,15.07162356376648,5154.4498830885705,31.597019991826993 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,135,True,R/R,0.0,1.0,4770.125,0.0,0.0,0.0,0.0,0.0,0.22912871837615967,0.0,0.0,0.0,0.0,0.0,38161.0,83.05987668037415,7.3149330615997314,110.27387976646423,15.06610631942749,5386.464478116082,33.6415855274577 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,136,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.29016029834747314,0.0,0.0,0.0,0.0,0.0,57344.0,106.35691976547241,9.826253414154053,136.02771186828613,15.046844005584717,5988.83303723953,38.82266899820101 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,137,True,R/R,0.0,1.0,4796.375,0.0,0.0,0.0,0.0,0.0,0.1824309527873993,0.0,0.0,0.0,0.0,0.0,38371.0,72.33375096321106,8.381801843643188,100.72745990753174,15.137195348739624,4753.258240979486,29.58561380294257 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,138,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.34472280740737915,0.0,0.0,0.0,0.0,0.0,57344.0,106.66308379173279,9.860812902450562,140.0180413722992,18.700576066970825,5984.022396958633,38.80207253539916 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,139,True,R/R,0.0,1.0,5233.0,0.0,0.0,0.0,0.0,0.0,0.2399088442325592,0.0,0.0,0.0,0.0,0.0,41864.0,87.77922749519348,8.533170938491821,116.19944596290588,15.120784997940063,5117.437557142433,32.1998972856662 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,140,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3161659240722656,0.0,0.0,0.0,0.0,0.0,57344.0,106.25823736190796,9.818536043167114,135.79692339897156,14.981595277786255,5976.256745112817,38.727263611035035 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,141,True,R/R,3.725290298461914e-09,0.125,7167.25,0.875,0.0,-3.725290298461914e-09,-3.725290298461914e-09,0.0,0.33183735609054565,0.0,0.8826371415430503,0.0,0.0,0.0,57338.0,107.23869204521179,9.891815662384033,136.9379904270172,15.12413239479065,6156.258285162805,40.08178776972741 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,142,True,R/R,0.0,1.0,3834.625,0.0,0.0,0.0,0.0,0.0,0.2444797158241272,0.0,0.0,0.0,0.0,0.0,30677.0,61.687840700149536,6.182439088821411,87.6212375164032,14.986059665679932,5210.7250149412,31.921556264040746 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,143,True,R/R,0.0,1.0,5016.125,0.0,0.0,0.0,0.0,0.0,0.3366270959377289,0.0,0.0,0.0,0.0,0.0,40129.0,83.13061332702637,8.561321258544922,111.51257109642029,15.138887405395508,4910.272585696943,30.750157012852974 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,144,True,R/R,0.0,1.0,3599.75,0.0,0.0,0.0,0.0,0.0,0.251398503780365,0.0,0.0,0.0,0.0,0.0,28798.0,56.036277294158936,5.6889190673828125,81.5456793308258,15.057341575622559,5521.622944114486,33.75625030285137 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,145,True,R/R,-1.4901161193847656e-08,0.625,6572.125,0.375,0.0,4.842877388000488e-08,4.842877388000488e-08,0.0,0.2958953082561493,0.0,0.8921654078859936,0.0,0.0,0.0,52577.0,100.99381303787231,9.788766860961914,130.56375932693481,15.039905071258545,5510.918414545211,35.37989617530043 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,146,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3647816777229309,0.0,0.0,0.0,0.0,0.0,57344.0,106.30568647384644,9.91199254989624,136.01618218421936,15.068060874938965,5939.216065327897,38.50299697969601 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,147,True,R/R,-3.725290298461914e-09,0.875,6358.625,0.125,0.0,-1.862645149230957e-09,-1.862645149230957e-09,0.0,0.2572970390319824,0.0,0.8926396254839866,0.0,0.0,0.0,50869.0,100.12709665298462,9.47467041015625,129.44930863380432,15.051716566085815,5736.609060956017,36.86734338172835 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,148,True,R/R,0.0,1.0,3989.875,0.0,0.0,0.0,0.0,0.0,0.20656436681747437,0.0,0.0,0.0,0.0,0.0,31919.0,64.42954301834106,6.551229953765869,90.71045899391174,15.067670106887817,5059.765887669779,31.064815432097173 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,149,True,R/R,1.4901161193847656e-08,0.375,6846.0,0.625,0.0,0.0,0.0,0.0,0.2770436108112335,0.0,0.8796056380727433,0.0,0.0,0.0,54768.0,103.56030344963074,9.498394250869751,132.9811041355133,15.13652229309082,5952.188379660559,38.413781425072 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,150,True,R/R,0.0,1.0,2079.75,0.0,0.0,0.0,0.0,0.0,0.3441318869590759,0.0,0.0,0.0,0.0,0.0,16638.0,34.750821590423584,4.4592390060424805,59.16800117492676,15.172546625137329,3988.9027066334756,23.657169820536772 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,151,True,R/R,-3.725290298461914e-09,0.875,5174.75,0.125,0.0,0.0,0.0,0.0,0.19138669967651367,0.0,0.9414780775345,0.0,0.0,0.0,41398.0,89.48445415496826,8.848736763000488,118.18456840515137,15.125012874603271,4820.584925351536,30.276126051324106 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,152,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3625980317592621,0.0,0.0,0.0,0.0,0.0,57344.0,106.29785180091858,9.826718807220459,136.09006595611572,14.982118606567383,5974.4196634376995,38.71535897752586 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,153,True,R/R,0.0,1.0,5617.5,0.0,0.0,0.0,0.0,0.0,0.2675483226776123,0.0,0.0,0.0,0.0,0.0,44940.0,91.59982752799988,8.772034406661987,120.4679913520813,14.959800481796265,5310.399469584792,33.62820291452803 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,154,True,R/R,-3.725290298461914e-09,0.875,5486.5,0.125,0.0,-5.587935447692871e-09,-5.587935447692871e-09,0.0,0.25076621770858765,0.0,0.746985416949197,0.0,0.0,0.0,43892.0,92.19539999961853,8.776500940322876,120.77260422706604,15.011022329330444,5171.500481954117,32.67772002519447 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,155,True,R/R,0.0,1.0,4016.5,0.0,0.0,0.0,0.0,0.0,0.16167952120304108,0.0,0.0,0.0,0.0,0.0,32132.0,61.88561749458313,6.438212871551514,88.16743755340576,15.045132160186768,5212.336945109171,32.02126789963637 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,156,True,R/R,0.0,1.0,5361.25,0.0,0.0,0.0,0.0,0.0,0.25006091594696045,0.0,0.0,0.0,0.0,0.0,42890.0,85.13082098960876,8.895331382751465,113.76079773902893,14.97626519203186,5034.930033710516,31.692821168023894 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,157,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.21544548869132996,0.0,0.0,0.0,0.0,0.0,57344.0,106.31903028488159,9.918592691421509,136.03642916679382,15.008241415023804,5926.886623979272,38.41651250196119 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,158,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.22067534923553467,0.0,0.0,0.0,0.0,0.0,57344.0,106.47345614433289,9.939860343933105,137.25790786743164,16.02967667579651,5898.101544455803,38.217541850888196 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,159,True,R/R,0.0,0.5,2076.625,0.0,0.0,1.4901161193847656e-08,1.4901161193847656e-08,0.0,0.2065499722957611,0.0,1.8307076494047627,0.0,0.0,0.0,16613.0,30.99752163887024,4.38522744178772,55.549827575683594,15.1596360206604,4115.679889165522,24.380747003220094 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,160,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.36881357431411743,0.0,0.0,0.0,0.0,0.0,57344.0,106.48476839065552,9.82790732383728,136.2992308139801,15.170084953308105,6018.496699081193,39.040256685809304 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,161,True,R/R,0.0,1.0,4605.5,0.0,0.0,0.0,0.0,0.0,0.23782871663570404,0.0,0.0,0.0,0.0,0.0,36844.0,70.9047920703888,7.588473081588745,98.32134461402893,14.971625566482544,5027.442226328106,31.205784124248453 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,162,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.25266557931900024,0.0,0.0,0.0,0.0,0.0,57344.0,106.90417337417603,9.82996916770935,136.96385192871094,15.393736600875854,6143.690605506192,39.95698776784767 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,163,True,R/R,1.4901161193847656e-08,0.375,6770.75,0.625,0.0,-1.4901161193847656e-08,-1.4901161193847656e-08,0.0,0.2774495780467987,0.0,0.8176011827245694,0.0,0.0,0.0,54166.0,103.0316801071167,9.662811994552612,132.9263617992401,15.108536958694458,5807.466394793424,37.484926236196124 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,164,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2929115891456604,0.0,0.0,0.0,0.0,0.0,57344.0,106.55230164527893,9.799777030944824,136.4532208442688,15.150434494018555,5999.676688452111,38.888981976248616 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,165,True,R/R,0.0,1.0,4781.25,0.0,0.0,0.0,0.0,0.0,0.2765621542930603,0.0,0.0,0.0,0.0,0.0,38250.0,83.5475697517395,7.815023422241211,111.28426361083984,15.153486013412476,5064.568019157702,31.646082695599965 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,166,True,R/R,1.4901161193847656e-08,0.375,6774.0,0.625,0.0,-5.21540641784668e-08,-5.21540641784668e-08,0.0,0.4095819592475891,0.0,1.0471015229614187,0.0,0.0,0.0,54192.0,102.79275250434875,9.531656980514526,132.09194469451904,14.954890012741089,5825.929269516545,37.545153081836084 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,167,True,R/R,0.0,1.0,4059.0,0.0,0.0,0.0,0.0,0.0,0.24671649932861328,0.0,0.0,0.0,0.0,0.0,32472.0,60.3428099155426,6.93825364112854,87.18427991867065,15.063477754592896,4907.487580927157,30.150022029311586 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,168,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3675587475299835,0.0,0.0,0.0,0.0,0.0,57344.0,106.2992115020752,9.839010953903198,135.86057257652283,14.941007852554321,5972.337996128129,38.70913480411239 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,169,True,R/R,0.0,1.0,5620.375,0.0,0.0,0.0,0.0,0.0,0.2560022473335266,0.0,0.0,0.0,0.0,0.0,44963.0,92.19208717346191,8.984221458435059,120.87776279449463,14.876828670501709,5228.679686492196,33.15856251443465 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,170,True,R/R,0.0,1.0,3718.0,0.0,0.0,0.0,0.0,0.0,0.27128690481185913,0.0,0.0,0.0,0.0,0.0,29744.0,57.998475074768066,6.311166286468506,84.13588428497314,15.066863536834717,4954.745859952201,30.272394770146878 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,171,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.30513694882392883,0.0,0.0,0.0,0.0,0.0,57344.0,106.316246509552,9.808502912521362,136.01705884933472,15.143722534179688,5995.992771110747,38.86642960802823 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,172,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.29002976417541504,0.0,0.0,0.0,0.0,0.0,57344.0,106.47453927993774,9.827337741851807,136.06204986572266,15.020173072814941,6045.217678297744,39.23431312501361 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,173,True,R/R,0.0,0.75,5302.25,0.25,0.0,-3.725290298461914e-09,-3.725290298461914e-09,0.0,0.31783053278923035,0.0,0.9145071464267686,0.0,0.0,0.0,42418.0,91.10562133789062,8.280426502227783,119.55698323249817,15.216026067733765,5312.464555211632,33.54692398276475 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,174,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.2994121313095093,0.0,0.0,0.0,0.0,0.0,57344.0,106.42725777626038,9.923699378967285,137.36109519004822,16.098244190216064,5958.576010345745,38.65025047428992 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,175,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3000553250312805,0.0,0.0,0.0,0.0,0.0,57344.0,106.73192071914673,9.843688249588013,136.58534622192383,15.232363224029541,6075.144015190798,39.46213237759358 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,176,True,R/R,0.0,1.0,1362.625,0.0,0.0,0.0,0.0,0.0,0.19049625098705292,0.0,0.0,0.0,0.0,0.0,10901.0,20.290615558624268,3.5965168476104736,43.761765003204346,15.151153326034546,3302.3378934300363,19.287776705099915 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,177,True,R/R,0.0,0.5,6728.0,0.5,0.0,7.450580596923828e-09,7.450580596923828e-09,0.0,0.3442314565181732,0.0,0.9330366060327009,0.0,0.0,0.0,53824.0,102.38975167274475,9.786652326583862,132.11154556274414,15.169187068939209,5649.990239447905,36.37699654619442 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,178,True,R/R,-3.725290298461914e-09,0.875,5304.375,0.125,0.0,-1.862645149230957e-08,-1.862645149230957e-08,0.0,0.19511249661445618,0.0,0.8266789749085477,0.0,0.0,0.0,42435.0,90.46649670600891,8.631792783737183,118.9227819442749,15.034082412719727,5048.933217448199,31.765387409945294 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,179,True,R/R,0.0,0.5,6691.75,0.5,0.0,3.725290298461914e-09,3.725290298461914e-09,0.0,0.2506338059902191,0.0,0.8685627591457393,0.0,0.0,0.0,53534.0,101.97011637687683,9.574954986572266,131.48887252807617,15.159355640411377,5725.256920239932,36.83088951400693 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,180,True,R/R,3.725290298461914e-09,0.125,7166.5,0.875,0.0,3.725290298461914e-09,3.725290298461914e-09,0.0,0.46545249223709106,0.0,0.9741947308770822,0.0,0.0,0.0,57332.0,106.1377649307251,9.876896858215332,135.97729587554932,15.129446029663086,5953.555426153215,38.59102009054996 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,181,True,R/R,0.0,0.0,4472.0,0.125,0.0,0.0,0.0,0.0,0.2897905111312866,0.0,0.0,0.0,0.0,0.0,35776.0,83.808758020401,7.258536100387573,110.86910629272461,15.02374005317688,5126.300264012414,31.906928505495177 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,182,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3127349019050598,0.0,0.0,0.0,0.0,0.0,57344.0,106.66589951515198,9.83966326713562,136.29004955291748,15.005625009536743,6003.8592701937605,38.939332086272394 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,183,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.1816769540309906,0.0,0.0,0.0,0.0,0.0,57344.0,106.42355847358704,9.843867778778076,136.07164359092712,15.046072006225586,6011.829283986949,38.99900166284745 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,184,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.294482946395874,0.0,0.0,0.0,0.0,0.0,57344.0,106.37854290008545,9.802569389343262,136.01149249076843,15.059925079345703,6020.65774822644,39.04362140461684 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,185,True,R/R,0.0,0.5,6048.25,0.5,0.0,-3.3527612686157227e-08,-3.3527612686157227e-08,0.0,0.28785669803619385,0.0,1.1218478760056991,0.0,0.0,0.0,48386.0,97.13527965545654,9.155336856842041,126.19809818267822,15.133671045303345,5500.632280137702,35.13849363451478 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,186,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3374626636505127,0.0,0.0,0.0,0.0,0.0,57344.0,106.31134581565857,9.832192420959473,135.96868252754211,15.056167364120483,5979.871779112788,38.76060959848023 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,187,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.28030502796173096,0.0,0.0,0.0,0.0,0.0,57344.0,106.72685861587524,9.838476419448853,136.388507604599,15.023741722106934,6021.6286827957365,39.06789834670864 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,188,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3280041217803955,0.0,0.0,0.0,0.0,0.0,57344.0,106.19614171981812,9.7950119972229,135.83292412757874,15.054868221282959,5984.43160140316,38.77560539988149 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,189,True,R/R,-1.4901161193847656e-08,0.625,6372.25,0.25,0.0,2.2351741790771484e-08,2.2351741790771484e-08,0.0,0.38804101943969727,0.0,0.9955949918411783,0.0,0.0,0.0,50978.0,99.34619069099426,9.455724716186523,128.70133709907532,15.108193635940552,5539.566077975089,35.453950205716254 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,190,True,R/R,0.0,1.0,4307.5,0.0,0.0,0.0,0.0,0.0,0.19965213537216187,0.0,0.0,0.0,0.0,0.0,34460.0,75.20471000671387,7.157049179077148,102.31583499908447,15.012760639190674,5009.179907558157,30.98079175358185 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,191,True,R/R,-3.725290298461914e-09,0.875,5103.75,0.125,0.0,-1.4901161193847656e-08,-1.4901161193847656e-08,0.0,0.1859758496284485,0.0,0.7104328166733137,0.0,0.0,0.0,40830.0,88.72375750541687,8.517722845077515,117.16335701942444,15.066148281097412,4963.734133055374,31.11465303625673 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,192,True,R/R,0.0,1.0,5242.625,0.0,0.0,0.0,0.0,0.0,0.23547077178955078,0.0,0.0,0.0,0.0,0.0,41941.0,89.08780813217163,8.884622573852539,117.83596706390381,15.112238883972168,4974.449859034837,31.339109922756247 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,193,True,R/R,0.0,0.0,7162.375,0.875,0.0,0.0,0.0,0.0,0.4127928614616394,0.0,0.0,0.0,0.0,0.0,57299.0,106.62063717842102,9.84789490699768,136.64904928207397,15.124083757400513,6008.920861914948,38.979739337299534 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,194,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.3544083535671234,0.0,0.0,0.0,0.0,0.0,57344.0,106.21925687789917,9.862560749053955,136.05077481269836,15.074256896972656,5945.8125152377615,38.52471882205807 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,195,True,R/R,0.0,0.0,7168.0,1.0,0.0,0.0,0.0,0.0,0.6454512476921082,0.0,0.0,0.0,0.0,0.0,57344.0,106.45256519317627,9.88426661491394,136.2524118423462,15.048710107803345,5986.056152908138,38.83114821512515 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,196,True,R/R,1.4901161193847656e-08,0.375,7118.375,0.625,0.0,-5.21540641784668e-08,-5.21540641784668e-08,0.0,0.2456616908311844,0.0,0.7628938658255296,0.0,0.0,0.0,56947.0,105.62290453910828,9.922688961029053,135.5261697769165,15.148266077041626,5899.123567265307,38.21644256886224 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,197,True,R/R,0.0,1.0,4896.875,0.0,0.0,0.0,0.0,0.0,0.23570671677589417,0.0,0.0,0.0,0.0,0.0,39175.0,76.90373492240906,8.311903238296509,105.46857190132141,15.146685361862183,4862.282461280152,30.34066812519422 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,198,True,R/R,0.0,1.0,3148.75,0.0,0.0,0.0,0.0,0.0,0.3606951832771301,0.0,0.0,0.0,0.0,0.0,25190.0,47.816829442977905,5.446022987365723,73.21564793586731,15.102548837661743,4950.632815625408,29.946802588244978 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,199,True,R/R,0.0,0.0,7081.125,0.625,0.0,0.0,0.0,0.0,0.35494035482406616,0.0,0.0,0.0,0.0,0.0,56649.0,105.74253225326538,9.820414304733276,135.48555397987366,15.07497525215149,5956.361536554825,38.58503640502098 diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/runs.csv b/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/runs.csv new file mode 100644 index 00000000..fde15fc5 --- /dev/null +++ b/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/runs.csv @@ -0,0 +1,3 @@ +phase,run_id,group,seed,rollout_seed,num_rollout,rl_kernel_revision,vime_revision,prompt_data_sha256,cudagraph_passed,validation_passed,offline_tensor_comparison,rounds_observed +convergence,g10-convergence-s1234-tp4-20260901j,G10,1234,1234,200,d2173e8d948e8cf062ac36be32bdf53bac75daa0,1a113710e80aa7cfc271caa9bd90bcf348a7af08,73e2166517fd635e1157aff17202f86a5cced44ca1669e6f49d2d63a59bf509d,True,True,unavailable,200 +convergence,g11-convergence-s1234-tp4-20260901e,G11,1234,1234,200,5403df6e3c5244343438916248ccfcc597dd96f6,a013293fb6dfdc5cd27152b54f64209ea2691d26,73e2166517fd635e1157aff17202f86a5cced44ca1669e6f49d2d63a59bf509d,True,True,unavailable,200 diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/summary.csv b/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/summary.csv new file mode 100644 index 00000000..c3957589 --- /dev/null +++ b/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/summary.csv @@ -0,0 +1,3 @@ +phase,group,run_count,round_count,active_token_exposure,bitwise_mismatch_count,bitwise_mismatch_rate,mean_abs_dlogp_token_weighted,max_abs_dlogp,reward_mean,raw_reward_mean,truncated_ratio_mean,step_time_mean,actor_tokens_per_second_mean,unweighted_mean_abs_dlogp +convergence,G10,1,200,9927045.0,5764529.0,0.5806893189262263,0.012566951307437265,1.1917808055877686,4.842877388000488e-10,0.379375,0.58,87.69379462599754,8576.64340160429,0.012304821419529616 +convergence,G11,1,200,9806995.0,0.0,0.0,0.0,0.0,5.029141902923584e-10,0.395,0.561875,122.63016048192978,5564.762254530005,0.0 diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/summary.json b/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/summary.json new file mode 100644 index 00000000..f21350e5 --- /dev/null +++ b/examples/vime_qwen3_8b_tp4_cp2_200/results/convergence_s1234_g10_g11/summary.json @@ -0,0 +1,41 @@ +{ + "round_count": 400, + "schema_version": "rlkernel.vime_qwen3_8b_tp4_cp2_200.results.v1", + "sealed_run_count": 2, + "summaries": [ + { + "active_token_exposure": 9927045.0, + "actor_tokens_per_second_mean": 8576.64340160429, + "bitwise_mismatch_count": 5764529.0, + "bitwise_mismatch_rate": 0.5806893189262263, + "group": "G10", + "max_abs_dlogp": 1.1917808055877686, + "mean_abs_dlogp_token_weighted": 0.012566951307437265, + "phase": "convergence", + "raw_reward_mean": 0.379375, + "reward_mean": 4.842877388000488e-10, + "round_count": 200, + "run_count": 1, + "step_time_mean": 87.69379462599754, + "truncated_ratio_mean": 0.58, + "unweighted_mean_abs_dlogp": 0.012304821419529616 + }, + { + "active_token_exposure": 9806995.0, + "actor_tokens_per_second_mean": 5564.762254530005, + "bitwise_mismatch_count": 0.0, + "bitwise_mismatch_rate": 0.0, + "group": "G11", + "max_abs_dlogp": 0.0, + "mean_abs_dlogp_token_weighted": 0.0, + "phase": "convergence", + "raw_reward_mean": 0.395, + "reward_mean": 5.029141902923584e-10, + "round_count": 200, + "run_count": 1, + "step_time_mean": 122.63016048192978, + "truncated_ratio_mean": 0.561875, + "unweighted_mean_abs_dlogp": 0.0 + } + ] +} diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/results/scale_reference_s1234_g10_g11_optimized/README.md b/examples/vime_qwen3_8b_tp4_cp2_200/results/scale_reference_s1234_g10_g11_optimized/README.md new file mode 100644 index 00000000..40df7c48 --- /dev/null +++ b/examples/vime_qwen3_8b_tp4_cp2_200/results/scale_reference_s1234_g10_g11_optimized/README.md @@ -0,0 +1,27 @@ +# Version-aligned G10 vs optimized G11 (200 steps) + +Both runs use the same Qwen3-8B workload: one 8×H100 node, TP4/CP2, 200 steps, +seed and rollout seed 1234, rollout batch 8 prompts × 16 samples, global batch +128, maximum response length 7,168, and maximum 4,096 tokens/GPU. + +| Metric | G10 | Optimized G11 | Result | +|---|---:|---:|---| +| Rollout time (s) | 130.22 | 82.75 | G11 36.5% faster | +| Rollout tokens/GPU/s | 672.39 | 1134.00 | G11 68.7% higher | +| Reference logp time (s) | 20.90 | 20.92 | approximately equal | +| Actor train time (s) | 80.51 | 107.18 | G11 33.1% slower | +| Total step time (s) | 251.99 | 231.27 | G11 8.2% faster | +| Mean raw reward | 0.528555 | 0.491445 | G10−G11 +0.037109 | + +G11 has exactly zero mismatch count and zero maximum absolute difference at all +200 steps. G10 is the production P/P comparison and has non-zero mismatch at +all 200 steps. + +The paired mean reward difference (G10−G11) has a 95% bootstrap interval of +[+0.030664, +0.043633], using seed 1234 and +20,000 paired-step resamples. This is a single-training-seed result, not a +multi-seed generalization interval. + +`rounds.csv` contains every paired step. `summary.json` records formulas, +distribution summaries, and bootstrap details. `plot_report.py` regenerates all +five PNG figures from the authoritative Ray logs and the sealed G10 CSV. diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/results/scale_reference_s1234_g10_g11_optimized/consistency-reward.png b/examples/vime_qwen3_8b_tp4_cp2_200/results/scale_reference_s1234_g10_g11_optimized/consistency-reward.png new file mode 100644 index 00000000..7fa1adff Binary files /dev/null and b/examples/vime_qwen3_8b_tp4_cp2_200/results/scale_reference_s1234_g10_g11_optimized/consistency-reward.png differ diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/results/scale_reference_s1234_g10_g11_optimized/performance-matrix.png b/examples/vime_qwen3_8b_tp4_cp2_200/results/scale_reference_s1234_g10_g11_optimized/performance-matrix.png new file mode 100644 index 00000000..9a5aa0d9 Binary files /dev/null and b/examples/vime_qwen3_8b_tp4_cp2_200/results/scale_reference_s1234_g10_g11_optimized/performance-matrix.png differ diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/results/scale_reference_s1234_g10_g11_optimized/performance-statistics.png b/examples/vime_qwen3_8b_tp4_cp2_200/results/scale_reference_s1234_g10_g11_optimized/performance-statistics.png new file mode 100644 index 00000000..192abdb0 Binary files /dev/null and b/examples/vime_qwen3_8b_tp4_cp2_200/results/scale_reference_s1234_g10_g11_optimized/performance-statistics.png differ diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/results/scale_reference_s1234_g10_g11_optimized/performance-summary.png b/examples/vime_qwen3_8b_tp4_cp2_200/results/scale_reference_s1234_g10_g11_optimized/performance-summary.png new file mode 100644 index 00000000..5e2c0894 Binary files /dev/null and b/examples/vime_qwen3_8b_tp4_cp2_200/results/scale_reference_s1234_g10_g11_optimized/performance-summary.png differ diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/results/scale_reference_s1234_g10_g11_optimized/performance-trajectories.png b/examples/vime_qwen3_8b_tp4_cp2_200/results/scale_reference_s1234_g10_g11_optimized/performance-trajectories.png new file mode 100644 index 00000000..677cae2e Binary files /dev/null and b/examples/vime_qwen3_8b_tp4_cp2_200/results/scale_reference_s1234_g10_g11_optimized/performance-trajectories.png differ diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/results/scale_reference_s1234_g10_g11_optimized/plot_report.py b/examples/vime_qwen3_8b_tp4_cp2_200/results/scale_reference_s1234_g10_g11_optimized/plot_report.py new file mode 100644 index 00000000..e5ce7c7b --- /dev/null +++ b/examples/vime_qwen3_8b_tp4_cp2_200/results/scale_reference_s1234_g10_g11_optimized/plot_report.py @@ -0,0 +1,854 @@ +#!/usr/bin/env python3 +"""Build the version-aligned G10 vs optimized-G11 PR #377 result bundle.""" + +from __future__ import annotations + +import argparse +import ast +import csv +import json +import math +import re +import subprocess +from pathlib import Path +from typing import Any + +import matplotlib.pyplot as plt +import numpy as np + +RECORD_RE = re.compile(r"(?:perf|step|rollout)\s+(\d+):\s+(\{.*\})") +ANSI_RE = re.compile(r"\x1b\[[0-?]*[ -/]*[@-~]") +BLUE = "#2F67D8" +LIGHT_BLUE = "#9AB8EE" +RED = "#E45756" +LIGHT_RED = "#F3AAA7" +GREEN = "#2CA56C" +ORANGE = "#E28A2B" +GRID = "#D7DCE2" +TEXT = "#20242A" + + +def parse_args() -> argparse.Namespace: + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--g10-csv", type=Path, required=True) + parser.add_argument("--ray-bin") + parser.add_argument("--ssh-host", help="Optional SSH host from which to stream Ray logs") + parser.add_argument("--ssh-key", type=Path, help="SSH identity used with --ssh-host") + parser.add_argument("--g11-job", required=True) + parser.add_argument("--output-dir", type=Path, required=True) + return parser.parse_args() + + +def read_g10(path: Path) -> dict[int, dict[str, float]]: + rows: dict[int, dict[str, float]] = {} + with path.open(newline="", encoding="utf-8") as handle: + for raw in csv.DictReader(handle): + index = int(raw["log_step_index"]) + rows[index] = { + key: float(value) + for key, value in raw.items() + if key.startswith("g10_") and value not in (None, "") + } + if sorted(rows) != list(range(200)): + raise RuntimeError(f"G10 input has {len(rows)} steps; expected 0..199") + return rows + + +def read_g11( + ray_bin: str | Path | None, + job: str, + *, + ssh_host: str | None = None, + ssh_key: Path | None = None, +) -> dict[int, dict[str, Any]]: + if ssh_host: + if not ssh_key or not ray_bin: + raise RuntimeError("--ssh-host requires both --ssh-key and --ray-bin") + command = [ + "ssh", + "-i", + str(ssh_key), + "-o", + "BatchMode=yes", + "-o", + "ConnectTimeout=20", + ssh_host, + str(ray_bin), + "job", + "logs", + job, + ] + elif ray_bin: + command = [str(ray_bin), "job", "logs", job] + else: + raise RuntimeError("--ray-bin is required") + process = subprocess.Popen( + command, + text=True, + stdout=subprocess.PIPE, + stderr=subprocess.STDOUT, + errors="replace", + ) + assert process.stdout is not None + rows: dict[int, dict[str, Any]] = {} + for raw_line in process.stdout: + line = ANSI_RE.sub("", raw_line) + match = RECORD_RE.search(line) + if not match: + continue + try: + payload = ast.literal_eval(match.group(2)) + except (SyntaxError, ValueError): + continue + if isinstance(payload, dict): + rows.setdefault(int(match.group(1)), {}).update(payload) + return_code = process.wait() + if return_code != 0: + raise RuntimeError(f"ray job logs failed with return code {return_code}") + if sorted(rows) != list(range(200)): + raise RuntimeError(f"G11 logs have {len(rows)} merged steps; expected 0..199") + return rows + + +def finite(value: Any, *, field: str, step: int) -> float: + number = float(value) + if not math.isfinite(number): + raise RuntimeError(f"non-finite {field} at step {step + 1}: {value!r}") + return number + + +def build_rows( + g10: dict[int, dict[str, float]], g11: dict[int, dict[str, Any]] +) -> list[dict[str, float | int]]: + g11_keys = { + "reward": "rollout/raw_reward", + "reference_kl": "rollout/kl", + "kl_loss": "train/kl_loss", + "response_len_mean_tokens": "rollout/response_len/mean", + "response_len_max_tokens": "rollout/response_len/max", + "rollout_time_s": "perf/rollout_time", + "tokens_per_gpu_s": "perf/tokens_per_gpu_per_sec", + "longest_sample_tokens_s": "perf/longest_sample_tokens_per_sec", + "ref_log_probs_time_s": "perf/ref_log_probs_time", + "actor_train_time_s": "perf/actor_train_time", + "actor_train_tokens_s": "perf/actor_train_tok_per_s", + "train_time_s": "perf/train_time", + "step_time_s": "perf/step_time", + "mismatch_count": "train/train_current_rollout_logprob_mismatch_count", + "max_abs_diff": "train/train_current_rollout_logprob_max_abs_diff", + "active_tokens_mean": "train/train_current_rollout_logprob_numel", + } + result: list[dict[str, float | int]] = [] + for index in range(200): + source10 = g10[index] + source11 = g11[index] + row: dict[str, float | int] = {"step": index + 1, "log_step_index": index} + for metric, log_key in g11_keys.items(): + if log_key not in source11: + raise RuntimeError(f"missing G11 {log_key} at step {index + 1}") + row[f"g11_{metric}"] = finite(source11[log_key], field=log_key, step=index) + g10_key = f"g10_{metric}" + if g10_key not in source10: + raise RuntimeError(f"missing {g10_key} at step {index + 1}") + row[g10_key] = finite(source10[g10_key], field=g10_key, step=index) + for metric in g11_keys: + row[f"delta_g11_minus_g10_{metric}"] = float(row[f"g11_{metric}"]) - float( + row[f"g10_{metric}"] + ) + result.append(row) + return result + + +def values(rows: list[dict[str, Any]], key: str) -> np.ndarray: + return np.asarray([float(row[key]) for row in rows], dtype=np.float64) + + +def describe(array: np.ndarray) -> dict[str, float | int]: + return { + "n": int(array.size), + "mean": float(np.mean(array)), + "min": float(np.min(array)), + "median": float(np.median(array)), + "p90": float(np.percentile(array, 90)), + "p95": float(np.percentile(array, 95)), + "p99": float(np.percentile(array, 99)), + "max": float(np.max(array)), + } + + +def paired_bootstrap(delta: np.ndarray, *, draws: int = 20_000) -> dict[str, Any]: + rng = np.random.default_rng(1234) + indices = rng.integers(0, delta.size, size=(draws, delta.size)) + means = delta[indices].mean(axis=1) + return { + "estimate": float(delta.mean()), + "ci95": [float(np.percentile(means, 2.5)), float(np.percentile(means, 97.5))], + "seed": 1234, + "draws": draws, + "method": "paired non-parametric bootstrap over the 200 aligned steps", + } + + +def summarize(rows: list[dict[str, Any]]) -> dict[str, Any]: + metrics = [ + "reward", + "reference_kl", + "kl_loss", + "response_len_mean_tokens", + "response_len_max_tokens", + "rollout_time_s", + "tokens_per_gpu_s", + "longest_sample_tokens_s", + "ref_log_probs_time_s", + "actor_train_time_s", + "actor_train_tokens_s", + "train_time_s", + "step_time_s", + "mismatch_count", + "max_abs_diff", + "active_tokens_mean", + ] + paired: dict[str, Any] = {} + for metric in metrics: + g10 = values(rows, f"g10_{metric}") + g11 = values(rows, f"g11_{metric}") + paired[metric] = { + "g10": describe(g10), + "g11": describe(g11), + "g11_minus_g10_mean": float((g11 - g10).mean()), + "g11_over_g10": float(g11.mean() / g10.mean()) if g10.mean() else None, + } + return { + "paired": paired, + "bootstrap": { + "reward_g10_minus_g11": paired_bootstrap( + values(rows, "g10_reward") - values(rows, "g11_reward") + ), + "step_time_g11_minus_g10_s": paired_bootstrap( + values(rows, "g11_step_time_s") - values(rows, "g10_step_time_s") + ), + }, + "formulas": { + "rollout_tokens_per_gpu_s": ( + "mean response tokens * 128 samples / " "(rollout seconds * 8 GPUs)" + ), + "longest_sample_tokens_s": "max response tokens / rollout seconds", + "step_time_s": "VIME perf/step_time wall-clock timer", + }, + "missing_value_policy": ( + "All required values must exist and be finite for all 200 paired " + "steps; no imputation or row deletion." + ), + } + + +def moving_average(array: np.ndarray, window: int = 10) -> np.ndarray: + totals = np.convolve(array, np.ones(window), mode="full")[: array.size] + counts = np.minimum(np.arange(1, array.size + 1), window) + return totals / counts + + +def style(axis: plt.Axes) -> None: + axis.set_facecolor("white") + axis.grid(True, color=GRID, linestyle="--", linewidth=0.8) + axis.set_axisbelow(True) + axis.tick_params(colors=TEXT, labelsize=9.5) + for spine in axis.spines.values(): + spine.set_color("#AEB5BF") + + +def save(figure: plt.Figure, output_dir: Path, name: str) -> None: + figure.savefig(output_dir / f"{name}.png", dpi=220, bbox_inches="tight", facecolor="white") + plt.close(figure) + + +def grouped_bars(axis: plt.Axes, labels: list[str], g10: np.ndarray, g11: np.ndarray) -> None: + x = np.arange(len(labels)) + width = 0.36 + bars10 = axis.bar( + x - width / 2, g10, width, color=LIGHT_RED, edgecolor=RED, label="G10 production P/P" + ) + bars11 = axis.bar( + x + width / 2, g11, width, color=LIGHT_BLUE, edgecolor=BLUE, label="G11 optimized R/R" + ) + axis.set_xticks(x, labels) + for bars in (bars10, bars11): + for bar in bars: + axis.text( + bar.get_x() + bar.get_width() / 2, + bar.get_height() * 1.015, + f"{bar.get_height():,.1f}", + ha="center", + va="bottom", + fontsize=8.8, + ) + + +def plot_performance_summary(rows: list[dict[str, Any]], output_dir: Path) -> None: + fig, axes = plt.subplots(2, 2, figsize=(14.5, 9.2)) + time_metrics = [ + "rollout_time_s", + "ref_log_probs_time_s", + "actor_train_time_s", + "train_time_s", + "step_time_s", + ] + time_labels = ["Rollout", "Ref logp", "Actor train", "Total train", "Total step"] + g10_time = np.asarray([values(rows, f"g10_{m}").mean() for m in time_metrics]) + g11_time = np.asarray([values(rows, f"g11_{m}").mean() for m in time_metrics]) + grouped_bars(axes[0, 0], time_labels, g10_time, g11_time) + style(axes[0, 0]) + axes[0, 0].set_title("Mean stage time (lower is better)") + axes[0, 0].set_ylabel("Seconds / step") + axes[0, 0].legend(frameon=True, fontsize=9) + + throughput_metrics = ["tokens_per_gpu_s", "actor_train_tokens_s"] + throughput_labels = ["Rollout\ntok/GPU/s", "Actor train\ntok/s"] + g10_thr = np.asarray([values(rows, f"g10_{m}").mean() for m in throughput_metrics]) + g11_thr = np.asarray([values(rows, f"g11_{m}").mean() for m in throughput_metrics]) + grouped_bars(axes[0, 1], throughput_labels, g10_thr, g11_thr) + style(axes[0, 1]) + axes[0, 1].set_title("Token-normalized throughput (higher is better)") + axes[0, 1].set_ylabel("Tokens / second") + axes[0, 1].legend(frameon=True, fontsize=9) + + effect_labels = [ + "Rollout time", + "Rollout throughput", + "Ref logp time", + "Actor time", + "Actor throughput", + "Total step time", + ] + effects = np.asarray( + [ + 100 * (1 - g11_time[0] / g10_time[0]), + 100 * (g11_thr[0] / g10_thr[0] - 1), + 100 * (1 - g11_time[1] / g10_time[1]), + 100 * (1 - g11_time[2] / g10_time[2]), + 100 * (g11_thr[1] / g10_thr[1] - 1), + 100 * (1 - g11_time[4] / g10_time[4]), + ] + ) + colors = [GREEN if value >= 0 else RED for value in effects] + y = np.arange(len(effect_labels)) + axes[1, 0].barh(y, effects, color=colors, alpha=0.85) + axes[1, 0].axvline(0, color=TEXT, linewidth=1) + axes[1, 0].set_yticks(y, effect_labels) + axes[1, 0].invert_yaxis() + axes[1, 0].set_xlabel("G11 improvement over G10 (%)") + for yi, value in zip(y, effects, strict=True): + axes[1, 0].text( + value + (1.2 if value >= 0 else -1.2), + yi, + f"{value:+.1f}%", + va="center", + ha="left" if value >= 0 else "right", + fontweight="bold", + fontsize=9, + ) + style(axes[1, 0]) + axes[1, 0].set_title("Direction-aware performance delta") + + axes[1, 1].axis("off") + g10_reward = values(rows, "g10_reward").mean() + g11_reward = values(rows, "g11_reward").mean() + g10_mismatch = values(rows, "g10_mismatch_count") + g11_mismatch = values(rows, "g11_mismatch_count") + headline = ( + "200-step matched workload\n\n" + f"Total step: {g11_time[4]:.2f}s vs {g10_time[4]:.2f}s\n" + f"G11 end-to-end: {100 * (1 - g11_time[4] / g10_time[4]):.1f}% faster\n\n" + f"G11 bitwise mismatches: {int(g11_mismatch.sum()):,}\n" + f"G10 per-step mismatch mean: {g10_mismatch.mean():,.1f}\n\n" + f"Mean raw reward: G11 {g11_reward:.4f} · G10 {g10_reward:.4f}\n" + "1 node · 8×H100 · TP4/CP2 · batch 128" + ) + axes[1, 1].text( + 0.04, + 0.93, + headline, + va="top", + ha="left", + fontsize=14, + linespacing=1.45, + color=TEXT, + bbox={"boxstyle": "round,pad=0.8", "facecolor": "#F5F7FA", "edgecolor": "#C9D0D9"}, + ) + fig.suptitle("VIME Qwen3-8B · G10 vs Optimized G11", fontsize=19, fontweight="bold", y=0.995) + fig.tight_layout(rect=(0, 0, 1, 0.97)) + save(fig, output_dir, "performance-summary") + + +def plot_performance_matrix(rows: list[dict[str, Any]], output_dir: Path) -> None: + """Render a publication-friendly table in the style of the reference matrix.""" + + def mean(group: str, metric: str) -> float: + return float(values(rows, f"{group}_{metric}").mean()) + + def lower_delta(metric: str) -> str: + return f"{100 * (1 - mean('g11', metric) / mean('g10', metric)):+.1f}%" + + def higher_delta(metric: str) -> str: + return f"{100 * (mean('g11', metric) / mean('g10', metric) - 1):+.1f}%" + + performance = [ + [ + "Rollout time", + "s / step", + f"{mean('g10', 'rollout_time_s'):.2f}", + f"{mean('g11', 'rollout_time_s'):.2f}", + lower_delta("rollout_time_s"), + "G11 faster", + ], + [ + "Rollout throughput", + "tok/GPU/s", + f"{mean('g10', 'tokens_per_gpu_s'):,.0f}", + f"{mean('g11', 'tokens_per_gpu_s'):,.0f}", + higher_delta("tokens_per_gpu_s"), + "G11 higher", + ], + [ + "Reference logp", + "s / step", + f"{mean('g10', 'ref_log_probs_time_s'):.2f}", + f"{mean('g11', 'ref_log_probs_time_s'):.2f}", + lower_delta("ref_log_probs_time_s"), + "Parity", + ], + [ + "Actor train", + "s / step", + f"{mean('g10', 'actor_train_time_s'):.2f}", + f"{mean('g11', 'actor_train_time_s'):.2f}", + lower_delta("actor_train_time_s"), + "G10 faster", + ], + [ + "Actor throughput", + "tok/s", + f"{mean('g10', 'actor_train_tokens_s'):,.0f}", + f"{mean('g11', 'actor_train_tokens_s'):,.0f}", + higher_delta("actor_train_tokens_s"), + "G10 higher", + ], + [ + "Total train", + "s / step", + f"{mean('g10', 'train_time_s'):.2f}", + f"{mean('g11', 'train_time_s'):.2f}", + lower_delta("train_time_s"), + "G10 faster", + ], + [ + "End-to-end step", + "s / step", + f"{mean('g10', 'step_time_s'):.2f}", + f"{mean('g11', 'step_time_s'):.2f}", + lower_delta("step_time_s"), + "G11 faster", + ], + ] + correctness = [ + [ + "Mean raw reward", + "score", + f"{mean('g10', 'reward'):.6f}", + f"{mean('g11', 'reward'):.6f}", + f"{mean('g11', 'reward') - mean('g10', 'reward'):+.6f}", + "Report", + ], + [ + "Mean KL loss", + "loss", + f"{mean('g10', 'kl_loss'):.6f}", + f"{mean('g11', 'kl_loss'):.6f}", + higher_delta("kl_loss"), + "Report", + ], + [ + "Mismatch count", + "mean / step", + f"{mean('g10', 'mismatch_count'):,.1f}", + f"{mean('g11', 'mismatch_count'):.1f}", + "exactly zero", + "G11 PASS", + ], + [ + "Max |Delta logp|", + "max / 200", + f"{values(rows, 'g10_max_abs_diff').max():.6f}", + f"{values(rows, 'g11_max_abs_diff').max():.1f}", + "exactly zero", + "G11 PASS", + ], + ] + columns = ["Metric", "Unit", "G10 P/P", "Optimized G11 R/R", "G11 vs G10", "Finding"] + fig = plt.figure(figsize=(15.2, 9.3), facecolor="#F7E8D2") + fig.text( + 0.5, + 0.965, + "VIME Qwen3-8B · 200-Step Performance Matrix", + ha="center", + va="center", + fontsize=20, + fontweight="bold", + color="#171717", + bbox={ + "boxstyle": "square,pad=0.62", + "facecolor": "#F1C58B", + "edgecolor": "#80633A", + "linewidth": 1.5, + }, + ) + fig.text( + 0.5, + 0.905, + "Matched workload · 1 node · 8× NVIDIA H100 80GB · TP4/CP2 · global batch 128 · seed 1234", + ha="center", + va="center", + fontsize=12.5, + fontweight="bold", + color="#2A241A", + bbox={"boxstyle": "square,pad=0.48", "facecolor": "#F3D879", "edgecolor": "#9A8132"}, + ) + + def add_table(bounds: list[float], title: str, rows_data: list[list[str]]) -> None: + axis = fig.add_axes(bounds) + axis.axis("off") + axis.text( + 0.5, + 1.10, + title, + ha="center", + va="center", + fontsize=13.5, + fontweight="bold", + color="#2A241A", + transform=axis.transAxes, + bbox={"boxstyle": "square,pad=0.35", "facecolor": "#F3D879", "edgecolor": "#9A8132"}, + ) + table = axis.table( + cellText=rows_data, + colLabels=columns, + cellLoc="center", + colLoc="center", + loc="center", + colWidths=[0.205, 0.13, 0.14, 0.19, 0.16, 0.16], + bbox=[0, 0, 1, 1], + ) + table.auto_set_font_size(False) + table.set_fontsize(10.3) + table.scale(1, 1.5) + for (row, col), cell in table.get_celld().items(): + cell.set_edgecolor("#8D7A5D") + cell.set_linewidth(0.8) + if row == 0: + cell.set_facecolor("#F3D879") + cell.set_text_props(weight="bold", color="#221F19") + else: + cell.set_facecolor("#F7E7D2" if row % 2 else "#F2DCC2") + if col == 0: + cell.set_text_props(weight="bold") + if col == 5: + text_value = rows_data[row - 1][5] + color = ( + "#D9EAD3" + if "G11" in text_value or text_value == "Parity" + else "#EAD1DC" if "G10" in text_value else "#E2E3E5" + ) + cell.set_facecolor(color) + cell.set_text_props(weight="bold") + + add_table([0.035, 0.40, 0.93, 0.40], "Mean performance over 200 paired steps", performance) + add_table( + [0.035, 0.065, 0.93, 0.225], "Quality and strict train/rollout consistency", correctness + ) + fig.text( + 0.5, + 0.018, + "Times are arithmetic means; percentage signs are direction-aware " + "(positive means G11 is better). No missing values or imputation.", + ha="center", + fontsize=9.5, + color="#51483B", + ) + save(fig, output_dir, "performance-matrix") + + +def plot_performance_statistics(rows: list[dict[str, Any]], output_dir: Path) -> None: + """Plot paired mean effects with deterministic paired-bootstrap intervals.""" + specs = [ + ("End-to-end step", "step_time_s", "lower"), + ("Rollout time", "rollout_time_s", "lower"), + ("Rollout throughput", "tokens_per_gpu_s", "higher"), + ("Reference logp", "ref_log_probs_time_s", "lower"), + ("Actor train", "actor_train_time_s", "lower"), + ("Actor throughput", "actor_train_tokens_s", "higher"), + ("Total train", "train_time_s", "lower"), + ] + rng = np.random.default_rng(1234) + indices = rng.integers(0, 200, size=(20_000, 200)) + estimates: list[float] = [] + intervals: list[tuple[float, float]] = [] + for _, metric, direction in specs: + g10 = values(rows, f"g10_{metric}") + g11 = values(rows, f"g11_{metric}") + sign = 1.0 if direction == "higher" else -1.0 + estimate = sign * 100.0 * (g11.mean() / g10.mean() - 1.0) + boot = sign * 100.0 * (g11[indices].mean(axis=1) / g10[indices].mean(axis=1) - 1.0) + estimates.append(float(estimate)) + intervals.append((float(np.percentile(boot, 2.5)), float(np.percentile(boot, 97.5)))) + estimates_array = np.asarray(estimates) + lower = estimates_array - np.asarray([item[0] for item in intervals]) + upper = np.asarray([item[1] for item in intervals]) - estimates_array + y = np.arange(len(specs)) + fig, axis = plt.subplots(figsize=(12.8, 7.2)) + colors = [GREEN if value >= 0 else RED for value in estimates_array] + axis.barh(y, estimates_array, color=colors, alpha=0.82, height=0.62) + axis.errorbar( + estimates_array, + y, + xerr=np.vstack([lower, upper]), + fmt="none", + ecolor=TEXT, + elinewidth=1.5, + capsize=4, + ) + axis.axvline(0, color=TEXT, linewidth=1.2) + axis.set_yticks(y, [item[0] for item in specs]) + axis.invert_yaxis() + axis.set_xlabel("G11 improvement over G10 (%) · positive is better") + for yi, value, interval in zip(y, estimates_array, intervals, strict=True): + axis.text( + interval[1] + 1.2, + yi, + f"{value:+.1f}% [{interval[0]:+.1f}, {interval[1]:+.1f}]", + va="center", + fontsize=10, + fontweight="bold", + ) + style(axis) + axis.set_title( + "Paired Performance Effects · Mean and 95% Bootstrap CI", + fontsize=17, + fontweight="bold", + pad=16, + ) + fig.text( + 0.5, + 0.02, + "200 aligned steps · paired non-parametric bootstrap · 20,000 resamples · seed 1234", + ha="center", + fontsize=10, + color="#4B525B", + ) + fig.tight_layout(rect=(0, 0.05, 1, 1)) + save(fig, output_dir, "performance-statistics") + + +def plot_performance_trajectories(rows: list[dict[str, Any]], output_dir: Path) -> None: + steps = np.arange(1, 201) + fig, axes = plt.subplots(3, 1, figsize=(14, 10), sharex=True) + panels = [ + ("rollout_time_s", "Rollout generation", "Seconds"), + ("actor_train_time_s", "Actor training", "Seconds"), + ("step_time_s", "Total step", "Seconds"), + ] + for axis, (metric, title, ylabel) in zip(axes, panels, strict=True): + for group, raw, strong in (("G10", LIGHT_RED, RED), ("G11", LIGHT_BLUE, BLUE)): + series = values(rows, f"{group.lower()}_{metric}") + axis.plot(steps, series, color=raw, linewidth=0.8, alpha=0.65) + axis.plot( + steps, + moving_average(series), + color=strong, + linewidth=2.2, + label=f"{group} 10-step MA", + ) + style(axis) + axis.set_title(title) + axis.set_ylabel(ylabel) + axis.legend(ncol=2, frameon=True) + axes[-1].set_xlabel("Training step") + axes[-1].set_xlim(1, 200) + fig.suptitle( + "G10 vs Optimized G11 · Performance Across 200 Steps", + fontsize=18, + fontweight="bold", + y=0.995, + ) + fig.tight_layout(rect=(0, 0, 1, 0.97)) + save(fig, output_dir, "performance-trajectories") + + +def plot_consistency_reward(rows: list[dict[str, Any]], output_dir: Path) -> None: + steps = np.arange(1, 201) + fig, axes = plt.subplots(2, 2, figsize=(14.5, 9.2)) + for group, raw, strong in (("G10", LIGHT_RED, RED), ("G11", LIGHT_BLUE, BLUE)): + reward = values(rows, f"{group.lower()}_reward") + axes[0, 0].plot(steps, reward, color=raw, linewidth=0.8, alpha=0.65) + axes[0, 0].plot( + steps, moving_average(reward), color=strong, linewidth=2.2, label=f"{group} 10-step MA" + ) + kl = values(rows, f"{group.lower()}_kl_loss") + axes[0, 1].plot(steps, kl, color=strong, linewidth=1.5, label=group) + axes[1, 0].plot( + steps, + values(rows, f"{group.lower()}_mismatch_count"), + color=strong, + linewidth=1.5, + label=group, + ) + axes[1, 1].plot( + steps, + values(rows, f"{group.lower()}_max_abs_diff"), + color=strong, + linewidth=1.5, + label=group, + ) + titles = [ + "Raw reward", + "Reference KL loss", + "Train/rollout mismatch count", + "Maximum absolute Δlogp", + ] + ylabels = [ + "Reward", + "KL loss", + "Mismatched active tokens", + "Absolute log-probability difference", + ] + for axis, title, ylabel in zip(axes.flat, titles, ylabels, strict=True): + style(axis) + axis.set_title(title) + axis.set_ylabel(ylabel) + axis.set_xlabel("Training step") + axis.legend(frameon=True) + axes[0, 1].set_yscale("symlog", linthresh=1e-5) + fig.suptitle( + "G10 vs Optimized G11 · Training and Bitwise Consistency", + fontsize=18, + fontweight="bold", + y=0.995, + ) + fig.tight_layout(rect=(0, 0, 1, 0.97)) + save(fig, output_dir, "consistency-reward") + + +def write_csv(path: Path, rows: list[dict[str, Any]]) -> None: + with path.open("w", newline="", encoding="utf-8") as handle: + writer = csv.DictWriter(handle, fieldnames=list(rows[0]), lineterminator="\n") + writer.writeheader() + writer.writerows(rows) + + +def write_readme(path: Path, summary: dict[str, Any]) -> None: + p = summary["paired"] + reward_ci = summary["bootstrap"]["reward_g10_minus_g11"] + rollout = p["rollout_time_s"] + rollout_tput = p["tokens_per_gpu_s"] + ref_logp = p["ref_log_probs_time_s"] + actor = p["actor_train_time_s"] + step_time = p["step_time_s"] + reward = p["reward"] + rollout_gain = 100 * (1 - rollout["g11"]["mean"] / rollout["g10"]["mean"]) + rollout_tput_gain = 100 * (rollout_tput["g11"]["mean"] / rollout_tput["g10"]["mean"] - 1) + actor_slowdown = 100 * (actor["g11"]["mean"] / actor["g10"]["mean"] - 1) + step_gain = 100 * (1 - step_time["g11"]["mean"] / step_time["g10"]["mean"]) + table_rows = "\n".join( + [ + ( + f"| Rollout time (s) | {rollout['g10']['mean']:.2f} | " + f"{rollout['g11']['mean']:.2f} | G11 {rollout_gain:.1f}% faster |" + ), + ( + f"| Rollout tokens/GPU/s | {rollout_tput['g10']['mean']:.2f} | " + f"{rollout_tput['g11']['mean']:.2f} | " + f"G11 {rollout_tput_gain:.1f}% higher |" + ), + ( + f"| Reference logp time (s) | {ref_logp['g10']['mean']:.2f} | " + f"{ref_logp['g11']['mean']:.2f} | approximately equal |" + ), + ( + f"| Actor train time (s) | {actor['g10']['mean']:.2f} | " + f"{actor['g11']['mean']:.2f} | G11 {actor_slowdown:.1f}% slower |" + ), + ( + f"| Total step time (s) | {step_time['g10']['mean']:.2f} | " + f"{step_time['g11']['mean']:.2f} | G11 {step_gain:.1f}% faster |" + ), + ( + f"| Mean raw reward | {reward['g10']['mean']:.6f} | " + f"{reward['g11']['mean']:.6f} | G10−G11 " + f"{reward['g10']['mean']-reward['g11']['mean']:+.6f} |" + ), + ] + ) + text = f"""# Version-aligned G10 vs optimized G11 (200 steps) + +Both runs use the same Qwen3-8B workload: one 8×H100 node, TP4/CP2, 200 steps, +seed and rollout seed 1234, rollout batch 8 prompts × 16 samples, global batch +128, maximum response length 7,168, and maximum 4,096 tokens/GPU. + +| Metric | G10 | Optimized G11 | Result | +|---|---:|---:|---| +{table_rows} + +G11 has exactly zero mismatch count and zero maximum absolute difference at all +200 steps. G10 is the production P/P comparison and has non-zero mismatch at +all 200 steps. + +The paired mean reward difference (G10−G11) has a 95% bootstrap interval of +[{reward_ci['ci95'][0]:+.6f}, {reward_ci['ci95'][1]:+.6f}], using seed 1234 and +20,000 paired-step resamples. This is a single-training-seed result, not a +multi-seed generalization interval. + +`rounds.csv` contains every paired step. `summary.json` records formulas, +distribution summaries, and bootstrap details. `plot_report.py` regenerates all +five PNG figures from the authoritative Ray logs and the sealed G10 CSV. +""" + path.write_text(text, encoding="utf-8", newline="\n") + + +def main() -> None: + args = parse_args() + args.output_dir.mkdir(parents=True, exist_ok=True) + rows = build_rows( + read_g10(args.g10_csv), + read_g11( + args.ray_bin, + args.g11_job, + ssh_host=args.ssh_host, + ssh_key=args.ssh_key, + ), + ) + summary = summarize(rows) + write_csv(args.output_dir / "rounds.csv", rows) + (args.output_dir / "summary.json").write_text( + json.dumps(summary, indent=2, sort_keys=True) + "\n", + encoding="utf-8", + newline="\n", + ) + write_readme(args.output_dir / "README.md", summary) + plot_performance_matrix(rows, args.output_dir) + plot_performance_statistics(rows, args.output_dir) + plot_performance_summary(rows, args.output_dir) + plot_performance_trajectories(rows, args.output_dir) + plot_consistency_reward(rows, args.output_dir) + print( + json.dumps( + { + "output_dir": str(args.output_dir), + "paired": summary["paired"], + "bootstrap": summary["bootstrap"], + }, + indent=2, + ) + ) + + +if __name__ == "__main__": + main() diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/results/scale_reference_s1234_g10_g11_optimized/rounds.csv b/examples/vime_qwen3_8b_tp4_cp2_200/results/scale_reference_s1234_g10_g11_optimized/rounds.csv new file mode 100644 index 00000000..2b0ce436 --- /dev/null +++ b/examples/vime_qwen3_8b_tp4_cp2_200/results/scale_reference_s1234_g10_g11_optimized/rounds.csv @@ -0,0 +1,201 @@ +step,log_step_index,g11_reward,g10_reward,g11_reference_kl,g10_reference_kl,g11_kl_loss,g10_kl_loss,g11_response_len_mean_tokens,g10_response_len_mean_tokens,g11_response_len_max_tokens,g10_response_len_max_tokens,g11_rollout_time_s,g10_rollout_time_s,g11_tokens_per_gpu_s,g10_tokens_per_gpu_s,g11_longest_sample_tokens_s,g10_longest_sample_tokens_s,g11_ref_log_probs_time_s,g10_ref_log_probs_time_s,g11_actor_train_time_s,g10_actor_train_time_s,g11_actor_train_tokens_s,g10_actor_train_tokens_s,g11_train_time_s,g10_train_time_s,g11_step_time_s,g10_step_time_s,g11_mismatch_count,g10_mismatch_count,g11_max_abs_diff,g10_max_abs_diff,g11_active_tokens_mean,g10_active_tokens_mean,delta_g11_minus_g10_reward,delta_g11_minus_g10_reference_kl,delta_g11_minus_g10_kl_loss,delta_g11_minus_g10_response_len_mean_tokens,delta_g11_minus_g10_response_len_max_tokens,delta_g11_minus_g10_rollout_time_s,delta_g11_minus_g10_tokens_per_gpu_s,delta_g11_minus_g10_longest_sample_tokens_s,delta_g11_minus_g10_ref_log_probs_time_s,delta_g11_minus_g10_actor_train_time_s,delta_g11_minus_g10_actor_train_tokens_s,delta_g11_minus_g10_train_time_s,delta_g11_minus_g10_step_time_s,delta_g11_minus_g10_mismatch_count,delta_g11_minus_g10_max_abs_diff,delta_g11_minus_g10_active_tokens_mean +1,0,0.265625,0.2890625,0.0,0.0,0.0,0.0,6084.4609375,6096.0703125,7168.0,7168.0,75.51525902748108,141.25830817222595,1289.1616377104983,690.4877048440939,94.92121317350528,50.74391795249721,23.863692045211792,26.11477017402649,116.55882525444031,95.08617210388184,6827.205046575271,8384.57351221369,140.66537046432495,121.47803211212158,236.34011459350586,282.79273891448975,0.0,3980.3046875,0.0,0.6022037267684937,6084.4609375,6096.0703125,-0.0234375,0.0,0.0,-11.609375,0.0,-65.74304914474487,598.6739328664045,44.177295221008066,-2.2510781288146973,21.47265315055847,-1557.3684656384194,19.18733835220337,-46.45262432098389,-3980.3046875,-0.6022037267684937,-11.609375 +2,1,0.3359375,0.3828125,0.0,0.0,0.000712860724888742,-1.4674065823783167e-05,5784.3828125,5768.6484375,7168.0,7168.0,90.09214925765991,132.13017749786377,1027.2829071411136,698.5412170621828,79.56298144802616,54.249529787514966,22.43020534515381,23.04313087463379,110.54298210144043,88.58934831619263,6855.7676443408045,8531.985101665374,133.20728874206543,111.86720776557922,243.8365352153778,263.9777925014496,0.0,3834.9453125,0.0,0.6568862199783325,5784.3828125,5768.6484375,-0.046875,0.0,0.0007275347907125251,15.734375,0.0,-42.03802824020386,328.7416900789308,25.3134516605112,-0.6129255294799805,21.953633785247803,-1676.2174573245693,21.340080976486206,-20.141257286071777,-3834.9453125,-0.6568862199783325,15.734375 +3,2,0.25,0.1875,0.0,0.0,0.0006669942522421479,-5.290961780701764e-06,6381.171875,6406.703125,7168.0,7168.0,73.77074527740479,143.26961016654968,1384.000522376067,715.4849509315773,97.16588836192068,50.03154536169438,23.570497512817383,25.40219259262085,119.60307955741882,96.56643486022949,6978.867125233859,8677.57001915695,143.40848922729492,122.20473909378052,237.49957275390625,285.64910101890564,0.0,4083.8671875,0.0,0.6736352443695068,6381.171875,6406.703125,0.0625,0.0,0.0006722852140228497,-25.53125,0.0,-69.4988648891449,668.5155714444898,47.1343430002263,-1.8316950798034668,23.03664469718933,-1698.7028939230913,21.203750133514404,-48.14952826499939,-4083.8671875,-0.6736352443695068,-25.53125 +4,3,0.4609375,0.4140625,0.0,0.0,0.0007254835218191147,-3.674428444355726e-05,5909.3515625,5958.46875,7168.0,7168.0,86.32255506515503,136.65579295158386,1095.3061448266365,697.6323355262126,83.03739381427827,52.45295384250245,21.96596360206604,25.01266384124756,110.86005091667175,92.60140895843506,7024.469081160144,8477.408808675504,133.05785655975342,117.85110354423523,239.7432084083557,274.68990540504456,0.0,3659.40625,0.0,0.6597182750701904,5909.3515625,5958.46875,0.046875,0.0,0.000762227806262672,-49.1171875,0.0,-50.33323788642883,397.67380930042384,30.584439971775822,-3.0467002391815186,18.258641958236694,-1452.9397275153597,15.206753015518188,-34.94669699668884,-3659.40625,-0.6597182750701904,-49.1171875 +5,4,0.09375,0.0546875,0.0,0.0,0.0006345350993797183,-9.201021384797059e-06,7063.2265625,7103.40625,7168.0,7168.0,85.00034356117249,160.1422040462494,1329.5431555364073,709.709852420767,84.32907091535908,44.760218224109536,23.69849395751953,25.934202671051025,127.20533752441406,100.66892504692078,7321.05285928952,9301.976747676048,151.13978481292725,126.84098148345947,256.1628189086914,306.910710811615,0.0,4342.1953125,0.0,0.7133777141571045,7063.2265625,7103.40625,0.0390625,0.0,0.0006437361207645154,-40.1796875,0.0,-75.1418604850769,619.8333031156403,39.568852691249546,-2.235708713531494,26.536412477493286,-1980.9238883865282,24.298803329467773,-50.747891902923584,-4342.1953125,-0.7133777141571045,-40.1796875 +6,5,0.3046875,0.2890625,0.0,0.0,0.0007149508455768228,2.658294943103101e-05,6473.1875,6526.09375,7168.0,7168.0,76.50349497795105,146.43432807922363,1353.8074310180211,713.0670886372234,93.69506585373489,48.950270705117596,23.34084963798523,26.129151344299316,119.49083304405212,97.23577761650085,7102.737326193963,8798.03731682015,143.06632900238037,123.60168290138245,239.67205619812012,290.24485754966736,0.0,4230.3125,0.0,0.7200061678886414,6473.1875,6526.09375,0.015625,0.0,0.0006883678961457917,-52.90625,0.0,-69.93083310127258,640.7403423807978,44.74479514861729,-2.788301706314087,22.25505542755127,-1695.2999906261866,19.464646100997925,-50.57280135154724,-4230.3125,-0.7200061678886414,-52.90625 +7,6,0.140625,0.1640625,0.0,0.0,0.0007667674799449742,3.9696584281045943e-05,7026.0703125,6974.140625,7168.0,7168.0,83.00801396369934,155.66282057762146,1354.2924307183366,716.8458697197856,86.35310806417648,46.04824693142232,23.82844853401184,25.64682126045227,126.35603857040405,100.06018543243408,7236.0055787168085,9071.200458776922,150.42072367668152,125.94506978988647,253.87423276901245,301.56313610076904,0.0,4402.0859375,0.0,0.6990071535110474,7026.0703125,6974.140625,-0.0234375,0.0,0.0007270708956639282,51.9296875,0.0,-72.65480661392212,637.446560998551,40.30486113275416,-1.8183727264404297,26.29585313796997,-1835.194880060114,24.475653886795044,-47.68890333175659,-4402.0859375,-0.6990071535110474,51.9296875 +8,7,0.3828125,0.359375,0.0,0.0,0.0007493073353543878,8.727327804081142e-05,6117.75,5945.8671875,7168.0,7168.0,72.90011692047119,137.07149863243103,1342.7138958746038,694.0456327475462,98.32631692236893,52.29387634566984,22.534987926483154,23.299910306930542,114.39365291595459,89.15067601203918,7038.852064559764,8785.115660752077,137.16255593299866,112.6868543624878,230.27771162986755,269.72412848472595,0.0,3762.15625,0.0,0.6611260771751404,6117.75,5945.8671875,0.0234375,0.0,0.0006620340573135763,171.8828125,0.0,-64.17138171195984,648.6682631270577,46.032440576699095,-0.7649223804473877,25.242976903915405,-1746.2635961923133,24.475701570510864,-39.4464168548584,-3762.15625,-0.6611260771751404,171.8828125 +9,8,0.2734375,0.28125,0.0,0.0,0.0007488494738936424,0.0001040944189298898,6332.078125,6283.3671875,7168.0,7168.0,76.89779472351074,144.5529181957245,1317.5052726060098,695.4814628085006,93.21463672362576,49.58737664703895,21.97661519050598,23.140987634658813,114.92500829696655,89.53767895698547,7219.542659122872,9196.92144795937,137.13611245155334,112.91694784164429,234.1822793483734,277.462443113327,0.0,4298.3828125,0.0,0.6516216993331909,6332.078125,6283.3671875,-0.0078125,0.0,0.0006447550549637526,48.7109375,0.0,-67.65512347221375,622.0238097975092,43.62726007658681,-1.164372444152832,25.38732933998108,-1977.3787888364986,24.219164609909058,-43.28016376495361,-4298.3828125,-0.6516216993331909,48.7109375 +10,9,0.28125,0.359375,0.0,0.0,0.000817956228274852,0.0001463514199713245,6532.109375,6423.078125,7168.0,7168.0,76.76470732688904,144.73364400863647,1361.4817751463122,710.0577803034353,93.37624345360075,49.52545794792726,23.73345708847046,25.465274333953857,121.4333393573761,98.60245275497437,7119.478098643641,8626.4182708892,145.39990258216858,124.30501127243042,242.50367140769958,289.1481671333313,0.0,4113.59375,0.0,0.723059892654419,6532.109375,6423.078125,-0.078125,0.0,0.0006716048083035275,109.03125,0.0,-67.96893668174744,651.4239948428769,43.85078550567349,-1.7318172454833984,22.830886602401733,-1506.9401722455596,21.09489130973816,-46.644495725631714,-4113.59375,-0.723059892654419,109.03125 +11,10,0.109375,0.0859375,0.0,0.0,0.0008709862595424056,0.00019862415501847863,6655.75,6646.7734375,7168.0,7168.0,80.98465633392334,152.9777708053589,1314.9651405681398,695.1884214296223,88.51059354310583,46.85648092702434,22.691545009613037,24.325316190719604,120.06068658828735,94.07751250267029,7259.928497568937,9252.827555100293,142.98743748664856,118.64040422439575,243.97635436058044,291.48740696907043,0.0,4758.0078125,0.0,0.6639820337295532,6655.75,6646.7734375,0.0234375,0.0,0.000672362104523927,8.9765625,0.0,-71.99311447143555,619.7767191385175,41.654112616081484,-1.6337711811065674,25.983174085617065,-1992.8990575313564,24.347033262252808,-47.51105260848999,-4758.0078125,-0.6639820337295532,8.9765625 +12,11,0.3203125,0.265625,0.0,0.0,0.0009113270207308233,0.00028101992211304605,6431.0,6485.34375,7168.0,7168.0,77.01611065864563,148.24735116958618,1336.0321511957468,699.9484252592035,93.07143581646626,48.35162276727787,22.850786685943604,24.585490942001343,117.0029046535492,95.03204703330994,7188.197613472562,8923.263535539403,140.08953881263733,119.85596489906311,237.24193239212036,287.9791316986084,0.0,4260.8125,0.0,0.6969559192657471,6431.0,6485.34375,0.0546875,0.0,0.0006303070986177772,-54.34375,0.0,-71.23124051094055,636.0837259365433,44.719813049188396,-1.7347042560577393,21.970857620239258,-1735.0659220668413,20.23357391357422,-50.73719930648804,-4260.8125,-0.6969559192657471,-54.34375 +13,12,0.1796875,0.171875,0.0,0.0,0.0008430213783867657,0.000376501731807366,6874.0,6848.8203125,7168.0,7168.0,78.62020516395569,152.85804653167725,1398.9279189826307,716.881626361038,91.17249166485577,46.8931807166236,23.75480628013611,25.506823539733887,124.88073682785034,100.39033365249634,7215.716553964468,8943.898952542957,148.8730766773224,126.13527703285217,247.67507147789001,299.42618560791016,0.0,4067.375,0.0,0.7591615319252014,6874.0,6848.8203125,0.0078125,0.0,0.0004665196465793997,25.1796875,0.0,-74.23784136772156,682.0462926215927,44.27931094823217,-1.7520172595977783,24.490403175354004,-1728.1823985784886,22.737799644470215,-51.75111413002014,-4067.375,-0.7591615319252014,25.1796875 +14,13,0.3671875,0.359375,0.0,0.0,0.0009466548217460513,0.0006614563753828406,6102.6484375,6076.328125,7168.0,7168.0,74.46841311454773,140.60563206672668,1311.1918317610173,691.4463423048522,96.25557602488645,50.979465720109346,21.460724592208862,22.161086320877075,111.03829741477966,85.14374160766602,7209.215366566403,9362.16784638261,132.73233485221863,107.54205775260925,227.34471225738525,268.3511025905609,0.0,3702.890625,0.0,0.6690464019775391,6102.6484375,6076.328125,0.0078125,0.0,0.0002851984463632107,26.3203125,0.0,-66.13721895217896,619.745489456165,45.276110304777106,-0.7003617286682129,25.894555807113647,-2152.9524798162074,25.190277099609375,-41.00639033317566,-3702.890625,-0.6690464019775391,26.3203125 +15,14,0.34375,0.3984375,0.0,0.0,0.0010082339867949486,0.0007493121665902436,6206.6484375,6081.3359375,7168.0,7168.0,74.95943784713745,139.4591670036316,1324.8014906743635,697.7051210801095,95.62505010533148,51.39855739862078,21.61498475074768,22.461721658706665,112.77708411216736,86.40737390518188,7218.514349868439,9235.820554802685,134.62438774108887,109.1053946018219,229.6506006717682,268.7142331600189,0.0,3762.96875,0.0,0.6645004153251648,6206.6484375,6081.3359375,-0.0546875,0.0,0.000258921820204705,125.3125,0.0,-64.49972915649414,627.096369594254,44.2264927067107,-0.8467369079589844,26.369710206985474,-2017.3062049342461,25.518993139266968,-39.06363248825073,-3762.96875,-0.6645004153251648,125.3125 +16,15,0.296875,0.2890625,0.0,0.0,0.0009228518465533853,0.0006607099203392863,6420.375,6347.4375,7168.0,7168.0,77.94829845428467,143.80149960517883,1317.8735397315572,706.2443735207229,91.95838962673326,49.84648991617229,22.726919889450073,24.317172527313232,118.04665565490723,93.48403859138489,7106.613866744697,8873.985468535913,141.00832295417786,118.04398322105408,238.980122089386,282.10977149009705,0.0,3754.03125,0.0,0.7192646265029907,6420.375,6347.4375,0.0078125,0.0,0.00026214192621409893,72.9375,0.0,-65.85320115089417,611.6291662108343,42.11189971056098,-1.5902526378631592,24.56261706352234,-1767.3716017912166,22.96433973312378,-43.12964940071106,-3754.03125,-0.7192646265029907,72.9375 +17,16,0.53125,0.515625,0.0,0.0,0.001262323698028922,0.001020960509777069,5795.2734375,5672.8984375,7168.0,7168.0,81.23074221611023,132.19050860404968,1141.493632463822,686.6330719089112,88.24245358893685,54.22477056556545,21.335344791412354,21.822431802749634,107.84097266197205,84.70865058898926,7029.341272506081,8764.004559606316,129.40917325019836,106.76566934585571,230.62474513053894,258.8312635421753,0.0,3380.828125,0.0,0.6724258661270142,5795.2734375,5672.8984375,0.015625,0.0,0.000241363188251853,122.375,0.0,-50.95976638793945,454.86056055491076,34.017683023371404,-0.4870870113372803,23.132322072982788,-1734.6632871002348,22.64350390434265,-28.206518411636353,-3380.828125,-0.6724258661270142,122.375 +18,17,0.46875,0.4609375,0.0,0.0,0.0010526957921683788,0.0011903722770512104,6193.1171875,6109.7578125,7168.0,7168.0,74.79426908493042,139.42461156845093,1324.8324532389163,701.1396617878066,95.83621964218395,51.4112961790885,22.463701963424683,24.098953247070312,115.52955055236816,93.26494526863098,7052.039898923502,8621.127666821634,138.22729778289795,117.59765338897705,233.0146963596344,277.1389527320862,0.0,3708.015625,0.0,0.7090646028518677,6193.1171875,6109.7578125,0.0078125,0.0,-0.00013767648488283157,83.359375,0.0,-64.63034248352051,623.6927914511097,44.424923463095446,-1.6352512836456299,22.264605283737183,-1569.0877678981324,20.6296443939209,-44.12425637245178,-3708.015625,-0.7090646028518677,83.359375 +19,18,0.609375,0.625,0.0,0.0,0.0018642160575836897,0.002102834638208151,5520.859375,5415.515625,7168.0,7168.0,85.51658606529236,127.83942937850952,1032.9428952245207,677.7897118380445,83.81999714683646,56.07033788282051,20.155421018600464,20.309549570083618,102.74633812904358,78.31824111938477,7086.9484329988945,9125.256004033383,123.1337821483612,98.8604326248169,228.63529896736145,246.51128435134888,0.0,3068.5859375,0.0,0.6968894004821777,5520.859375,5415.515625,-0.015625,0.0,-0.00023861858062446117,105.34375,0.0,-42.32284331321716,355.15318338647626,27.749659264015946,-0.1541285514831543,24.428097009658813,-2038.3075710344883,24.27334952354431,-17.875985383987427,-3068.5859375,-0.6968894004821777,105.34375 +20,19,0.4765625,0.5390625,0.0,0.0,0.0013392632827162743,0.0015923543833196163,5996.4765625,5819.0390625,7168.0,7168.0,74.75797581672668,135.59859204292297,1283.3898183012755,686.6194080431768,95.88274591025242,52.86190580600579,21.307152032852173,22.376373767852783,109.9072778224945,85.65980100631714,7134.714056574588,8889.175448164371,131.445214509964,108.27153134346008,226.42052960395813,263.8059124946594,0.0,3118.5703125,0.0,0.6564945578575134,5996.4765625,5819.0390625,-0.0625,0.0,-0.00025309110060334206,177.4375,0.0,-60.84061622619629,596.7704102580987,43.02084010424664,-1.0692217350006104,24.247476816177368,-1754.4613915897835,23.173683166503906,-37.385382890701294,-3118.5703125,-0.6564945578575134,177.4375 +21,20,0.3828125,0.3984375,0.0,0.0,0.0013123640092089772,0.0019212639890611172,6325.2890625,6290.75,7168.0,7168.0,75.545982837677,144.91725301742554,1339.6427076401246,694.5480810894002,94.88260964718177,49.46270958598757,23.38938808441162,24.94454550743103,118.7495493888855,96.52980518341064,6985.76966623541,8547.991974418754,142.37055039405823,121.71083927154541,237.99326968193054,286.6295602321625,0.0,3842.4140625,0.0,0.7862330675125122,6325.2890625,6290.75,-0.015625,0.0,-0.00060889997985214,34.5390625,0.0,-69.37127017974854,645.0946265507245,45.4199000611942,-1.5551574230194092,22.219744205474854,-1562.2223081833436,20.659711122512817,-48.636290550231934,-3842.4140625,-0.7862330675125122,34.5390625 +22,21,0.3984375,0.3671875,0.0,0.0,0.001361468923278153,0.0019440214382484555,6432.859375,6524.671875,7168.0,7168.0,75.47169494628906,147.75604248046875,1363.766244725909,706.534556878102,94.97600398535172,48.51239840798733,23.170296669006348,25.418222665786743,120.21986865997314,97.19155740737915,7022.716040290412,8807.575707548109,143.62660670280457,122.84597945213318,239.3022711277008,290.7105915546417,0.0,3838.296875,0.0,0.7943591475486755,6432.859375,6524.671875,0.03125,0.0,-0.0005825525149703026,-91.8125,0.0,-72.28434753417969,657.231687847807,46.4636055773644,-2.2479259967803955,23.028311252593994,-1784.8596672576969,20.780627250671387,-51.40832042694092,-3838.296875,-0.7943591475486755,-91.8125 +23,22,0.296875,0.3515625,0.0,0.0,0.0013427475932985544,0.0030021369457244873,6008.4765625,5907.328125,7168.0,7168.0,72.73514175415039,137.52131366729736,1321.7218346111815,687.2916457783682,98.54933704849735,52.12282960982618,21.72939395904541,22.737223386764526,111.33932709693909,86.70462083816528,7046.10868823517,8898.741411257946,133.30308318138123,109.6762363910675,226.55291390419006,267.34770250320435,0.0,3282.890625,0.0,0.7751240134239197,6008.4765625,5907.328125,-0.0546875,0.0,-0.0016593893524259329,101.1484375,0.0,-64.78617191314697,634.4301888328133,46.42650743867117,-1.0078294277191162,24.634706258773804,-1852.6327230227762,23.62684679031372,-40.79478859901428,-3282.890625,-0.7751240134239197,101.1484375 +24,23,0.3828125,0.3828125,0.0,0.0,0.0012435333337634802,0.002531500067561865,6522.0703125,6498.921875,7168.0,7168.0,76.27400612831116,146.69217538833618,1368.1348377644283,708.8500100617357,93.97697018748046,48.86422865448857,23.724435329437256,25.69256353378296,120.92605495452881,98.05967283248901,7047.422495676839,8660.583657572904,144.88602995872498,123.98894739151001,241.63490891456604,290.58310437202454,0.0,3681.8671875,0.0,0.7919431924819946,6522.0703125,6498.921875,0.0,0.0,-0.0012879667337983847,23.1484375,0.0,-70.41816926002502,659.2848277026926,45.11274153299189,-1.9681282043457031,22.866382122039795,-1613.161161896065,20.897082567214966,-48.948195457458496,-3681.8671875,-0.7919431924819946,23.1484375 +25,24,0.375,0.3984375,0.0,0.0,0.0013757410924881697,0.004087383393198252,6292.6796875,6213.1484375,7168.0,7168.0,74.27027106285095,142.04829621315002,1355.6282151548562,699.8350395616865,96.51237160470447,50.46171049629547,22.379383325576782,23.928799867630005,115.69404768943787,92.55846953392029,7132.5277011233375,8805.363832224122,138.30823230743408,116.72415471076965,232.9859459400177,278.8049690723419,0.0,3851.890625,0.0,0.7779906988143921,6292.6796875,6213.1484375,-0.0234375,0.0,-0.002711642300710082,79.53125,0.0,-67.77802515029907,655.7931755931697,46.050661108409,-1.5494165420532227,23.135578155517578,-1672.8361311007848,21.58407759666443,-45.81902313232422,-3851.890625,-0.7779906988143921,79.53125 +26,25,0.3671875,0.375,0.0,0.0,0.0012049735523760319,0.0039017912931740284,6158.2265625,6037.875,7168.0,7168.0,73.27543640136719,140.2966320514679,1344.6746937171647,688.5838853534277,97.8226859098755,51.09174678812258,21.854750394821167,22.644676208496094,113.22535037994385,87.67426919937134,7151.728807044929,9060.263715385447,135.3156316280365,110.55447459220886,228.84834098815918,270.8762581348419,0.0,3513.6171875,0.0,0.7183178663253784,6158.2265625,6037.875,-0.0078125,0.0,-0.0026968177407979965,120.3515625,0.0,-67.02119565010071,656.090808363737,46.730939121752925,-0.7899258136749268,25.55108118057251,-1908.5349083405172,24.761157035827637,-42.02791714668274,-3513.6171875,-0.7183178663253784,120.3515625 +27,26,0.4296875,0.4609375,0.0,0.0,0.001379612716846168,0.005677415058016777,6149.6171875,5916.65625,7168.0,7168.0,73.96776461601257,139.75210690498352,1330.2264237778472,677.3887141777626,96.90707887700947,51.29081885594378,21.577576398849487,21.66386842727661,113.68681287765503,83.60731339454651,7124.546633844441,9331.097583753804,135.49795198440552,105.50913977622986,230.0724902153015,265.20688486099243,0.0,3373.6328125,0.0,0.6944106817245483,6149.6171875,5916.65625,-0.03125,0.0,-0.004297802341170609,232.9609375,0.0,-65.78434228897095,652.8377096000846,45.616260021065685,-0.08629202842712402,30.07949948310852,-2206.550949909363,29.98881220817566,-35.13439464569092,-3373.6328125,-0.6944106817245483,232.9609375 +28,27,0.4921875,0.5703125,0.0,0.0,0.0014020290691405535,0.005338232032954693,5968.109375,5848.9765625,7168.0,7168.0,72.61299800872803,134.10947608947754,1315.050371402131,697.8151561606371,98.71510881754823,53.44887034841241,21.81940531730652,21.995765447616577,112.84546041488647,86.14387249946594,6931.231412626842,8902.652942665824,134.89857459068298,108.37852573394775,227.84667372703552,262.49282598495483,0.0,3180.609375,0.0,0.705493688583374,5968.109375,5848.9765625,-0.078125,0.0,-0.003936202963814139,119.1328125,0.0,-61.49647808074951,617.235215241494,45.266238469135814,-0.1763601303100586,26.701587915420532,-1971.4215300389815,26.52004885673523,-34.64615225791931,-3180.609375,-0.705493688583374,119.1328125 +29,28,0.5625,0.578125,0.0,0.0,0.001420644111931324,0.0058601126074790955,6103.46875,5906.640625,7168.0,7168.0,82.25525307655334,134.27402257919312,1187.2250871213532,703.8312265074312,87.14337056782117,53.38337127550047,22.46091628074646,23.372512102127075,115.8578588962555,90.14373540878296,6905.530687533341,8595.894062811409,138.5509488582611,113.75086736679077,241.18285012245178,269.253133058548,0.0,2798.0,0.0,0.7543010115623474,6103.46875,5906.640625,-0.015625,0.0,-0.0044394684955477715,196.828125,0.0,-52.01876950263977,483.393860613922,33.75999929232069,-0.9115958213806152,25.714123487472534,-1690.363375278068,24.800081491470337,-28.07028293609619,-2798.0,-0.7543010115623474,196.828125 +30,29,0.4609375,0.4296875,0.0,0.0,0.0012942436151206493,0.0060240658931434155,5833.8984375,5835.3125,7168.0,7168.0,71.58825373649597,136.3617250919342,1303.878361715278,684.6862632241849,100.12815826440149,52.566070098976674,20.479743480682373,21.10917019844055,108.07364535331726,81.95430207252502,7104.96067278655,9371.564159259475,128.78803515434265,103.29877591133118,220.63688564300537,259.6613461971283,0.0,2941.5859375,0.0,0.6769297122955322,5833.8984375,5835.3125,0.03125,0.0,-0.004729822278022766,-1.4140625,0.0,-64.77347135543823,619.1920984910931,47.56208816542482,-0.6294267177581787,26.119343280792236,-2266.6034864729245,25.489259243011475,-39.024460554122925,-2941.5859375,-0.6769297122955322,-1.4140625 +31,30,0.4296875,0.4609375,0.0,0.0,0.0015144185163080692,0.007590139284729958,6192.765625,5992.4921875,7168.0,7168.0,73.51472210884094,138.68396258354187,1347.8150655769675,691.3551734018481,97.50427933860027,51.68586090610201,21.96708917617798,22.470194816589355,115.4867513179779,86.29063510894775,7037.776980687096,9121.905279827746,137.6874976158142,108.99810481071472,231.50495052337646,268.047385931015,0.0,3189.59375,0.0,0.8203096389770508,6192.765625,5992.4921875,-0.03125,0.0,-0.006075720768421888,200.2734375,0.0,-65.16924047470093,656.4598921751194,45.81841843249826,-0.503105640411377,29.19611620903015,-2084.1282991406506,28.689392805099487,-36.54243540763855,-3189.59375,-0.8203096389770508,200.2734375 +32,31,0.546875,0.578125,0.0,0.0,0.0013798079453408718,0.00758768804371357,6106.484375,5949.453125,7168.0,7168.0,79.49454712867737,135.5788402557373,1229.062288283088,702.1099296943704,90.16970671456495,52.86960698645356,22.875407457351685,23.53608012199402,115.6891496181488,91.52526307106018,6962.364255062701,8580.90950681277,138.798828125,115.30067896842957,238.53971219062805,271.3368785381317,0.0,2856.0859375,0.0,0.7732164859771729,6106.484375,5949.453125,-0.03125,0.0,-0.006207880098372698,157.03125,0.0,-56.08429312705994,526.9523585887176,37.30009972811139,-0.660672664642334,24.163886547088623,-1618.545251750069,23.498149156570435,-32.79716634750366,-2856.0859375,-0.7732164859771729,157.03125 +33,32,0.203125,0.2890625,0.0,0.0,0.0010425497312098742,0.0075919367372989655,6974.8046875,6789.9921875,7168.0,7168.0,82.07424473762512,151.43631505966187,1359.706389705477,717.3964511563742,87.33555846800243,47.333428558242446,23.698503732681274,25.939606189727783,125.64225888252258,100.14602279663086,7248.444974644467,8857.615861604067,149.57572197914124,126.32321190834045,251.76937675476074,297.9715495109558,0.0,3638.0390625,0.0,0.8618718385696411,6974.8046875,6789.9921875,-0.0859375,0.0,-0.006549387006089091,184.8125,0.0,-69.36207032203674,642.3099385491029,40.00212990975999,-2.241102457046509,25.496236085891724,-1609.1708869595996,23.25251007080078,-46.20217275619507,-3638.0390625,-0.8618718385696411,184.8125 +34,33,0.625,0.7109375,0.0,0.0,0.0016999590443447232,0.009697234258055687,5673.8671875,5412.578125,7168.0,7168.0,87.90410661697388,125.92486023902893,1032.737587511872,687.7216288794336,81.54340309984894,56.922834668180656,21.147135496139526,21.240893125534058,106.12948369979858,82.56478905677795,7006.5261233472975,8601.172583528834,127.51026320457458,104.03811550140381,235.54982805252075,250.2824251651764,0.0,2697.34375,0.0,0.7204962968826294,5673.8671875,5412.578125,-0.0859375,0.0,-0.007997275213710964,261.2890625,0.0,-38.020753622055054,345.01595863243847,24.62056843166829,-0.09375762939453125,23.56469464302063,-1594.646460181537,23.472147703170776,-14.73259711265564,-2697.34375,-0.7204962968826294,261.2890625 +35,34,0.4765625,0.578125,0.0,0.0,0.0014655782142654061,0.012641005218029022,6279.03125,6044.609375,7168.0,7168.0,74.02645182609558,136.5413224697113,1357.143257872378,708.3112148811495,96.83025220281,52.4969281851658,23.23618507385254,25.98934054374695,118.91491317749023,96.7929196357727,6934.874507868727,8209.836039559985,142.38428735733032,123.01590299606323,236.66528248786926,279.69403314590454,0.0,3308.703125,0.0,0.8351297378540039,6279.03125,6044.609375,-0.1015625,0.0,-0.011175427003763616,234.421875,0.0,-62.51487064361572,648.8320429912285,44.3333240176442,-2.753155469894409,22.12199354171753,-1274.9615316912577,19.36838436126709,-43.02875065803528,-3308.703125,-0.8351297378540039,234.421875 +36,35,0.3125,0.421875,0.0,0.0,0.0018182576168328524,0.012619839049875736,5783.6171875,5619.8984375,7168.0,7168.0,73.97461414337158,133.1270182132721,1250.9409622691728,675.4329527305193,96.89810596521079,53.843315175261594,20.23202157020569,21.195914268493652,106.95720911026001,81.77313208580017,7067.639538170092,8988.02554399929,127.4229257106781,103.2038881778717,221.97860145568848,256.49823212623596,0.0,2713.0703125,0.0,0.7314876317977905,5783.6171875,5619.8984375,-0.109375,0.0,-0.010801581433042884,163.71875,0.0,-59.15240406990051,575.5080095386535,43.054790789949195,-0.9638926982879639,25.18407702445984,-1920.3860058291984,24.219037532806396,-34.519630670547485,-2713.0703125,-0.7314876317977905,163.71875 +37,36,0.171875,0.2421875,0.0,0.0,0.0011753416620194912,0.01196231134235859,6926.6484375,6835.3359375,7168.0,7168.0,82.41997861862183,153.31985354423523,1344.654255648643,713.315154377227,86.96920479885291,46.75193612764519,23.702624797821045,25.03836441040039,127.30228090286255,97.23776769638062,7133.281458585246,9218.588838844411,151.24044394493103,122.51350331306458,253.81723022460938,296.27348470687866,0.0,3305.4453125,0.0,0.8874380588531494,6926.6484375,6835.3359375,-0.0703125,0.0,-0.010786969680339098,91.3125,0.0,-70.8998749256134,631.3391012714159,40.21726867120773,-1.3357396125793457,30.064513206481934,-2085.3073802591653,28.726940631866455,-42.45625448226929,-3305.4453125,-0.8874380588531494,91.3125 +38,37,0.46875,0.484375,0.0,0.0,0.002175199566408992,0.01271863840520382,5531.984375,5449.703125,7168.0,7168.0,69.15472793579102,128.8740575313568,1279.9088745194927,676.5927268083742,103.65162605592731,55.620193367900505,19.953352212905884,20.447848081588745,104.11892080307007,78.47632503509521,6977.847968422062,9123.694307548194,124.30545496940613,99.15881633758545,213.8094642162323,248.16125392913818,0.0,2957.1171875,0.0,0.7147747278213501,5531.984375,5449.703125,-0.015625,0.0,-0.010543438838794827,82.28125,0.0,-59.719329595565796,603.3161477111184,48.031432688026804,-0.49449586868286133,25.642595767974854,-2145.846339126132,25.14663863182068,-34.351789712905884,-2957.1171875,-0.7147747278213501,82.28125 +39,38,0.453125,0.5,0.0,0.0,0.0017289501847699285,0.01888348162174225,6202.2109375,5894.1796875,7168.0,7168.0,73.09492945671082,134.7831907272339,1357.6232405938692,699.6931478707354,98.06425771633205,53.181705829372795,22.60435700416565,22.171786308288574,117.55799674987793,85.68691802024841,6916.16923117435,9028.484369308131,140.39578580856323,108.09402465820312,233.8988699913025,263.09600353240967,0.0,3268.4140625,0.0,0.8372789025306702,6202.2109375,5894.1796875,-0.046875,0.0,-0.01715453143697232,308.03125,0.0,-61.68826127052307,657.9300927231338,44.882551886959256,0.4325706958770752,31.871078729629517,-2112.315138133781,32.30176115036011,-29.197133541107178,-3268.4140625,-0.8372789025306702,308.03125 +40,39,0.5859375,0.5859375,0.0,0.0,0.0016666848678141832,0.014175266027450562,5938.1328125,5836.0625,7168.0,7168.0,83.75661611557007,133.63771677017212,1134.359641140492,698.7323807738214,85.58129891624756,53.637552131539366,22.007693767547607,22.83385181427002,113.20562243461609,89.55574488639832,6901.998179916193,8578.790796442652,135.44507384300232,112.62217974662781,239.488130569458,266.14931321144104,0.0,2796.7109375,0.0,0.8285865783691406,5938.1328125,5836.0625,0.0,0.0,-0.012508581159636378,102.0703125,0.0,-49.88110065460205,435.6272603666705,31.943746784708196,-0.8261580467224121,23.649877548217773,-1676.792616526459,22.82289409637451,-26.661182641983032,-2796.7109375,-0.8285865783691406,102.0703125 +41,40,0.3671875,0.390625,0.0,0.0,0.0017815136816352606,0.019767679274082184,5913.265625,5856.171875,7168.0,7168.0,72.71093106269836,137.42293167114258,1301.2108168222494,681.8276168363521,98.58215120115929,52.16014469225013,20.80664873123169,22.04793667793274,109.64676928520203,84.82650899887085,7046.938136318444,9022.721894758017,130.6865372657776,107.10970950126648,223.39451718330383,264.52131605148315,0.0,3273.109375,0.0,0.7956519722938538,5913.265625,5856.171875,-0.0234375,0.0,-0.017986165592446923,57.09375,0.0,-64.71200060844421,619.3831999858972,46.42200650890916,-1.2412879467010498,24.820260286331177,-1975.783758439573,23.57682776451111,-41.12679886817932,-3273.109375,-0.7956519722938538,57.09375 +42,41,0.3125,0.453125,0.0,0.0,0.0013291820650920272,0.015728659927845,6742.4921875,6376.296875,7168.0,7168.0,78.49755501747131,142.9364354610443,1374.3087281634316,713.7490848357177,91.31494603118031,50.14816534971978,23.46744441986084,24.952731132507324,124.62277817726135,96.82289910316467,7077.349846473968,8625.294302643988,148.32533049583435,122.01162934303284,247.14692068099976,285.08054757118225,0.0,3059.09375,0.0,0.8898141384124756,6742.4921875,6376.296875,-0.140625,0.0,-0.014399477862752974,366.1953125,0.0,-64.438880443573,660.559643327714,41.16678068146053,-1.4852867126464844,27.79987907409668,-1547.9444561700202,26.313701152801514,-37.933626890182495,-3059.09375,-0.8898141384124756,366.1953125 +43,42,0.6875,0.6875,0.0,0.0,0.002072055358439684,0.019971981644630432,5444.2578125,5253.6875,7168.0,7168.0,86.82285451889038,123.83264064788818,1003.2856611624944,678.8113340731987,82.55890732595564,57.88457681672027,20.29501748085022,21.31100082397461,102.38407468795776,79.7760705947876,6960.77981045448,8627.649806118312,122.91022229194641,101.32143092155457,229.93872928619385,245.1762089729309,0.0,2529.8359375,0.0,0.8514864444732666,5444.2578125,5253.6875,0.0,0.0,-0.017899926286190748,190.5703125,0.0,-37.0097861289978,324.4743270892957,24.674330509235375,-1.0159833431243896,22.608004093170166,-1666.869995663832,21.588791370391846,-15.23747968673706,-2529.8359375,-0.8514864444732666,190.5703125 +44,43,0.0625,0.1484375,0.0,0.0,0.0013166103744879365,0.01697881892323494,7129.2421875,7017.0078125,7168.0,7168.0,85.38206958770752,155.52499222755432,1335.9699003644482,721.8912111291447,83.95205263368292,46.089055510205306,23.716686010360718,25.551151752471924,128.76038551330566,100.15981340408325,7248.308525012582,9174.627715136476,152.71362829208374,125.94951915740967,258.1519515514374,301.4726793766022,0.0,3311.140625,0.0,0.9239728450775146,7129.2421875,7017.0078125,-0.0859375,0.0,-0.015662208548747003,112.234375,0.0,-70.1429226398468,614.0786892353035,37.86299712347762,-1.834465742111206,28.600572109222412,-1926.3191901238933,26.764109134674072,-43.320727825164795,-3311.140625,-0.9239728450775146,112.234375 +45,44,0.5625,0.609375,0.0,0.0,0.0019004116766154766,0.023962941020727158,6160.640625,5761.9296875,7168.0,7168.0,86.04612708091736,132.76705408096313,1145.5512681855514,694.3806627190867,83.30415607502297,53.98929764329076,23.286268949508667,23.658780574798584,118.85289597511292,91.48754858970642,6773.028064613259,8241.110529491556,142.37439393997192,115.38058137893677,248.44135975837708,268.24540400505066,0.0,2790.125,0.0,0.893954873085022,6160.640625,5761.9296875,-0.046875,0.0,-0.02206252934411168,398.7109375,0.0,-46.720927000045776,451.17060546646474,29.314858431732212,-0.372511625289917,27.365347385406494,-1468.0824648782973,26.993812561035156,-19.804044246673584,-2790.125,-0.893954873085022,398.7109375 +46,45,0.390625,0.4375,0.0,0.0,0.0017820914508774877,0.025402288883924484,5934.2421875,5707.0625,7168.0,7168.0,71.88719034194946,134.75031995773315,1320.7899007925696,677.6458863225108,99.7117840592129,53.19467888646477,21.085479021072388,21.955585956573486,110.70132517814636,83.42964720726013,7022.562726769133,8969.569272431008,132.02196741104126,105.61977291107178,224.22100377082825,260.3512599468231,0.0,2663.2890625,0.0,0.8057901859283447,5934.2421875,5707.0625,-0.046875,0.0,-0.023620197433046997,227.1796875,0.0,-62.86312961578369,643.1440144700588,46.51710517274813,-0.8701069355010986,27.27167797088623,-1947.006545661875,26.402194499969482,-36.13025617599487,-2663.2890625,-0.8057901859283447,227.1796875 +47,46,0.3828125,0.5,0.0,0.0,0.001665993593633175,0.022048229351639748,6776.2890625,6448.484375,7168.0,7168.0,76.8701012134552,143.9238224029541,1410.439472415086,716.8775000370076,93.24821857715112,49.804124712107935,23.662214994430542,25.81503200531006,124.54255270957947,99.57074761390686,7121.268841085686,8485.845695126513,148.4392740726471,125.62244772911072,245.77935934066772,289.44649171829224,0.0,3009.8046875,0.0,0.9592458009719849,6776.2890625,6448.484375,-0.1171875,0.0,-0.020382235758006573,327.8046875,0.0,-67.0537211894989,693.5619723780783,43.444093865043186,-2.1528170108795166,24.971805095672607,-1364.5768540408262,22.816826343536377,-43.66713237762451,-3009.8046875,-0.9592458009719849,327.8046875 +48,47,0.4921875,0.65625,0.0,0.0,0.0019905450753867626,0.022940408438444138,6449.734375,5982.546875,7168.0,7168.0,77.81224775314331,134.67485189437866,1326.2147409927675,710.7544478687887,92.11917412719183,53.22448771372432,23.698041915893555,25.712996244430542,120.69192409515381,97.7624409198761,6999.755835642711,8029.811782659763,144.62469601631165,123.7444977760315,242.6808226108551,278.3009946346283,0.0,2787.5390625,0.0,0.9266676902770996,6449.734375,5982.546875,-0.1640625,0.0,-0.020949863363057375,467.1875,0.0,-56.86260414123535,615.4602931239788,38.894686413467504,-2.0149543285369873,22.92948317527771,-1030.055947017052,20.88019824028015,-35.62017202377319,-2787.5390625,-0.9266676902770996,467.1875 +49,48,0.1640625,0.2421875,0.0,0.0,0.0013237509410828352,0.026233401149511337,7037.3203125,6847.8671875,7168.0,7168.0,83.26385235786438,153.58689165115356,1352.2930036441564,713.3803791593114,86.08777755312414,46.670649577835654,23.704346656799316,26.038492679595947,126.54122805595398,100.12424302101135,7277.256702399068,8955.103908369536,150.4801230430603,126.4011070728302,253.89857172966003,299.85666251182556,0.0,3407.5,0.0,0.9539734125137329,7037.3203125,6847.8671875,-0.078125,0.0,-0.024909650208428502,189.453125,0.0,-70.32303929328918,638.9126244848451,39.41712797528849,-2.334146022796631,26.416985034942627,-1677.8472059704682,24.079015970230103,-45.95809078216553,-3407.5,-0.9539734125137329,189.453125 +50,49,0.2578125,0.3359375,0.0,0.0,0.0016027060337364674,0.025359654799103737,6422.0390625,6334.546875,7168.0,7168.0,78.13214111328125,145.48700881004333,1315.1133904166572,696.6446752117387,91.74201420651393,49.269003869334995,22.073874473571777,23.30060648918152,116.4170401096344,89.46053218841553,7225.3769655013575,9277.364885914163,138.72740030288696,112.9991807937622,237.03785157203674,278.4230058193207,0.0,2836.515625,0.0,0.8862186670303345,6422.0390625,6334.546875,-0.078125,0.0,-0.02375694876536727,87.4921875,0.0,-67.35486769676208,618.4687152049186,42.473010337178934,-1.2267320156097412,26.956507921218872,-2051.9879204128056,25.728219509124756,-41.385154247283936,-2836.515625,-0.8862186670303345,87.4921875 +51,50,0.3671875,0.359375,0.0,0.0,0.0024111364036798477,0.02409507893025875,5584.3359375,5458.3671875,7168.0,7168.0,84.73390865325928,131.69740104675293,1054.4701220573659,663.140459157552,84.59423286292936,54.42780148300224,20.38700222969055,20.197001695632935,104.0275366306305,78.19530892372131,7046.701515223192,9168.414446694682,124.64582300186157,98.62713289260864,229.92590856552124,250.558922290802,0.0,2279.21875,0.0,0.8059177398681641,5584.3359375,5458.3671875,0.0078125,0.0,-0.021683942526578903,125.96875,0.0,-46.96349239349365,391.3296628998139,30.16643137992712,0.1900005340576172,25.83222770690918,-2121.7129314714903,26.01869010925293,-20.63301372528076,-2279.21875,-0.8059177398681641,125.96875 +52,51,0.3671875,0.3671875,0.0,0.0,0.001672852784395218,0.030810344964265823,6287.40625,6230.5390625,7168.0,7168.0,74.65390872955322,142.26130294799805,1347.5315855789356,700.7430898931103,96.01640586519491,50.386154572337766,22.261980295181274,22.8878173828125,115.54599738121033,89.7183358669281,7149.118262181617,9126.038641804718,138.04305386543274,112.84406399726868,232.80843806266785,275.08760690689087,0.0,2927.390625,0.0,0.951687753200531,6287.40625,6230.5390625,0.0,0.0,-0.029137492179870605,56.8671875,0.0,-67.60739421844482,646.7884956858253,45.630251292857146,-0.6258370876312256,25.827661514282227,-1976.9203796231013,25.198989868164062,-42.27916884422302,-2927.390625,-0.951687753200531,56.8671875 +53,52,0.3984375,0.3984375,0.0,0.0,0.0016998036298900843,0.029189204797148705,6302.234375,6137.21875,7168.0,7168.0,74.65992617607117,141.21662139892578,1350.6007193497373,695.3536986457527,96.00866712747133,50.75889742292423,23.05222773551941,24.09495735168457,118.08777475357056,92.23829197883606,6973.778629655309,8699.16368555598,141.37536883354187,116.57039046287537,236.27747344970703,277.8569760322571,0.0,2779.3671875,0.0,0.9398800134658813,6302.234375,6137.21875,0.0,0.0,-0.02748940116725862,165.015625,0.0,-66.55669522285461,655.2470207039846,45.2497697045471,-1.0427296161651611,25.849482774734497,-1725.3850559006705,24.804978370666504,-41.57950258255005,-2779.3671875,-0.9398800134658813,165.015625 +54,53,0.15625,0.1640625,0.0,0.0,0.0016851560212671757,0.031399063766002655,6663.546875,6661.0390625,7168.0,7168.0,80.51664781570435,151.9789125919342,1324.1578343405022,701.2592943480254,89.02506741719965,47.16443799835701,22.761798858642578,24.56570315361023,120.4573323726654,94.0684142112732,7223.586832456318,9246.600012268105,143.45497250556946,118.87002515792847,244.6762819290161,290.96141839027405,0.0,2961.1640625,0.0,0.920147180557251,6663.546875,6661.0390625,-0.0078125,0.0,-0.02971390774473548,2.5078125,0.0,-71.46226477622986,622.8985399924768,41.86062941884264,-1.8039042949676514,26.388918161392212,-2023.0131798117873,24.58494734764099,-46.285136461257935,-2961.1640625,-0.920147180557251,2.5078125 +55,54,0.2109375,0.2578125,0.0,0.0,0.0014902394032105803,0.030728641897439957,6961.5859375,6845.515625,7168.0,7168.0,81.37171149253845,154.87792038917542,1368.846408130591,707.1908618399492,88.0895813609289,46.28161316983295,23.702333450317383,26.149888277053833,126.24105715751648,100.23993515968323,7193.689758688216,8911.438326221907,150.17978525161743,126.6276547908783,251.88689637184143,301.4969186782837,0.0,2929.859375,0.0,1.0271353721618652,6961.5859375,6845.515625,-0.046875,0.0,-0.029238402494229376,116.0703125,0.0,-73.50620889663696,661.6555462906418,41.80796819109595,-2.44755482673645,26.001121997833252,-1717.7485675336911,23.552130460739136,-49.61002230644226,-2929.859375,-1.0271353721618652,116.0703125 +56,55,0.4375,0.4765625,0.0,0.0,0.0022621690295636654,0.030439790338277817,5428.53125,5337.828125,7168.0,7168.0,94.60689282417297,129.42915320396423,918.0779265356797,659.8609964279991,75.76614965382846,55.38164951681422,19.328219413757324,19.938042879104614,99.84437704086304,75.4192545413971,7137.387413497955,9294.947348163143,119.40588903427124,95.59052181243896,234.1366264820099,245.06268954277039,0.0,2288.5,0.0,0.8165364265441895,5428.53125,5337.828125,-0.0390625,0.0,-0.02817762130871415,90.703125,0.0,-34.82226037979126,258.2169301076806,20.384500137014243,-0.60982346534729,24.425122499465942,-2157.559934665188,23.815367221832275,-10.926063060760498,-2288.5,-0.8165364265441895,90.703125 +57,56,0.4296875,0.5390625,0.0,0.0,0.0024752996396273375,0.03197157382965088,5717.109375,5367.6640625,7168.0,7168.0,73.06116223335266,127.88423132896423,1252.0160808260716,671.5654002648614,98.10958080718547,56.050694643980975,20.794092416763306,21.261232137680054,106.96289706230164,81.64930891990662,7009.365121845055,8634.647486013362,127.98913264274597,103.1429672241211,221.65795922279358,251.43477034568787,0.0,2414.03125,0.0,0.8752177953720093,5717.109375,5367.6640625,-0.109375,0.0,-0.02949627419002354,349.4453125,0.0,-54.82306909561157,580.4506805612102,42.05888616320449,-0.46713972091674805,25.31358814239502,-1625.2823641683071,24.846165418624878,-29.776811122894287,-2414.03125,-0.8752177953720093,349.4453125 +58,57,0.5078125,0.484375,0.0,0.0,0.002101929159834981,0.03409340977668762,5897.59375,5882.8671875,7168.0,7168.0,81.07691764831543,137.01941299438477,1163.8515959537172,686.9528407909426,88.40987309227009,52.31375498808883,21.174758195877075,22.925553560256958,109.34984493255615,87.69443130493164,7049.191523549252,8768.435903600612,130.75818061828613,110.85466742515564,231.99302554130554,267.9116585254669,0.0,2392.7421875,0.0,0.941716194152832,5897.59375,5882.8671875,0.0234375,0.0,-0.03199148061685264,14.7265625,0.0,-55.942495346069336,476.8987551627746,36.09611810418126,-1.7507953643798828,21.65541362762451,-1719.2443800513593,19.903513193130493,-35.91863298416138,-2392.7421875,-0.941716194152832,14.7265625 +59,58,0.3984375,0.421875,0.0,0.0,0.0017224873881787062,0.029396476224064827,6194.2890625,6059.34375,7168.0,7168.0,74.14904189109802,138.98903393745422,1336.6136968507276,697.5334474490107,96.67016345979995,51.57241400228479,21.696770668029785,22.51665186882019,113.35883283615112,89.25593090057373,7164.267483009967,8905.402610000572,135.2908034324646,112.00896739959717,229.59352922439575,270.9420094490051,0.0,2435.140625,0.0,0.9517225027084351,6194.2890625,6059.34375,-0.0234375,0.0,-0.02767398883588612,134.9453125,0.0,-64.8399920463562,639.0802494017169,45.097749457515164,-0.8198812007904053,24.102901935577393,-1741.1351269906054,23.28183603286743,-41.348480224609375,-2435.140625,-0.9517225027084351,134.9453125 +60,59,0.7421875,0.7890625,0.0,0.0,0.002289219293743372,0.03898944705724716,5359.7890625,4910.640625,7168.0,7168.0,91.3830828666687,118.70761251449585,938.4299840827441,661.8804669364021,78.43902585841164,60.38365904397823,21.099823713302612,19.33893871307373,104.46255993843079,73.83624744415283,6767.180513445678,8795.49032460247,125.79222059249878,93.40821814537048,237.32350873947144,232.16530394554138,0.0,2186.3046875,0.0,0.815165638923645,5359.7890625,4910.640625,-0.046875,0.0,-0.03670022776350379,449.1484375,0.0,-27.32452964782715,276.54951714634205,18.055366814433413,1.7608850002288818,30.626312494277954,-2028.3098111567915,32.384002447128296,5.158204793930054,-2186.3046875,-0.815165638923645,449.1484375 +61,60,0.2890625,0.3359375,0.0,0.0,0.002105850726366043,0.033121757209300995,5863.6953125,5641.9375,7168.0,7168.0,72.88557028770447,133.924498796463,1287.2112357722328,674.043963660397,98.345941064952,53.522694237548336,20.50202441215515,20.85079574584961,107.43855237960815,80.18392539024353,7168.311401654506,9250.831714585222,128.17388558387756,101.26953625679016,221.50826287269592,255.11287760734558,0.0,2388.6640625,0.0,0.8576261401176453,5863.6953125,5641.9375,-0.046875,0.0,-0.031015906482934952,221.7578125,0.0,-61.038928508758545,613.1672721118358,44.823246827403665,-0.348771333694458,27.254626989364624,-2082.5203129307156,26.904349327087402,-33.60461473464966,-2388.6640625,-0.8576261401176453,221.7578125 +62,61,0.6015625,0.625,0.0,0.0,0.0030279222410172224,0.04049159586429596,5630.1953125,5217.59375,7168.0,7168.0,89.06223249435425,125.39767003059387,1011.4626871239778,665.7340601275336,80.48304875418867,57.16214661924092,20.102020740509033,20.05941653251648,103.79407286643982,79.27688002586365,7128.663319263958,8667.091840342826,124.13001656532288,99.56886196136475,233.29967713356018,245.48942184448242,0.0,2292.6015625,0.0,0.8180185556411743,5630.1953125,5217.59375,-0.0234375,0.0,-0.03746367362327874,412.6015625,0.0,-36.335437536239624,345.7286269964443,23.32090213494775,0.04260420799255371,24.517192840576172,-1538.4285210788676,24.56115460395813,-12.189744710922241,-2292.6015625,-0.8180185556411743,412.6015625 +63,62,0.2890625,0.375,0.0,0.0,0.0024309256114065647,0.03630291670560837,6128.2109375,5973.1875,7168.0,7168.0,75.67534637451172,138.77247190475464,1295.6845220734235,688.6884602415555,94.72041217394758,51.65289557513754,21.133028745651245,22.42874264717102,111.74169826507568,86.10330939292908,7187.674900865764,9097.443588670325,133.10800957679749,108.76719951629639,228.95385575294495,267.8028745651245,0.0,2594.703125,0.0,0.9090120792388916,6128.2109375,5973.1875,-0.0859375,0.0,-0.0338719910942018,155.0234375,0.0,-63.09712553024292,606.996061831868,43.06751659881004,-1.2957139015197754,25.638388872146606,-1909.7686878045606,24.3408100605011,-38.849018812179565,-2594.703125,-0.9090120792388916,155.0234375 +64,63,0.46875,0.5625,0.0,0.0,0.0024836231023073196,0.04656257480382919,5913.1640625,5597.859375,7168.0,7168.0,73.71829628944397,130.65686964988708,1283.4076445354272,685.5035654841851,97.23501980913815,54.86125619883313,21.6168372631073,21.9595685005188,111.19676518440247,84.01827645301819,6997.658598337951,8780.922808057645,133.04352688789368,106.21284174919128,227.0935251712799,256.869131565094,0.0,2558.8046875,0.0,0.9484083652496338,5913.1640625,5597.859375,-0.09375,0.0,-0.044078951701521873,315.3046875,0.0,-56.938573360443115,597.9040790512421,42.37376361030502,-0.342731237411499,27.178488731384277,-1783.264209719694,26.830685138702393,-29.775606393814087,-2558.8046875,-0.9484083652496338,315.3046875 +65,64,0.5078125,0.53125,0.0,0.0,0.0025276949163526297,0.04560329392552376,6139.71875,5724.7265625,7168.0,7168.0,75.04967474937439,133.44054698944092,1308.9397166350661,686.4152393443644,95.51007414672043,53.71680618610774,21.864505529403687,21.34281849861145,112.92568445205688,82.85911011695862,7116.8574616096685,9058.23148402827,135.02137088775635,104.43701481819153,230.42391681671143,257.9188423156738,0.0,2644.4140625,0.0,0.9482257962226868,6139.71875,5724.7265625,-0.0234375,0.0,-0.04307559900917113,414.9921875,0.0,-58.39087224006653,622.5244772907017,41.79326796061269,0.5216870307922363,30.066574335098267,-1941.3740224186022,30.58435606956482,-27.494925498962402,-2644.4140625,-0.9482257962226868,414.9921875 +66,65,0.453125,0.4765625,0.0,0.0,0.0027204290963709354,0.04304514080286026,5797.8203125,5534.0078125,7168.0,7168.0,71.85756778717041,132.04652333259583,1290.9583201417856,670.5524898749287,99.75288923263263,54.28389797090986,21.000763654708862,20.559786319732666,107.75184082984924,79.82262110710144,7075.749185612003,9128.452434834602,128.98508214950562,100.61790728569031,220.89747166633606,252.5872642993927,0.0,2401.3671875,0.0,0.9617340564727783,5797.8203125,5534.0078125,-0.0234375,0.0,-0.040324711706489325,263.8125,0.0,-60.188955545425415,620.4058302668569,45.46899126172276,0.4409773349761963,27.929219722747803,-2052.703249222599,28.367174863815308,-31.68979263305664,-2401.3671875,-0.9617340564727783,263.8125 +67,66,0.421875,0.5390625,0.0,0.0,0.0022645341232419014,0.03353730961680412,6413.3984375,5907.9296875,7168.0,7168.0,73.8682177066803,135.32429003715515,1389.1546078377905,698.5211226605833,97.03767361036192,52.969056760112515,23.630884408950806,24.079463005065918,120.8043920993805,93.2121250629425,6960.831352126909,8327.232100715042,144.66837000846863,117.52693486213684,238.71158623695374,273.4303925037384,0.0,2251.40625,0.0,0.9707183837890625,6413.3984375,5907.9296875,-0.1171875,0.0,-0.03127277549356222,505.46875,0.0,-61.45607233047485,690.6334851772073,44.068616850249406,-0.4485785961151123,27.59226703643799,-1366.400748588133,27.141435146331787,-34.71880626678467,-2251.40625,-0.9707183837890625,505.46875 +68,67,0.609375,0.59375,0.0,0.0,0.0035136621445417404,0.040111981332302094,5106.8203125,4940.8828125,7168.0,7168.0,86.67696022987366,121.00642013549805,942.6856316061547,653.3052123307046,82.69787012592433,59.23652639234816,18.352999210357666,17.759806394577026,94.8217842578888,67.80481433868408,7056.026262702776,9554.262574396611,113.4082658290863,85.79796147346497,220.0423128604889,226.8457760810852,0.0,1864.5625,0.0,0.7557700872421265,5106.8203125,4940.8828125,0.015625,0.0,-0.03659831918776035,165.9375,0.0,-34.32945990562439,289.3804192754501,23.461343733576165,0.5931928157806396,27.016969919204712,-2498.236311693835,27.610304355621338,-6.8034632205963135,-1864.5625,-0.7557700872421265,165.9375 +69,68,0.34375,0.546875,0.0,0.0,0.0025010015815496445,0.04340093582868576,5987.703125,5492.453125,7168.0,7168.0,74.10842275619507,129.929372549057,1292.744420093482,676.3616900160102,96.72314877866961,55.16843389121733,21.162949562072754,19.772753953933716,109.8695764541626,76.7754054069519,7144.998873528089,9399.181888717943,131.265784740448,96.78300666809082,225.88799214363098,246.8119032382965,0.0,2284.9765625,0.0,0.868461012840271,5987.703125,5492.453125,-0.203125,0.0,-0.040899934247136116,495.25,0.0,-55.82094979286194,616.3827300774717,41.55471488745228,1.390195608139038,33.09417104721069,-2254.183015189855,34.48277807235718,-20.923911094665527,-2284.9765625,-0.868461012840271,495.25 +70,69,0.4375,0.453125,0.0,0.0,0.002600564621388912,0.05367598682641983,5960.5390625,5643.9296875,7168.0,7168.0,72.17984533309937,132.59180092811584,1321.2639145995372,681.0592688831293,99.30750013276885,54.060657973007736,21.434947967529297,20.68468475341797,110.73090124130249,79.90657830238342,7075.504590111329,9297.745139186363,132.40033555030823,100.82724380493164,224.8952305316925,253.50887322425842,0.0,2552.5546875,0.0,1.0244548320770264,5960.5390625,5643.9296875,-0.015625,0.0,-0.05107542220503092,316.609375,0.0,-60.41195559501648,640.2046457164079,45.246842159761115,0.7502632141113281,30.824322938919067,-2222.2405490750343,31.573091745376587,-28.613642692565918,-2552.5546875,-1.0244548320770264,316.609375 +71,70,0.3828125,0.4140625,0.0,0.0,0.002795071341097355,0.05180183798074722,5642.859375,5503.453125,7168.0,7168.0,72.12057948112488,131.8772692680359,1251.8722207941942,667.7060458465425,99.38910712546316,54.35356706872125,20.117834091186523,19.93895673751831,104.05524897575378,78.03302335739136,7131.576804673388,9281.121874299388,124.40613436698914,98.20690965652466,216.77461314201355,250.3336215019226,0.0,2413.703125,0.0,0.923332691192627,5642.859375,5503.453125,-0.03125,0.0,-0.04900676663964987,139.40625,0.0,-59.75668978691101,584.1661749476517,45.03554005674191,0.1788773536682129,26.022225618362427,-2149.5450696259995,26.199224710464478,-33.55900835990906,-2413.703125,-0.923332691192627,139.40625 +72,71,0.390625,0.3984375,0.0,0.0,0.0024853572249412537,0.04350922256708145,6026.734375,5954.8515625,7168.0,7168.0,74.09526586532593,140.18451142311096,1301.4023078784162,679.6587157366405,96.7403236399531,51.132610352118235,21.438775062561035,21.881216764450073,111.20845365524292,84.43753337860107,7126.346729510856,9276.78685837255,132.88150906562805,106.56365633010864,227.3091812133789,266.71507453918457,0.0,2365.078125,0.0,0.9633471965789795,6026.734375,5954.8515625,-0.0078125,0.0,-0.0410238653421402,71.8828125,0.0,-66.08924555778503,621.7435921417757,45.60771328783487,-0.4424417018890381,26.770920276641846,-2150.4401288616946,26.31785273551941,-39.405893325805664,-2365.078125,-0.9633471965789795,71.8828125 +73,72,0.5078125,0.5859375,0.0,0.0,0.00241178460419178,0.04660977050662041,5866.2890625,5617.671875,7168.0,7168.0,76.18635654449463,130.9425172805786,1231.9873171147538,686.429067247903,94.08508721392549,54.741577822585185,21.177287340164185,21.117196083068848,109.0224540233612,80.71910667419434,7048.428756110555,9125.646087402667,130.43294215202332,102.07062935829163,227.20173048973083,252.93299198150635,0.0,2229.7109375,0.0,0.9853517413139343,5866.2890625,5617.671875,-0.078125,0.0,-0.04419798590242863,248.6171875,0.0,-54.756160736083984,545.5582498668507,39.343509391340305,0.060091257095336914,28.30334734916687,-2077.2173312921122,28.36231279373169,-25.731261491775513,-2229.7109375,-0.9853517413139343,248.6171875 +74,73,0.3046875,0.34375,0.0,0.0,0.0025853195693343878,0.04785541445016861,6378.0234375,6217.3125,7168.0,7168.0,77.20556998252869,143.44839668273926,1321.7747763936352,693.4688870731017,92.84304230409917,49.96918868220787,21.954601287841797,22.897759199142456,115.88160824775696,89.14224553108215,7220.084469411003,9155.075633757067,138.07057666778564,112.27529788017273,235.48641800880432,275.77930665016174,0.0,2388.03125,0.0,0.9976344704627991,6378.0234375,6217.3125,-0.0390625,0.0,-0.04527009488083422,160.7109375,0.0,-66.24282670021057,628.3058893205335,42.8738536218913,-0.9431579113006592,26.739362716674805,-1934.9911643460646,25.795278787612915,-40.29288864135742,-2388.03125,-0.9976344704627991,160.7109375 +75,74,0.5625,0.65625,0.0,0.0,0.0033046130556613207,0.052180588245391846,5625.8203125,5257.171875,7168.0,7168.0,82.75330519676208,125.459956407547,1087.7284573223544,670.4509742276629,86.61889676740627,57.13376765982051,20.20428991317749,20.689284563064575,103.78132653236389,78.34434628486633,7109.034203469388,8814.90538562834,124.2177221775055,99.2669460773468,227.09605526924133,244.66874146461487,0.0,2012.59375,0.0,0.9125559329986572,5625.8203125,5257.171875,-0.09375,0.0,-0.048875975189730525,368.6484375,0.0,-42.70665121078491,417.27748309469155,29.485129107585756,-0.48499464988708496,25.43698024749756,-1705.8711821589532,24.95077610015869,-17.572686195373535,-2012.59375,-0.9125559329986572,368.6484375 +76,75,0.484375,0.484375,0.0,0.0,0.003323609009385109,0.0492306724190712,5536.9609375,5342.0546875,7168.0,7168.0,94.52343940734863,128.67099046707153,937.2423978164356,664.2746332311287,75.83304252302452,55.70797251175569,20.387977361679077,19.753803730010986,104.08518409729004,76.65125107765198,7028.349004179227,9218.362258486504,124.70374488830566,96.63997149467468,239.35156106948853,245.34655594825745,0.0,2271.4140625,0.0,0.8546422719955444,5536.9609375,5342.0546875,0.0,0.0,-0.04590706340968609,194.90625,0.0,-34.1475510597229,272.9677645853069,20.125070011268825,0.6341736316680908,27.43393301963806,-2190.013254307277,28.06377339363098,-5.994994878768921,-2271.4140625,-0.8546422719955444,194.90625 +77,76,0.3984375,0.421875,0.0,0.0,0.00257853209041059,0.04478643089532852,5945.390625,5680.2734375,7168.0,7168.0,84.27521228790283,132.06894063949585,1128.7571685376195,688.1585826305976,85.05466560573613,54.274683852929876,22.796850204467773,21.076946020126343,110.67542767524719,81.06307458877563,7045.484407686599,9200.57626463715,133.70230770111084,102.37499165534973,238.11534595489502,254.66926741600037,0.0,2042.7265625,0.0,0.9095518589019775,5945.390625,5680.2734375,-0.0234375,0.0,-0.04220789880491793,265.1171875,0.0,-47.79372835159302,440.59858590702186,30.779981752806258,1.7199041843414307,29.612353086471558,-2155.09185695055,31.32731604576111,-16.553921461105347,-2042.7265625,-0.9095518589019775,265.1171875 +78,77,0.640625,0.59375,0.0,0.0,0.002625118475407362,0.04508449137210846,5611.765625,5295.9453125,7168.0,7168.0,83.37149500846863,128.283531665802,1076.965814165616,660.5300298462925,85.97662785430315,55.87622906012382,20.193556785583496,19.651464700698853,104.09139013290405,75.3624005317688,7096.398646005782,9265.217072081658,124.51730108261108,95.24972796440125,227.9469985961914,243.71069598197937,0.0,1887.5,0.0,0.889992356300354,5611.765625,5295.9453125,0.046875,0.0,-0.0424593728967011,315.8203125,0.0,-44.912036657333374,416.4357843193235,30.100398794179327,0.5420920848846436,28.728989601135254,-2168.8184260758753,29.26757311820984,-15.763697385787964,-1887.5,-0.889992356300354,315.8203125 +79,78,0.25,0.296875,0.0,0.0,0.0028637656942009926,0.06773677468299866,6373.21875,6313.3828125,7168.0,7168.0,77.67335224151611,146.29784488677979,1312.8247598086364,690.468988645561,92.28390166182027,48.995937059410075,23.04257082939148,22.87431764602661,116.76722168922424,89.00555682182312,7147.862113405269,9291.296291258468,140.04475569725037,112.11530613899231,237.87781381607056,278.8924911022186,0.0,2872.3671875,0.0,1.040191650390625,6373.21875,6313.3828125,-0.046875,0.0,-0.06487300898879766,59.8359375,0.0,-68.62449264526367,622.3557711630754,43.28796460241019,0.16825318336486816,27.761664867401123,-2143.434177853199,27.929449558258057,-41.01467728614807,-2872.3671875,-1.040191650390625,59.8359375 +80,79,0.609375,0.6171875,0.0,0.0,0.003367701778188348,0.05134127289056778,5268.109375,4943.296875,7168.0,7168.0,97.6230525970459,121.6342408657074,863.4205523967669,650.2506978057589,73.42528029303712,58.93077433610135,19.74734330177307,18.040510892868042,99.53201103210449,70.38155055046082,7020.836741403724,9337.986942029609,119.51153469085693,88.65616273880005,237.3109118938446,230.15904808044434,0.0,1913.40625,0.0,0.8338128328323364,5268.109375,4943.296875,-0.0078125,0.0,-0.04797357111237943,324.8125,0.0,-24.0111882686615,213.16985459100795,14.494505956935768,1.7068324089050293,29.150460481643677,-2317.150200625885,30.855371952056885,7.1518638134002686,-1913.40625,-0.8338128328323364,324.8125 +81,80,0.3671875,0.390625,0.0,0.0,0.002277788706123829,0.05330277234315872,6613.9765625,6445.8828125,7168.0,7168.0,76.58322405815125,146.76756930351257,1381.8120913745536,702.7037750194021,93.59752201810134,48.83912729505461,23.656949520111084,24.74256706237793,122.25574779510498,96.86928844451904,7081.785647011236,8715.58998271727,146.14840006828308,121.84807205200195,243.0346007347107,288.63014340400696,0.0,2416.4609375,0.0,1.306652545928955,6613.9765625,6445.8828125,-0.0234375,0.0,-0.05102498363703489,168.09375,0.0,-70.18434524536133,679.1083163551515,44.758394723046734,-1.0856175422668457,25.386459350585938,-1633.804335706035,24.300328016281128,-45.595542669296265,-2416.4609375,-1.306652545928955,168.09375 +82,81,0.609375,0.703125,0.0,0.0,0.003348322119563818,0.05149313062429428,5411.5859375,5127.9609375,7168.0,7168.0,95.51388788223267,122.07451224327087,906.5213124478672,672.1089725633756,75.04667812117592,58.71823584038217,19.58725380897522,19.18023681640625,100.23277258872986,74.06468439102173,7157.36960548435,9196.00219187006,120.05269837379456,93.48610377311707,236.16118240356445,235.63641119003296,0.0,1910.4140625,0.0,0.8550747632980347,5411.5859375,5127.9609375,-0.09375,0.0,-0.04814480850473046,283.625,0.0,-26.560624361038208,234.41233988449164,16.328442280793745,0.4070169925689697,26.16808819770813,-2038.63258638571,26.56659460067749,0.5247712135314941,-1910.4140625,-0.8550747632980347,283.625 +83,82,0.5078125,0.515625,0.0,0.0,0.003156757215037942,0.060070641338825226,5964.59375,5686.3828125,7168.0,7168.0,74.76973557472229,132.68651819229126,1276.365353795147,685.6923087554936,95.86766550533736,54.02206718253039,22.032607793807983,22.790008068084717,111.94546246528625,88.75120949745178,6995.227700656742,8422.11620813417,134.2117462158203,111.7749650478363,229.22729301452637,264.6090099811554,0.0,2423.3515625,0.0,1.0481499433517456,5964.59375,5686.3828125,-0.0078125,0.0,-0.056913884123787284,278.2109375,0.0,-57.91678261756897,590.6730450396534,41.845598322806964,-0.7574002742767334,23.194252967834473,-1426.888507477428,22.43678116798401,-35.38171696662903,-2423.3515625,-1.0481499433517456,278.2109375 +84,83,0.484375,0.53125,0.0,0.0,0.0030816581565886736,0.050385478883981705,6016.1953125,5929.15625,7168.0,7168.0,93.35068941116333,136.0602638721466,1031.1560161706623,697.2388359407001,76.78572108266418,52.68253784025909,22.77182674407959,23.310746431350708,114.59602069854736,90.42062854766846,6878.502370283369,8594.366268868833,137.60077214241028,113.96474695205688,251.30492043495178,270.1660044193268,0.0,2190.203125,0.0,1.1127607822418213,6016.1953125,5929.15625,-0.046875,0.0,-0.04730382072739303,87.0390625,0.0,-42.709574460983276,333.9171802299621,24.103183242405095,-0.5389196872711182,24.175392150878906,-1715.863898585464,23.636025190353394,-18.861083984375,-2190.203125,-1.1127607822418213,87.0390625 +85,84,0.59375,0.6328125,0.0,0.0,0.0030895713716745377,0.05773267149925232,5696.65625,5357.0546875,7168.0,7168.0,83.29400277137756,126.48072338104248,1094.2744611538924,677.6754015058633,86.05661586074177,56.672667647585364,22.113557815551758,20.81862187385559,107.06492257118225,81.19642210006714,7006.178886472219,8702.932736729734,129.41826176643372,102.25190925598145,233.05269241333008,248.68501925468445,0.0,2075.5546875,0.0,0.996340811252594,5696.65625,5357.0546875,-0.0390625,0.0,-0.05464310012757778,339.6015625,0.0,-43.18672060966492,416.5990596480291,29.38394821315641,1.294935941696167,25.868500471115112,-1696.7538502575144,27.16635251045227,-15.63232684135437,-2075.5546875,-0.996340811252594,339.6015625 +86,85,0.1484375,0.2421875,0.0,0.0,0.0025229104794561863,0.060861509293317795,6778.65625,6582.796875,7168.0,7168.0,81.5180344581604,151.1938579082489,1330.484729187956,696.6205602341049,87.93146262228657,47.40933328356406,22.863046646118164,24.31824231147766,121.82742524147034,94.05258250236511,7277.34332596078,9159.876072284733,144.92725229263306,118.6089243888855,247.75180315971375,289.7472381591797,0.0,2662.3828125,0.0,1.1775352954864502,6778.65625,6582.796875,-0.09375,0.0,-0.05833859881386161,195.859375,0.0,-69.6758234500885,633.864168953851,40.522129338722515,-1.455195665359497,27.774842739105225,-1882.5327463239528,26.31832790374756,-41.99543499946594,-2662.3828125,-1.1775352954864502,195.859375 +87,86,0.5078125,0.546875,0.0,0.0,0.002939848229289055,0.05863761901855469,6129.015625,5911.8125,7168.0,7168.0,78.61101460456848,135.78864288330078,1247.4619554687822,696.5899208617285,91.18315080980308,52.78791987162217,21.990288496017456,22.600712776184082,113.42349028587341,87.54368281364441,7077.070172826227,8851.626697605927,135.65013146400452,110.3806459903717,234.62561798095703,266.20281767845154,0.0,2306.265625,0.0,1.0338592529296875,6129.015625,5911.8125,-0.0390625,0.0,-0.05569777078926563,217.203125,0.0,-57.1776282787323,550.8720346070537,38.39523093818091,-0.610424280166626,25.879807472229004,-1774.5565247797003,25.269485473632812,-31.577199697494507,-2306.265625,-1.0338592529296875,217.203125 +88,87,0.359375,0.3828125,0.0,0.0,0.0032676784321665764,0.06401734054088593,5915.8125,5566.890625,7168.0,7168.0,75.25092315673828,131.8213667869568,1257.8317451714129,675.6890189429681,95.25464538248853,54.37661719578867,21.28290343284607,20.552404403686523,109.504403591156,79.70811080932617,7075.7337110649105,9160.447946717213,131.0193738937378,100.49708199501038,226.83786010742188,252.243426322937,0.0,2227.953125,0.0,0.9506258964538574,5915.8125,5566.890625,-0.0234375,0.0,-0.06074966210871935,348.921875,0.0,-56.570443630218506,582.1427262284448,40.878028186699865,0.7304990291595459,29.796292781829834,-2084.7142356523027,30.522291898727417,-25.405566215515137,-2227.953125,-0.9506258964538574,348.921875 +89,88,0.7109375,0.765625,0.0,0.0,0.0033626402728259563,0.049085937440395355,5241.84375,4825.3046875,7168.0,7168.0,92.22923469543457,117.68317198753357,909.3591666130524,656.0400582011715,77.71939151041033,60.90930316493612,19.831007719039917,18.561187982559204,99.26457595825195,69.56559562683105,6953.175322990162,9155.200846930667,119.32651615142822,88.35657477378845,231.6993329524994,225.8910493850708,0.0,1726.984375,0.0,0.8024530410766602,5241.84375,4825.3046875,-0.0546875,0.0,-0.0457232971675694,416.5390625,0.0,-25.453937292099,253.31910841188085,16.810088345474213,1.269819736480713,29.6989803314209,-2202.025523940505,30.96994137763977,5.808283567428589,-1726.984375,-0.8024530410766602,416.5390625 +90,89,0.4453125,0.53125,0.0,0.0,0.003033608430996537,0.060414716601371765,6126.2578125,5825.0625,7168.0,7168.0,73.60523414611816,133.76704716682434,1331.7004712655946,696.7411030892132,97.38437874907609,53.58569357564275,21.624712228775024,22.535499811172485,112.3143298625946,84.82846808433533,7161.054168100967,9026.875261247385,134.17373919487,107.60248851776123,227.83800601959229,261.4341776371002,0.0,2288.59375,0.0,0.952215313911438,6126.2578125,5825.0625,-0.0859375,0.0,-0.05738110817037523,301.1953125,0.0,-60.16181302070618,634.9593681763814,43.79868517343334,-0.9107875823974609,27.485861778259277,-1865.8210931464182,26.571250677108765,-33.596171617507935,-2288.59375,-0.952215313911438,301.1953125 +91,90,0.4140625,0.4609375,0.0,0.0,0.00305469473823905,0.06480663269758224,6111.0,5812.203125,7168.0,7168.0,73.50594878196716,134.62363743782043,1330.1780552485961,690.7795077439679,97.51591699416971,53.24473574197351,21.525495529174805,22.16942071914673,112.16860508918762,84.77176928520203,7159.65041520707,9022.366838030744,133.92790579795837,107.176913022995,227.68297839164734,261.92730712890625,0.0,2461.1640625,0.0,0.9817542433738708,6111.0,5812.203125,-0.046875,0.0,-0.061751937959343195,298.796875,0.0,-61.11768865585327,639.3985475046283,44.27118125219621,-0.6439251899719238,27.396835803985596,-1862.7164228236743,26.75099277496338,-34.24432873725891,-2461.1640625,-0.9817542433738708,298.796875 +92,91,0.4296875,0.4609375,0.0,0.0,0.003873650450259447,0.07219116389751434,5417.203125,5202.2890625,7168.0,7168.0,69.04252457618713,127.6612000465393,1255.38934927496,652.0119266437713,103.82007384579697,56.148618353790205,19.328886032104492,19.161120891571045,100.24809050559998,73.37577700614929,7119.078242793396,9351.383085762698,119.80834913253784,92.77202224731445,209.42620730400085,240.6375994682312,0.0,2286.1796875,0.0,0.957042396068573,5417.203125,5202.2890625,-0.03125,0.0,-0.0683175134472549,214.9140625,0.0,-58.61867547035217,603.3774226311888,47.671455492006764,0.16776514053344727,26.872313499450684,-2232.3048429693017,27.03632688522339,-31.211392164230347,-2286.1796875,-0.957042396068573,214.9140625 +93,92,0.2578125,0.28125,0.0,0.0,0.002991404850035906,0.06960391998291016,6298.7734375,6160.96875,7168.0,7168.0,78.55419754981995,144.06466341018677,1282.9406720892791,684.2448221971812,91.24910219411223,49.75543502705434,21.63537096977234,22.750226974487305,113.97647595405579,87.58995938301086,7257.102775606274,9241.926879543826,135.84888291358948,110.57700705528259,234.47677564620972,274.57552552223206,0.0,2508.9296875,0.0,1.0879515409469604,6298.7734375,6160.96875,-0.0234375,0.0,-0.06661251513287425,137.8046875,0.0,-65.51046586036682,598.695849892098,41.493667167057886,-1.1148560047149658,26.386516571044922,-1984.8241039375516,25.271875858306885,-40.09874987602234,-2508.9296875,-1.0879515409469604,137.8046875 +94,93,0.671875,0.75,0.0,0.0,0.004158555530011654,0.05783912539482117,5069.5078125,4692.875,7168.0,7168.0,91.08276152610779,117.02773666381836,890.5321231037793,641.6085804999974,78.69765782129232,61.25043689934184,18.951768159866333,18.592655658721924,96.18557167053223,71.31425952911377,6969.038997824681,8723.528844130047,115.36727333068848,90.14009499549866,226.70958542823792,227.1968698501587,0.0,1727.796875,0.0,0.882881224155426,5069.5078125,4692.875,-0.078125,0.0,-0.05368056986480951,376.6328125,0.0,-25.94497513771057,248.92354260378193,17.447220921950482,0.3591125011444092,24.871312141418457,-1754.489846305366,25.22717833518982,-0.48728442192077637,-1727.796875,-0.882881224155426,376.6328125 +95,94,0.4375,0.546875,0.0,0.0,0.0032526510767638683,0.05503734201192856,5894.7109375,5598.828125,7168.0,7168.0,72.86558127403259,130.79356622695923,1294.3748385853,684.9056309433015,98.37292003535406,54.80391892948117,20.966469526290894,20.751298904418945,108.69650459289551,80.00089311599731,7102.739898505084,9177.022548179431,129.89422869682312,100.98760390281677,222.89393949508667,252.06781458854675,0.0,1935.203125,0.0,0.9974695444107056,5894.7109375,5598.828125,-0.109375,0.0,-0.05178469093516469,295.8828125,0.0,-57.927984952926636,609.4692076419985,43.569001105872886,0.21517062187194824,28.695611476898193,-2074.2826496743473,28.906624794006348,-29.173875093460083,-1935.203125,-0.9974695444107056,295.8828125 +96,95,0.5,0.5234375,0.0,0.0,0.003820810467004776,0.07296374440193176,6048.96875,5751.8203125,7168.0,7168.0,75.58201789855957,134.48657822608948,1280.5096065296304,684.2996990025803,94.83737268857182,53.298999012002945,21.212933778762817,22.401620388031006,110.51605129241943,85.94500207901001,7164.4615487118,8770.17839044402,131.96567845344543,108.58190751075745,227.65438508987427,262.9817771911621,0.0,2261.515625,0.0,1.0711255073547363,6048.96875,5751.8203125,-0.0234375,0.0,-0.06914293393492699,297.1484375,0.0,-58.90456032752991,596.2099075270502,41.538373676568874,-1.1886866092681885,24.571049213409424,-1605.71684173222,23.38377094268799,-35.32739210128784,-2261.515625,-1.0711255073547363,297.1484375 +97,96,0.71875,0.7421875,0.0,0.0,0.00416222820058465,0.06330355256795883,5211.4375,4793.15625,7168.0,7168.0,92.88889789581299,117.76237034797668,897.6637885565711,651.2309473169299,77.16745663232916,60.86834002083379,20.802613019943237,19.998307704925537,101.94027328491211,77.25141763687134,6735.159499534316,8194.5939552293,122.9748260974884,97.48343276977539,235.90153169631958,235.33713173866272,0.0,1880.875,0.0,0.9136902689933777,5211.4375,4793.15625,-0.0234375,0.0,-0.05914132436737418,418.28125,0.0,-24.873472452163696,246.43284123964122,16.29911661149537,0.8043053150177002,24.68885564804077,-1459.4344556949845,25.491393327713013,0.5643999576568604,-1880.875,-0.9136902689933777,418.28125 +98,97,0.828125,0.7578125,0.0,0.0,0.004585837945342064,0.07013313472270966,4836.0,4569.7734375,7168.0,7168.0,93.25037169456482,116.17204928398132,829.7661295489498,629.3800914303216,76.8683263105727,61.701588671108844,18.78017807006836,17.801530361175537,92.58408379554749,67.24532008171082,7017.7072922683765,9155.298826028533,111.59335470199585,85.2790138721466,224.99350023269653,221.22352170944214,0.0,1828.8203125,0.0,0.920142650604248,4836.0,4569.7734375,0.0703125,0.0,-0.06554729677736759,266.2265625,0.0,-22.921677589416504,200.38603811862822,15.166737639463854,0.9786477088928223,25.33876371383667,-2137.5915337601564,26.314340829849243,3.7699785232543945,-1828.8203125,-0.920142650604248,266.2265625 +99,98,0.609375,0.640625,0.0,0.0,0.004177144728600979,0.06724589318037033,5084.5625,4786.21875,7168.0,7168.0,94.65811443328857,117.76825475692749,859.4403183187691,650.2558788704081,75.72515090665294,60.86529867318388,18.4582941532135,17.675926685333252,94.15119075775146,67.62129020690918,7108.141645514998,9332.14965388996,112.84063935279846,85.53007388114929,227.5397334098816,223.24626636505127,0.0,1830.0546875,0.0,0.8510475158691406,5084.5625,4786.21875,-0.03125,0.0,-0.06306874845176935,298.34375,0.0,-23.110140323638916,209.184439448361,14.859852233469063,0.782367467880249,26.529900550842285,-2224.008008374961,27.31056547164917,4.293467044830322,-1830.0546875,-0.8510475158691406,298.34375 +100,99,0.34375,0.390625,0.0,0.0,0.0034313166979700327,0.06733152270317078,5846.609375,5711.9765625,7168.0,7168.0,74.32882928848267,138.74481415748596,1258.539262564371,658.7029976937627,96.43633659531739,51.66319219588108,20.202749729156494,21.003921270370483,106.38867950439453,80.47352170944214,7217.741620416309,9327.9501636614,126.82407307624817,101.71230864524841,221.32832622528076,260.5760066509247,0.0,2075.0234375,0.0,1.1081247329711914,5846.609375,5711.9765625,-0.046875,0.0,-0.06390020600520074,134.6328125,0.0,-64.4159848690033,599.8362648706084,44.773144399436305,-0.8011715412139893,25.915157794952393,-2110.2085432450913,25.111764430999756,-39.24768042564392,-2075.0234375,-1.1081247329711914,134.6328125 +101,100,0.265625,0.28125,0.0,0.0,0.0036121434532105923,0.07873289287090302,5661.234375,5561.765625,7168.0,7168.0,93.74137806892395,131.38273692131042,966.2728654724988,677.3207202503178,76.4656990078563,54.55815709101233,21.582914352416992,20.691041469573975,105.88982105255127,79.79546976089478,7002.580537292678,9132.974618530876,127.70471167564392,100.72114539146423,241.61266493797302,252.0437343120575,0.0,2094.375,0.0,1.0663609504699707,5661.234375,5561.765625,-0.015625,0.0,-0.07512074941769242,99.46875,0.0,-37.641358852386475,288.952145222181,21.907541916843968,0.8918728828430176,26.094351291656494,-2130.394081238198,26.983566284179688,-10.431069374084473,-2094.375,-1.0663609504699707,99.46875 +102,101,0.3671875,0.4296875,0.0,0.0,0.0029467607382684946,0.053985580801963806,6404.8203125,6129.515625,7168.0,7168.0,75.4841537475586,140.86076426506042,1357.5978521626396,696.2353960784676,94.96032801761173,50.8871298363243,21.84903573989868,22.633140087127686,115.40716576576233,87.67283082008362,7258.830025340813,9153.143482349742,137.49005031585693,110.54175591468811,233.24423122406006,271.48811960220337,0.0,1893.0234375,0.0,1.0898200273513794,6404.8203125,6129.515625,-0.0625,0.0,-0.05103882006369531,275.3046875,0.0,-65.37661051750183,661.362456084172,44.07319818128743,-0.7841043472290039,27.73433494567871,-1894.3134570089287,26.948294401168823,-38.24388837814331,-1893.0234375,-1.0898200273513794,275.3046875 +103,102,0.6953125,0.796875,0.0,0.0,0.004246615804731846,0.06361675262451172,5283.1796875,5034.1796875,7168.0,7168.0,91.60191893577576,120.409264087677,922.8068143339505,668.9425071259394,78.2516358093508,59.530303206409116,21.263068437576294,19.825845956802368,102.01584148406982,76.26808214187622,6821.91108631569,8707.063051155197,123.51047110557556,96.32734179496765,235.0486936569214,236.77469944953918,0.0,1801.0234375,0.0,1.0069422721862793,5283.1796875,5034.1796875,-0.1015625,0.0,-0.05937013681977987,249.0,0.0,-28.807345151901245,253.86430720801104,18.721332602941686,1.4372224807739258,25.747759342193604,-1885.151964839507,27.18312931060791,-1.7260057926177979,-1801.0234375,-1.0069422721862793,249.0 +104,103,0.5234375,0.515625,0.0,0.0,0.0040838452987372875,0.07183686643838882,5373.1875,5019.4453125,7168.0,7168.0,95.88099980354309,122.4566261768341,896.6427152006305,655.8332326094492,74.75933724811995,58.535011324328124,20.039458751678467,18.31595230102539,101.06401443481445,70.92465662956238,6952.504350133477,9268.553860379207,121.33668184280396,89.47461438179016,237.2939395904541,231.9485936164856,0.0,1908.4375,0.0,1.0243078470230103,5373.1875,5019.4453125,0.0078125,0.0,-0.06775302113965154,353.7421875,0.0,-26.575626373291016,240.80948259118134,16.224325923791824,1.7235064506530762,30.139357805252075,-2316.04951024573,31.862067461013794,5.345345973968506,-1908.4375,-1.0243078470230103,353.7421875 +105,104,0.609375,0.6484375,0.0,0.0,0.004841086454689503,0.06662565469741821,4717.921875,4389.2109375,7168.0,7168.0,106.56363606452942,115.68868017196655,708.3725066803195,607.0375675097152,67.26497203660948,61.9593895387609,17.554866790771484,15.995535373687744,88.9650046825409,61.57203960418701,6993.558896783842,9421.597915696651,106.7512354850769,77.80124688148499,233.33837270736694,213.50018692016602,0.0,1569.40625,0.0,0.8760596513748169,4717.921875,4389.2109375,-0.0390625,0.0,-0.06178456824272871,328.7109375,0.0,-9.125044107437134,101.33493917060423,5.305582497848576,1.5593314170837402,27.392965078353882,-2428.0390189128093,28.94998860359192,19.838185787200928,-1569.40625,-0.8760596513748169,328.7109375 +106,105,0.5,0.515625,0.0,0.0,0.004048597998917103,0.07435937225818634,5900.8203125,5615.78125,7168.0,7168.0,77.84810161590576,131.91194248199463,1212.7864782859356,681.1551578225334,92.07674755341047,54.33928016774068,21.83966064453125,20.990456581115723,111.36162495613098,81.09777927398682,6959.605701742508,9106.883155269074,133.4352822303772,102.32351517677307,231.5693278312683,254.02950048446655,0.0,2128.84375,0.0,1.0669434070587158,5900.8203125,5615.78125,-0.015625,0.0,-0.07031077425926924,285.0390625,0.0,-54.06384086608887,531.6313204634022,37.73746738566979,0.8492040634155273,30.263845682144165,-2147.2774535265653,31.111767053604126,-22.460172653198242,-2128.84375,-1.0669434070587158,285.0390625 +107,106,0.546875,0.609375,0.0,0.0,0.0031697081867605448,0.06376191973686218,6045.328125,5726.4375,7168.0,7168.0,86.59958600997925,131.90617418289185,1116.9250854023012,694.6073644207281,82.77175827577283,54.341656441808055,22.32366633415222,24.072847604751587,113.65227508544922,93.04445219039917,7025.4466916714255,8142.774578860677,136.21054363250732,117.35303258895874,243.04666209220886,269.21928191185,0.0,2096.0234375,0.0,1.0791597366333008,6045.328125,5726.4375,-0.0625,0.0,-0.06059221155010164,318.890625,0.0,-45.3065881729126,422.3177209815731,28.43010183396477,-1.7491812705993652,20.60782289505005,-1117.3278871892517,18.857511043548584,-26.172619819641113,-2096.0234375,-1.0791597366333008,318.890625 +108,107,0.53125,0.5703125,0.0,0.0,0.0039755310863256454,0.07308894395828247,5931.9453125,5733.1171875,7168.0,7168.0,78.10310959815979,132.99516367912292,1215.2028963804053,689.7233888994372,91.77611540538314,53.89669670465773,21.118566751480103,22.110548496246338,109.096360206604,84.71096968650818,7152.36510661119,8910.85301931355,130.44887399673462,107.05737352371216,228.7165334224701,259.98730993270874,0.0,2119.203125,0.0,1.1013646125793457,5931.9453125,5733.1171875,-0.0390625,0.0,-0.06911341287195683,198.828125,0.0,-54.892054080963135,525.4795074809681,37.87941870072541,-0.9919817447662354,24.385390520095825,-1758.4879127023605,23.39150047302246,-31.270776510238647,-2119.203125,-1.1013646125793457,198.828125 +109,108,0.265625,0.3671875,0.0,0.0,0.0036583440378308296,0.07278694212436676,6648.640625,6302.21875,7168.0,7168.0,78.4008960723877,142.47200775146484,1356.8499255643808,707.7565733186157,91.42752645813859,50.311637444628495,22.711830139160156,23.884265899658203,119.99580383300781,91.24057841300964,7248.4034625946315,9046.808057962775,142.943528175354,115.36114382743835,241.37930250167847,277.93513226509094,0.0,2242.75,0.0,1.1905896663665771,6648.640625,6302.21875,-0.1015625,0.0,-0.06912859808653593,346.421875,0.0,-64.07111167907715,649.0933522457651,41.11588901351009,-1.1724357604980469,28.75522541999817,-1798.4045953681434,27.58238434791565,-36.555829763412476,-2242.75,-1.1905896663665771,346.421875 +110,109,0.5078125,0.6015625,0.0,0.0,0.003844281192868948,0.08454635739326477,5916.3828125,5395.2109375,7168.0,7168.0,73.38776111602783,126.85289788246155,1289.8898066986521,680.4998264997059,97.67296196251604,56.50639535757137,21.3175847530365,20.496980667114258,109.59511208534241,78.31551432609558,7084.12068044991,9061.742186165131,131.14599871635437,99.04633665084839,224.72203755378723,246.22314834594727,0.0,2196.390625,0.0,1.0549137592315674,5916.3828125,5395.2109375,-0.09375,0.0,-0.08070207620039582,521.171875,0.0,-53.465136766433716,609.3899801989462,41.16656660494467,0.8206040859222412,31.279597759246826,-1977.6215057152212,32.09966206550598,-21.501110792160034,-2196.390625,-1.0549137592315674,521.171875 +111,110,0.3984375,0.40625,0.0,0.0,0.0037394859828054905,0.07690070569515228,5876.3515625,5598.15625,7168.0,7168.0,72.65618538856506,134.25654125213623,1294.0622260468622,667.1592993877666,98.65643181878538,53.39032223791886,20.810767650604248,20.994709730148315,108.14483094215393,81.60544204711914,7136.919936680494,9021.604215744703,129.1892387866974,102.83570790290833,221.88037514686584,260.8293755054474,0.0,2145.6796875,0.0,1.0942015647888184,5876.3515625,5598.15625,-0.0078125,0.0,-0.07316121971234679,278.1953125,0.0,-61.60035586357117,626.9029266590956,45.26610958086652,-0.18394207954406738,26.53938889503479,-1884.6842790642086,26.353530883789062,-38.94900035858154,-2145.6796875,-1.0942015647888184,278.1953125 +112,111,0.4921875,0.578125,0.0,0.0,0.0041178688406944275,0.0626387894153595,5888.8515625,5584.828125,7168.0,7168.0,75.1332848072052,130.23498678207397,1254.059705252821,686.1232316130543,95.40378832621725,55.038973605413915,21.348065853118896,21.84487509727478,109.73607134819031,84.08380460739136,7052.676394294511,8741.493126196961,131.3162682056427,106.16187047958374,226.4701223373413,256.62686371803284,0.0,1886.4765625,0.0,1.096635341644287,5888.8515625,5584.828125,-0.0859375,0.0,-0.05852092057466507,304.0234375,0.0,-55.101701974868774,567.9364736397667,40.36481472080334,-0.4968092441558838,25.65226674079895,-1688.8167319024506,25.15439772605896,-30.15674138069153,-1886.4765625,-1.096635341644287,304.0234375 +113,112,0.6328125,0.7109375,0.0,0.0,0.004884717985987663,0.06804326176643372,4702.15625,4366.3046875,7168.0,7168.0,107.35833144187927,112.1139223575592,700.7793339330149,623.1239932646035,66.76705853872691,63.9349676585167,17.800419092178345,16.23530411720276,89.42784905433655,62.94965577125549,6937.659873973168,9172.901629490252,107.45944118499756,79.4152729511261,234.7709183692932,211.60994410514832,0.0,1562.3828125,0.0,0.8425335884094238,4702.15625,4366.3046875,-0.078125,0.0,-0.06315854378044605,335.8515625,0.0,-4.755590915679932,77.65534066841133,2.8320908802102167,1.565114974975586,26.478193283081055,-2235.241755517084,28.04416823387146,23.160974264144897,-1562.3828125,-0.8425335884094238,335.8515625 +114,113,0.65625,0.6953125,0.0,0.0,0.004629327915608883,0.07933022081851959,5604.5234375,5288.921875,7168.0,7168.0,89.89357805252075,124.5242087841034,997.5392785857126,679.5686624013533,79.73873279147996,57.56310415453174,20.977179050445557,20.403534412384033,105.47219848632812,78.40679574012756,6976.350266325204,8869.307735835662,126.68039155006409,99.04354071617126,236.9314911365509,243.7873978614807,0.0,1854.8046875,0.0,1.1360549926757812,5604.5234375,5288.921875,-0.0390625,0.0,-0.07470089290291071,315.6015625,0.0,-34.63063073158264,317.9706161843593,22.17562863694822,0.5736446380615234,27.06540274620056,-1892.9574695104584,27.636850833892822,-6.85590672492981,-1854.8046875,-1.1360549926757812,315.6015625 +115,114,0.8671875,0.921875,0.0,0.0,0.00558727839961648,0.0692211464047432,4352.390625,3797.921875,7168.0,7168.0,84.96411681175232,103.95403027534485,819.619536024737,584.5540556633163,84.36502689578379,68.95355553809692,18.346563577651978,13.890090942382812,86.33058023452759,55.44534134864807,6702.630729783677,9156.224628642827,104.9069550037384,69.56525826454163,209.9309949874878,193.42623114585876,0.0,1389.796875,0.0,0.7635071277618408,4352.390625,3797.921875,-0.0546875,0.0,-0.06363386800512671,554.46875,0.0,-18.98991346359253,235.06548036142078,15.411471357686864,4.456472635269165,30.885238885879517,-2453.5938988591497,35.34169673919678,16.50476384162903,-1389.796875,-0.7635071277618408,554.46875 +116,115,0.296875,0.296875,0.0,0.0,0.003766424022614956,0.07239103317260742,6338.8125,6311.65625,7168.0,7168.0,73.88105297088623,142.58776187896729,1372.7606188824384,708.2410065859673,97.02081537501424,50.2707939695723,22.640639066696167,24.42678141593933,116.75765657424927,94.61723160743713,7124.295094694936,8754.642108286866,139.63363122940063,119.28190612792969,233.63447189331055,282.1760284900665,0.0,2161.8828125,0.0,1.2690949440002441,6338.8125,6311.65625,0.0,0.0,-0.06862460914999247,27.15625,0.0,-68.70670890808105,664.5196122964711,46.750021405441935,-1.786142349243164,22.140424966812134,-1630.3470135919297,20.351725101470947,-48.54155659675598,-2161.8828125,-1.2690949440002441,27.15625 +117,116,0.4765625,0.5390625,0.0,0.0,0.004141475073993206,0.07835642248392105,5598.2421875,5325.9921875,7168.0,7168.0,90.53212594985962,125.76549816131592,989.3932574786608,677.5775252024681,79.1763136543367,56.994963680784736,20.555387020111084,20.913206815719604,104.41884303092957,80.19333505630493,7041.784592290499,8734.478987664072,125.20597791671753,101.34614181518555,235.91728138923645,246.9815230369568,0.0,1937.046875,0.0,1.111623764038086,5598.2421875,5325.9921875,-0.0625,0.0,-0.07421494740992785,272.25,0.0,-35.2333722114563,311.81573227619265,22.18134997355196,-0.3578197956085205,24.225507974624634,-1692.6943953735727,23.859836101531982,-11.064241647720337,-1937.046875,-1.111623764038086,272.25 +118,117,0.6328125,0.6484375,0.0,0.0,0.004722542595118284,0.07317814975976944,5339.1328125,5184.84375,7168.0,7168.0,93.50070095062256,126.18199563026428,913.6415463357145,657.4432397081455,76.66252688079204,56.806836539529115,19.801528692245483,20.008702754974365,100.30531072616577,76.62986040115356,6996.728238208081,8900.707849778784,120.33820343017578,96.87278985977173,233.89962029457092,242.87260007858276,0.0,1711.234375,0.0,1.064851999282837,5339.1328125,5184.84375,-0.015625,0.0,-0.06845560716465116,154.2890625,0.0,-32.681294679641724,256.19830662756897,19.855690341262928,-0.20717406272888184,23.675450325012207,-1903.9796115707031,23.465413570404053,-8.97297978401184,-1711.234375,-1.064851999282837,154.2890625 +119,118,0.2890625,0.3125,0.0,0.0,0.0036868329625576735,0.08272591978311539,6356.3515625,6191.8046875,7168.0,7168.0,76.31097459793091,143.1322956085205,1332.7260664124387,692.1489980916825,93.93144351473599,50.079543331053074,21.730194330215454,22.689363718032837,115.16589093208313,88.1476743221283,7236.977834795833,9216.249960619325,137.12897849082947,111.0741548538208,233.5460524559021,274.2335169315338,0.0,2177.4375,0.0,1.1932907104492188,6356.3515625,6191.8046875,-0.0234375,0.0,-0.07903908682055771,164.546875,0.0,-66.8213210105896,640.5770683207562,43.851900183682915,-0.9591693878173828,27.018216609954834,-1979.2721258234924,26.054823637008667,-40.687464475631714,-2177.4375,-1.1932907104492188,164.546875 +120,119,0.8125,0.796875,0.0,0.0,0.005694529507309198,0.07402831315994263,5084.09375,4613.453125,7168.0,7168.0,84.07317590713501,116.60512018203735,967.555931155165,633.0360955399195,85.25906060593668,61.47242924504277,19.23640489578247,18.027021169662476,96.6011712551117,68.6473023891449,6990.515655515581,8959.565468624402,116.06756615638733,86.90587997436523,220.1521601676941,223.5569818019867,0.0,1559.484375,0.0,0.9745934009552002,5084.09375,4613.453125,0.015625,0.0,-0.06833378365263343,470.640625,0.0,-32.531944274902344,334.5198356152455,23.786631360893914,1.2093837261199951,27.953868865966797,-1969.0498131088207,29.161686182022095,-3.4048216342926025,-1559.484375,-0.9745934009552002,470.640625 +121,120,0.375,0.4453125,0.0,0.0,0.004252546466886997,0.07866380363702774,5980.0078125,5726.5859375,7168.0,7168.0,72.10374045372009,133.22993326187134,1326.9786615496382,687.723642553396,99.41231834707372,53.80172326522802,21.19683003425598,21.719905376434326,110.06376194953918,83.24465870857239,7128.676924197074,9035.666812368621,131.49196600914001,105.20540475845337,223.65085625648499,258.4434959888458,0.0,1951.4609375,0.0,1.1309354305267334,5980.0078125,5726.5859375,-0.0703125,0.0,-0.07441125717014074,253.421875,0.0,-61.126192808151245,639.2550189962423,45.6105950818457,-0.5230753421783447,26.819103240966797,-1906.9898881715471,26.286561250686646,-34.79263973236084,-1951.4609375,-1.1309354305267334,253.421875 +122,121,0.4296875,0.4609375,0.0,0.0,0.003954859916120768,0.07408130913972855,6414.0,6243.265625,7168.0,7168.0,74.95979571342468,141.9290895462036,1369.0538911330154,703.8180144703955,95.62459358085297,50.504093437917305,23.061588287353516,24.55138659477234,117.79100155830383,93.89139175415039,7138.2023149180495,8722.43966885067,141.09101057052612,118.67883968353271,236.64294743537903,280.68576741218567,0.0,2007.671875,0.0,1.3970439434051514,6414.0,6243.265625,-0.03125,0.0,-0.07012644922360778,170.734375,0.0,-66.96929383277893,665.23587666262,45.12050014293567,-1.4897983074188232,23.899609804153442,-1584.23735393262,22.412170886993408,-44.04281997680664,-2007.671875,-1.3970439434051514,170.734375 +123,122,0.4765625,0.453125,0.0,0.0,0.004625740461051464,0.07687224447727203,5418.0,5186.875,7168.0,7168.0,91.0422842502594,126.47402143478394,952.1729459435549,656.1821871283947,78.73264669300713,56.67567077161505,20.09965229034424,18.98230767250061,100.95114135742188,74.07478475570679,7076.849160828992,9245.143300227273,121.28383708000183,93.2912950515747,232.39758729934692,240.15364003181458,0.0,1739.734375,0.0,1.116510033607483,5418.0,5186.875,0.0234375,0.0,-0.07224650401622057,231.125,0.0,-35.431737184524536,295.9907588151601,22.056975921392073,1.117344617843628,26.876356601715088,-2168.294139398281,27.992542028427124,-7.756052732467651,-1739.734375,-1.116510033607483,231.125 +124,123,0.5234375,0.6015625,0.0,0.0,0.005339169409126043,0.07803019881248474,5146.09375,4882.9765625,7168.0,7168.0,99.91386485099792,121.2181658744812,824.0848266933758,644.5207649891311,71.74179490193555,59.13305112554095,18.485332012176514,17.952850103378296,95.00328063964844,69.09854054450989,7097.481218123282,9270.890455166034,113.7185971736908,87.28472757339478,233.75064659118652,228.5335259437561,0.0,1640.5859375,0.0,0.9857065677642822,5146.09375,4882.9765625,-0.078125,0.0,-0.0726910294033587,263.1171875,0.0,-21.304301023483276,179.56406170424475,12.608743776394597,0.5324819087982178,25.90474009513855,-2173.4092370427525,26.43386960029602,5.21712064743042,-1640.5859375,-0.9857065677642822,263.1171875 +125,124,0.3359375,0.4609375,0.0,0.0,0.004663742147386074,0.07813021540641785,6469.7421875,6221.828125,7168.0,7168.0,75.90293550491333,143.12142658233643,1363.7927744349129,695.5579774264634,94.43640028304311,50.08334650630607,22.44672131538391,24.18750762939453,117.79836964607239,92.77471923828125,7175.370954110502,8768.70344291296,140.4806101322174,117.19969010353088,236.54250240325928,280.61132884025574,0.0,2052.5078125,0.0,1.3073701858520508,6469.7421875,6221.828125,-0.125,0.0,-0.07346647325903177,247.9140625,0.0,-67.2184910774231,668.2347970084495,44.35305377673704,-1.7407863140106201,25.023650407791138,-1593.332488802458,23.280920028686523,-44.06882643699646,-2052.5078125,-1.3073701858520508,247.9140625 +126,125,0.59375,0.7421875,0.0,0.0,0.006263937801122665,0.09237298369407654,5540.0703125,4940.1328125,7168.0,7168.0,78.97179651260376,120.83701825141907,1122.4402750652005,654.1217761227881,90.76658144475667,59.31957030821407,20.20163607597351,19.563694953918457,103.53171038627625,75.01809334754944,7084.756904559257,8753.954822039637,123.96575212478638,94.81629705429077,223.03588557243347,235.66773390769958,0.0,1853.2890625,0.0,1.125150442123413,5540.0703125,4940.1328125,-0.1484375,0.0,-0.08610904589295387,599.9375,0.0,-41.86522173881531,468.31849894241236,31.4470111365426,0.6379411220550537,28.513617038726807,-1669.19791748038,29.149455070495605,-12.631848335266113,-1853.2890625,-1.125150442123413,599.9375 +127,126,0.359375,0.5078125,0.0,0.0,0.004869729280471802,0.07827115058898926,6170.765625,5799.1875,7168.0,7168.0,82.13327503204346,132.9409520626068,1202.0980529691606,697.9564879022619,87.27278922219378,53.91867508684851,22.261507272720337,22.13596796989441,114.6684877872467,85.70926403999329,7053.81239090483,8882.213708482797,137.16374897956848,108.07984495162964,239.5765664577484,261.13457322120667,0.0,1955.9765625,0.0,1.2440245151519775,6170.765625,5799.1875,-0.1484375,0.0,-0.07340142130851746,371.578125,0.0,-50.807677030563354,504.1415650668987,33.35411413534527,0.12553930282592773,28.959223747253418,-1828.4013175779673,29.083904027938843,-21.558006763458252,-1955.9765625,-1.2440245151519775,371.578125 +128,127,0.4296875,0.515625,0.0,0.0,0.004756408743560314,0.08267960697412491,5884.234375,5455.8515625,7168.0,7168.0,72.17803049087524,129.5961515903473,1304.3823634381679,673.5819229874561,99.3099971175603,55.310284387595146,20.433533191680908,20.452435970306396,107.25325870513916,78.67826771736145,7209.235498622203,9130.61536357999,127.91927981376648,99.36478185653687,220.225515127182,248.85659742355347,0.0,1986.5234375,0.0,1.120634913444519,5884.234375,5455.8515625,-0.0859375,0.0,-0.0779231982305646,428.3828125,0.0,-57.418121099472046,630.8004404507118,43.999712729965154,-0.01890277862548828,28.57499098777771,-1921.3798649577866,28.554497957229614,-28.63108229637146,-1986.5234375,-1.120634913444519,428.3828125 +129,128,0.765625,0.8125,0.0,0.0,0.005897243972867727,0.0777820497751236,4612.875,4217.5234375,7168.0,7168.0,95.55702638626099,109.89947056770325,772.3764833541502,614.019109022267,75.01279886028979,65.22324414278388,17.77504253387451,16.29570746421814,88.52249574661255,62.55459141731262,6885.30067819878,8934.579977854655,106.52857780456543,79.08116030693054,222.09120178222656,209.07893896102905,0.0,1313.25,0.0,1.0015547275543213,4612.875,4217.5234375,-0.046875,0.0,-0.07188480580225587,395.3515625,0.0,-14.34244418144226,158.3573743318832,9.789554717505908,1.479335069656372,25.967904329299927,-2049.2792996558755,27.447417497634888,13.01226282119751,-1313.25,-1.0015547275543213,395.3515625 +130,129,0.421875,0.4765625,0.0,0.0,0.004681185819208622,0.08066342025995255,6383.4765625,6101.03125,7168.0,7168.0,73.50231766700745,138.49942541122437,1389.556523410758,704.8151984036239,97.52073441375929,51.75472734790917,22.114240169525146,22.95432186126709,116.32373332977295,88.02919840812683,7169.070112595463,9062.674821838991,138.67279982566833,111.22865271568298,232.45295238494873,269.6622202396393,0.0,2025.03125,0.0,1.2909252643585205,6383.4765625,6101.03125,-0.0546875,0.0,-0.07598223444074392,282.4453125,0.0,-64.99710774421692,684.7413250071342,45.76600706585012,-0.8400816917419434,28.294534921646118,-1893.6047092435283,27.44414710998535,-37.20926785469055,-2025.03125,-1.2909252643585205,282.4453125 +131,130,0.5078125,0.59375,0.0,0.0,0.005255087278783321,0.08328816294670105,5640.6953125,5083.1796875,7168.0,7168.0,81.07238936424255,120.80159544944763,1113.2165427432903,673.259940793041,88.41481121020824,59.33696465954064,20.733874559402466,19.87248969078064,105.7122004032135,76.59913563728333,6997.650197218986,8725.620654062313,126.67755603790283,96.70333313941956,227.91082644462585,237.39089608192444,0.0,1841.9453125,0.0,1.1693823337554932,5640.6953125,5083.1796875,-0.0859375,0.0,-0.07803307566791773,557.515625,0.0,-39.72920608520508,439.95660195024936,29.0778465506676,0.8613848686218262,29.113064765930176,-1727.9704568433272,29.974222898483276,-9.480069637298584,-1841.9453125,-1.1693823337554932,557.515625 +132,131,0.625,0.6953125,0.0,0.0,0.005158852320164442,0.07806266844272614,5745.03125,5532.359375,7168.0,7168.0,82.68296098709106,127.0871090888977,1111.722402084163,696.5124207686686,86.6925895544441,56.402258666423585,20.89279270172119,22.002830266952515,106.69609045982361,84.38228273391724,7075.535727190206,8623.966743050658,127.82179284095764,106.62176871299744,230.79511952400208,257.2693314552307,0.0,1738.671875,0.0,1.278160810470581,5745.03125,5532.359375,-0.0703125,0.0,-0.0729038161225617,212.671875,0.0,-44.40414810180664,415.20998131549436,30.290330888020513,-1.1100375652313232,22.313807725906372,-1548.4310158604521,21.200024127960205,-26.474211931228638,-1738.671875,-1.278160810470581,212.671875 +133,132,0.671875,0.7421875,0.0,0.0,0.005110069178044796,0.09850157052278519,5215.2421875,4544.7890625,7168.0,7168.0,90.37609124183655,113.92680096626282,923.2959055145825,638.2749658838713,79.31301189845902,62.91759216624245,19.96185851097107,17.473338842391968,99.4376962184906,67.12414216995239,6889.61054060107,8927.771448947591,119.63164186477661,84.82887506484985,230.51304030418396,218.6557011604309,0.0,1655.703125,0.0,1.0348765850067139,5215.2421875,4544.7890625,-0.0703125,0.0,-0.09339150134474039,670.453125,0.0,-23.55070972442627,285.02093963071115,16.395419732216574,2.4885196685791016,32.31355404853821,-2038.160908346521,34.80276679992676,11.857339143753052,-1655.703125,-1.0348765850067139,670.453125 +134,133,0.5703125,0.6171875,0.0,0.0,0.006342677399516106,0.07979446649551392,4611.265625,4254.0,7168.0,7168.0,102.39644837379456,112.01653742790222,720.5352448423587,607.6245665405287,70.00242795368716,63.990551436332126,16.914050102233887,15.416209697723389,85.9440426826477,59.69638705253601,7061.362033444709,9400.099867118528,103.08779406547546,75.34503030776978,225.47550654411316,207.1798243522644,0.0,1345.296875,0.0,0.8975697755813599,4611.265625,4254.0,-0.046875,0.0,-0.07345178909599781,357.265625,0.0,-9.620089054107666,112.91067830182999,6.011876517355034,1.497840404510498,26.247655630111694,-2338.7378336738193,27.74276375770569,18.295682191848755,-1345.296875,-0.8975697755813599,357.265625 +135,134,0.453125,0.4296875,0.0,0.0,0.004833338316529989,0.10322726517915726,5729.7578125,5378.984375,7168.0,7168.0,71.32682037353516,132.80798435211182,1285.2966740967354,648.0314449455122,100.49515683527635,53.97265860910583,20.28565788269043,20.485159873962402,104.87684178352356,78.562992811203,7210.447865706064,9054.00844019995,125.39629936218262,99.28219485282898,216.8814356327057,252.44005918502808,0.0,2038.796875,0.0,1.2478289604187012,5729.7578125,5378.984375,0.0234375,0.0,-0.09839392686262727,350.7734375,0.0,-61.48116397857666,637.2652291512231,46.52249822617052,-0.19950199127197266,26.313848972320557,-1843.5605744938866,26.114104509353638,-35.55862355232239,-2038.796875,-1.2478289604187012,350.7734375 +136,135,0.75,0.7578125,0.0,0.0,0.005694401916116476,0.07593520730733871,5346.1484375,5031.0234375,7168.0,7168.0,87.37766671180725,121.01380157470703,978.9500935304935,665.1834249691439,82.03469227030064,59.23291316135445,20.58405303955078,19.877583742141724,102.42954277992249,76.42735505104065,6858.206928730875,8663.743492860598,123.24648380279541,96.5377950668335,230.67497873306274,237.5981011390686,0.0,1519.3828125,0.0,1.417217493057251,5346.1484375,5031.0234375,-0.0078125,0.0,-0.07024080539122224,315.125,0.0,-33.63613486289978,313.76666856134966,22.801779108946185,0.7064692974090576,26.002187728881836,-1805.536564129723,26.708688735961914,-6.923122406005859,-1519.3828125,-1.417217493057251,315.125 +137,136,0.328125,0.421875,0.0,0.0,0.004218222573399544,0.09620165079832077,6277.1640625,5976.28125,7168.0,7168.0,76.23289632797241,138.0221655368805,1317.4709323374764,692.790897955079,94.02764876152055,51.93368740533679,22.224551916122437,22.07953715324402,115.3943350315094,85.21946597099304,7129.266785716656,9201.700469079604,137.85260891914368,107.53670740127563,234.15970945358276,265.40975284576416,0.0,1952.640625,0.0,1.2984943389892578,6277.1640625,5976.28125,-0.09375,0.0,-0.09198342822492123,300.8828125,0.0,-61.78926920890808,624.6800343823974,42.09396135618376,0.14501476287841797,30.174869060516357,-2072.4336833629477,30.315901517868042,-31.250043392181396,-1952.640625,-1.2984943389892578,300.8828125 +138,137,0.8359375,0.9140625,0.0,0.0,0.005951998755335808,0.07409875094890594,5147.8984375,4431.765625,7168.0,7168.0,86.94544792175293,111.31993770599365,947.3339544368798,636.9770901891385,82.44249896154292,64.39098105616405,20.55040740966797,17.491096019744873,100.90635991096497,67.68025588989258,6729.119954372821,8678.247330440645,121.68774580955505,85.4098961353302,228.62155270576477,217.3167634010315,0.0,1295.1171875,0.0,1.0313926935195923,5147.8984375,4431.765625,-0.078125,0.0,-0.06814675219357014,716.1328125,0.0,-24.374489784240723,310.3568642477413,18.051517905378873,3.0593113899230957,33.22610402107239,-1949.1273760678241,36.27784967422485,11.304789304733276,-1295.1171875,-1.0313926935195923,716.1328125 +139,138,0.4609375,0.5390625,0.0,0.0,0.005353535525500774,0.08476299047470093,6137.015625,5754.8828125,7168.0,7168.0,72.97232794761658,134.20034265518188,1345.6093941594904,686.1243658415099,98.22901641764221,53.412680311984445,21.422632694244385,21.54769468307495,111.71978092193604,83.4072949886322,7198.027004384358,9054.951369696559,133.3786268234253,105.18949174880981,226.56150460243225,259.4150538444519,0.0,1832.2265625,0.0,1.2844314575195312,6137.015625,5754.8828125,-0.078125,0.0,-0.07940945494920015,382.1328125,0.0,-61.22801470756531,659.4850283179806,44.81633610565777,-0.1250619888305664,28.312485933303833,-1856.9243653122003,28.18913507461548,-32.85354924201965,-1832.2265625,-1.2844314575195312,382.1328125 +140,139,0.4296875,0.4921875,0.0,0.0,0.005276798270642757,0.08921639621257782,5566.1953125,5162.03125,7168.0,7168.0,70.62432599067688,128.37253999710083,1261.0261938890108,643.3813649076762,101.49477392458581,55.83748674102641,19.840506076812744,18.910818338394165,102.65329074859619,73.39111733436584,7160.033493714875,9309.955002961313,122.72747254371643,92.53623104095459,213.36836791038513,241.0462212562561,0.0,1745.5703125,0.0,1.122989535331726,5566.1953125,5162.03125,-0.0625,0.0,-0.08393959794193506,404.1640625,0.0,-57.74821400642395,617.6448289813345,45.657287183559404,0.9296877384185791,29.262173414230347,-2149.9215092464383,30.19124150276184,-27.67785334587097,-1745.5703125,-1.122989535331726,404.1640625 +141,140,0.3984375,0.484375,0.0,0.0,0.004780725575983524,0.08022038638591766,6181.25,5905.6484375,7168.0,7168.0,74.62091779708862,136.40001845359802,1325.3656336542492,692.744591029104,96.05885603674072,52.55131253841058,22.618371725082397,21.54559302330017,112.45236539840698,85.14542627334595,7225.388253251543,9128.323552046233,135.3031566143036,106.92821669578552,230.17918491363525,263.40882420539856,0.0,1806.0859375,0.0,1.3218774795532227,6181.25,5905.6484375,-0.0859375,0.0,-0.07543966080993414,275.6015625,0.0,-61.7791006565094,632.6210426251453,43.50754349833014,1.0727787017822266,27.306939125061035,-1902.9352987946904,28.374939918518066,-33.229639291763306,-1806.0859375,-1.3218774795532227,275.6015625 +142,141,0.453125,0.4921875,0.0,0.0,0.005041320808231831,0.1029646247625351,5901.7890625,5444.2734375,7168.0,7168.0,75.68108010292053,133.05736255645752,1247.7177237902026,654.6678314252553,94.71323599309184,53.871502202356886,20.635672330856323,19.988988399505615,107.57447385787964,77.07958626747131,7199.226472844819,9287.68607444013,128.4440860748291,97.30146670341492,224.1645221710205,250.36524200439453,0.0,1829.8515625,0.0,1.1916933059692383,5901.7890625,5444.2734375,-0.0390625,0.0,-0.09792330395430326,457.515625,0.0,-57.37628245353699,593.0498923649474,40.841733790734956,0.646683931350708,30.494887590408325,-2088.4596015953102,31.142619371414185,-26.200719833374023,-1829.8515625,-1.1916933059692383,457.515625 +143,142,0.484375,0.53125,0.0,0.0,0.00447797728702426,0.07975813001394272,6141.4140625,5777.6640625,7168.0,7168.0,80.22330284118652,134.11858892440796,1224.8638677283693,689.2603459472914,89.35059697292792,53.445238706172454,23.002623319625854,23.29643440246582,116.31926441192627,89.10913705825806,7003.93872088886,8620.137343466895,139.55543613433838,112.64115977287292,240.1136496067047,266.8887372016907,0.0,1732.5625,0.0,1.4024378061294556,6141.4140625,5777.6640625,-0.046875,0.0,-0.07528015272691846,363.75,0.0,-53.895286083221436,535.6035217810779,35.905358266755464,-0.2938110828399658,27.210127353668213,-1616.1986225780347,26.914276361465454,-26.775087594985962,-1732.5625,-1.4024378061294556,363.75 +144,143,0.5078125,0.515625,0.0,0.0,0.0055483621545135975,0.10059496760368347,5711.8828125,5258.265625,7168.0,7168.0,88.04528188705444,126.55463075637817,1037.9900324157786,664.7899764486481,81.41265319810319,56.63957104658332,22.767038345336914,19.455571174621582,109.64467239379883,75.04383897781372,6842.621566740459,9223.861804360024,132.64175295829773,94.73381328582764,240.65274381637573,241.16183161735535,0.0,1810.3046875,0.0,1.204777479171753,5711.8828125,5258.265625,-0.0078125,0.0,-0.09504660544916987,453.6171875,0.0,-38.50934886932373,373.20005596713054,24.773082151519866,3.311467170715332,34.60083341598511,-2381.240237619565,37.90793967247009,-0.5090878009796143,-1810.3046875,-1.204777479171753,453.6171875 +145,144,0.7265625,0.734375,0.0,0.0,0.00737136323004961,0.08143479377031326,4413.7109375,4153.203125,7168.0,7168.0,108.97623443603516,110.45659112930298,648.025465051748,601.6051131091912,65.77580916697337,64.89427137588356,16.246091842651367,14.892272472381592,82.51703977584839,57.083441495895386,7078.235011660595,9647.806536674922,98.99399662017822,72.20702362060547,228.0984182357788,202.85553550720215,0.0,1186.4296875,0.0,0.8659325838088989,4413.7109375,4153.203125,-0.0078125,0.0,-0.07406343054026365,260.5078125,0.0,-1.4803566932678223,46.42035194255686,0.8815377910898121,1.3538193702697754,25.433598279953003,-2569.5715250143267,26.786972999572754,25.24288272857666,-1186.4296875,-0.8659325838088989,260.5078125 +146,145,0.71875,0.6875,0.0,0.0,0.007275203708559275,0.07623441517353058,4579.859375,4286.5234375,7168.0,7168.0,95.50727128982544,110.57497644424438,767.2478651141864,620.2522234728447,75.05187723611176,64.82479337098793,17.216740131378174,15.73029899597168,86.76715040206909,60.390535831451416,6940.299378388251,9349.859083481317,104.21326518058777,76.35331559181213,219.89133405685425,207.00306630134583,0.0,1219.84375,0.0,0.9748073220252991,4579.859375,4286.5234375,0.03125,0.0,-0.0689592114649713,293.3359375,0.0,-15.067705154418945,146.99564164134176,10.22708386512383,1.4864411354064941,26.376614570617676,-2409.5597050930655,27.859949588775635,12.888267755508423,-1219.84375,-0.9748073220252991,293.3359375 +147,146,0.6875,0.625,0.0,0.0,0.006282266229391098,0.11164486408233643,5195.1953125,4941.78125,7168.0,7168.0,99.61715960502625,120.1426842212677,834.4257689094558,658.1216368895083,71.9554746232529,59.662392649715066,19.45773983001709,19.37446689605713,98.26133561134338,73.98255109786987,6971.216050933902,8820.512273721646,117.94897317886353,93.5908670425415,237.9932746887207,233.86782908439636,0.0,1843.109375,0.0,1.2461260557174683,5195.1953125,4941.78125,0.0625,0.0,-0.10536259785294533,253.4140625,0.0,-20.525524616241455,176.30413201994747,12.29308197353783,0.08327293395996094,24.27878451347351,-1849.2962227877442,24.35810613632202,4.125445604324341,-1843.109375,-1.2461260557174683,253.4140625 +148,147,0.6328125,0.671875,0.0,0.0,0.006139333359897137,0.0957472026348114,5393.59375,5032.5546875,7168.0,7168.0,87.28847980499268,121.72193670272827,988.6470722458843,661.5149017604746,82.1185111255656,58.888317045972016,19.567003965377808,19.133076190948486,100.04075264930725,73.42068815231323,7067.159945091595,9000.065466958999,119.83912634849548,92.78792476654053,227.367516040802,234.61328220367432,0.0,1752.671875,0.0,1.1437468528747559,5393.59375,5032.5546875,-0.0390625,0.0,-0.08960786927491426,361.0390625,0.0,-34.433456897735596,327.1321704854097,23.230194079593588,0.4339277744293213,26.62006449699402,-1932.9055218674039,27.051201581954956,-7.2457661628723145,-1752.671875,-1.1437468528747559,361.0390625 +149,148,0.4375,0.5234375,0.0,0.0,0.005077664740383625,0.09818699210882187,6158.015625,5750.7265625,7168.0,7168.0,75.10468602180481,133.1461889743805,1311.8788616118404,691.057143345684,95.44011671813588,53.835562663977115,21.939968585968018,22.803693532943726,113.96531319618225,87.41255187988281,7109.549188929377,8672.770485430385,136.13680958747864,110.45222735404968,231.3485701084137,263.4377279281616,0.0,1879.78125,0.0,1.4533660411834717,6158.015625,5750.7265625,-0.0859375,0.0,-0.09310932736843824,407.2890625,0.0,-58.041502952575684,620.8217182661564,41.60455405415877,-0.863724946975708,26.55276131629944,-1563.2212965010076,25.684582233428955,-32.089157819747925,-1879.78125,-1.4533660411834717,407.2890625 +150,149,0.7734375,0.828125,0.0,0.0,0.0062384032644331455,0.08007346838712692,4355.171875,4040.390625,7168.0,7168.0,91.83850574493408,107.39971899986267,758.7530898371981,601.9219659232317,78.05005037764778,66.74132918363749,16.52606225013733,14.819783449172974,82.8162624835968,58.34015464782715,6937.647060730368,9157.637706397445,99.57059144973755,73.39154243469238,211.6320559978485,200.7787160873413,0.0,1121.1171875,0.0,0.9179511070251465,4355.171875,4040.390625,-0.0546875,0.0,-0.07383506512269378,314.78125,0.0,-15.561213254928589,156.8311239139664,11.308721194010289,1.7062788009643555,24.476107835769653,-2219.990645667077,26.179049015045166,10.853339910507202,-1121.1171875,-0.9179511070251465,314.78125 +151,150,0.8828125,0.875,0.0,0.0,0.007502119988203049,0.08231836557388306,3733.890625,3559.25,7168.0,7168.0,88.98360872268677,100.75278449058533,671.3848860207942,565.2250733112137,80.55416163597877,71.14443572197055,15.613138914108276,13.43297791481018,75.00430226325989,51.76145315170288,6669.510746786035,9232.507414337886,90.84495854377747,65.4232292175293,199.75356698036194,186.46449184417725,0.0,1220.4765625,0.0,0.7871733903884888,3733.890625,3559.25,0.0078125,0.0,-0.07481624558568001,174.640625,0.0,-11.76917576789856,106.15981270958048,9.40972591400822,2.1801609992980957,23.242849111557007,-2562.996667551851,25.42172932624817,13.289075136184692,-1220.4765625,-0.7871733903884888,174.640625 +152,151,0.6875,0.7265625,0.0,0.0,0.0059100184589624405,0.08873619139194489,5354.8203125,4992.609375,7168.0,7168.0,89.68351411819458,122.34434795379639,955.3274739779429,652.9255444654257,79.9255032597545,58.588730251004414,20.22316336631775,18.78895926475525,101.75820875167847,71.96219205856323,6909.64403388638,9126.320102443979,122.2113356590271,90.98433828353882,231.85794615745544,233.15596652030945,0.0,1504.5234375,0.0,1.2065268754959106,5354.8203125,4992.609375,-0.0390625,0.0,-0.08282617293298244,362.2109375,0.0,-32.66083383560181,302.4019295125172,21.336773008750093,1.4342041015625,29.796016693115234,-2216.6760685575982,31.22699737548828,-1.298020362854004,-1504.5234375,-1.2065268754959106,362.2109375 +153,152,0.8046875,0.78125,0.0,0.0,0.006254359148442745,0.08744889497756958,4803.5859375,4387.203125,7168.0,7168.0,87.42914175987244,112.28371977806091,879.0818879486635,625.1596414756062,81.98639327476407,63.838284073311875,18.325032711029053,16.306845664978027,91.34886431694031,62.51698684692383,6942.472736165067,9291.714609060416,109.9041793346405,79.05391931533813,217.354745388031,213.09977269172668,0.0,1431.640625,0.0,1.0314526557922363,4803.5859375,4387.203125,0.0234375,0.0,-0.08119453582912683,416.3828125,0.0,-24.854578018188477,253.92224647305727,18.148109201452193,2.0181870460510254,28.83187747001648,-2349.24187289535,30.850260019302368,4.254972696304321,-1431.640625,-1.0314526557922363,416.3828125 +154,153,0.6171875,0.6640625,0.0,0.0,0.005628013052046299,0.07898208498954773,5552.3046875,5293.6484375,7168.0,7168.0,92.1932966709137,124.40653252601624,963.5936473462433,680.8193531339493,77.74968743754069,57.61755314979949,20.482747077941895,19.877278327941895,103.93482780456543,76.47408127784729,7024.931059422299,9114.552124759059,124.65035223960876,96.58385109901428,236.99006724357605,241.02419209480286,0.0,1478.890625,0.0,1.0990629196166992,5552.3046875,5293.6484375,-0.046875,0.0,-0.07335407193750143,258.65625,0.0,-32.21323585510254,282.77429421229397,20.132134287741195,0.60546875,27.46074652671814,-2089.62106533676,28.066501140594482,-4.034124851226807,-1478.890625,-1.0990629196166992,258.65625 +155,154,0.3671875,0.375,0.0,0.0,0.005701000802218914,0.08769209682941437,5594.09375,5412.3359375,7168.0,7168.0,71.57802367210388,131.31523537635803,1250.460621964378,659.461750586718,100.14246876717813,54.58620227467166,19.83438205718994,20.198233604431152,102.72650384902954,77.28265857696533,7142.343723468706,9192.786752961845,122.79372882843018,97.71449327468872,214.64751148223877,248.95665049552917,0.0,1516.625,0.0,1.2124712467193604,5594.09375,5412.3359375,-0.0078125,0.0,-0.08199109602719545,181.7578125,0.0,-59.73721170425415,590.9988713776601,45.55626649250647,-0.36385154724121094,25.44384527206421,-2050.443029493139,25.079235553741455,-34.309139013290405,-1516.625,-1.2124712467193604,181.7578125 +156,155,0.3671875,0.390625,0.0,0.0,0.005012281704694033,0.10173424333333969,5952.875,5609.7578125,7168.0,7168.0,72.26680755615234,132.1189570426941,1317.977135298145,679.3584131230722,99.18799850720349,54.25413703261122,21.314120054244995,21.086573362350464,110.34638714790344,81.44296503067017,7046.030414732726,9007.346426100072,131.89385986328125,102.76483869552612,224.297123670578,254.8620262145996,0.0,1919.953125,0.0,1.3119258880615234,5952.875,5609.7578125,-0.0234375,0.0,-0.09672196162864566,343.1171875,0.0,-59.85214948654175,638.6187221750728,44.93386147459227,0.22754669189453125,28.903422117233276,-1961.3160113673457,29.129021167755127,-30.564902544021606,-1919.953125,-1.3119258880615234,343.1171875 +157,156,0.3203125,0.34375,0.0,0.0,0.004776536487042904,0.09939177334308624,6677.7578125,6122.84375,7168.0,7168.0,78.67833352088928,139.29703617095947,1357.9866301010638,703.2848845381519,91.10513249618943,51.45838129106137,23.62806463241577,24.64979386329651,122.72823548316956,95.44939994812012,7109.309414944313,8396.95168786428,146.59108901023865,120.33655166625977,245.44176626205444,279.54559540748596,0.0,1971.640625,0.0,1.5118088722229004,6677.7578125,6122.84375,-0.0234375,0.0,-0.09461523685604334,554.9140625,0.0,-60.61870265007019,654.7017455629119,39.646751205128055,-1.0217292308807373,27.27883553504944,-1287.6422729199676,26.254537343978882,-34.10382914543152,-1971.640625,-1.5118088722229004,554.9140625 +158,157,0.484375,0.515625,0.0,0.0,0.005132999271154404,0.09950485825538635,5461.4296875,5112.921875,7168.0,7168.0,93.46647500991821,125.39221739768982,934.9114213489634,652.4069172534402,76.69059948221398,57.164632293455725,19.749046325683594,18.688716173171997,100.8894202709198,72.41037106513977,7108.525334709623,9288.255122943165,120.86946058273315,91.33299827575684,234.50493335723877,236.65158677101135,0.0,1694.3125,0.0,1.591546893119812,5461.4296875,5112.921875,-0.03125,0.0,-0.09437185898423195,348.5078125,0.0,-31.925742387771606,282.5045040955232,19.52596718875826,1.0603301525115967,28.47904920578003,-2179.7297882335415,29.53646230697632,-2.146653413772583,-1694.3125,-1.591546893119812,348.5078125 +159,158,0.5625,0.6328125,0.0,0.0,0.007353040389716625,0.09238594770431519,5221.25,4839.6640625,7168.0,7168.0,94.15075039863586,121.16070413589478,887.3004160486266,639.1067595080062,76.1332221957931,59.161095597136146,19.867631673812866,17.8976628780365,97.51997375488281,69.55385756492615,7023.46374417175,9145.21526582816,117.61779761314392,87.68497037887573,231.9200189113617,228.79486083984375,0.0,1561.5859375,0.0,1.0840492248535156,5221.25,4839.6640625,-0.0703125,0.0,-0.08503290731459856,381.5859375,0.0,-27.00995373725891,248.1936565406204,16.97212659865695,1.9699687957763672,27.966116189956665,-2121.7515216564097,29.93282723426819,3.1251580715179443,-1561.5859375,-1.0840492248535156,381.5859375 +160,159,0.4453125,0.4453125,0.0,0.0,0.0048034656792879105,0.07856027781963348,5830.53125,5589.4296875,7168.0,7168.0,71.33686137199402,132.13656997680664,1307.7180325265042,676.806390658524,100.48101166971259,54.24690531363246,20.72514033317566,20.9065420627594,107.84582018852234,80.8251736164093,7103.36291810714,9096.26255168027,128.80459117889404,101.96842956542969,220.45350980758667,254.12358212471008,0.0,1532.578125,0.0,1.3107150793075562,5830.53125,5589.4296875,0.0,0.0,-0.07375681214034557,241.1015625,0.0,-60.79970860481262,630.9116418679802,46.234106356080126,-0.18140172958374023,27.020646572113037,-1992.899633573129,26.836161613464355,-33.67007231712341,-1532.578125,-1.3107150793075562,241.1015625 +161,160,0.6171875,0.703125,0.0,0.0,0.005742167588323355,0.08357317745685577,5556.453125,5092.859375,7168.0,7168.0,94.05539727210999,121.65843796730042,945.222204981981,669.7911904959843,76.21040586604921,58.91905337405884,20.673632621765137,19.396034240722656,105.11107730865479,74.99265789985657,6960.65551541785,8964.904283000295,126.01545476913452,94.62386536598206,240.16348218917847,236.31707191467285,0.0,1468.6875,0.0,1.1496686935424805,5556.453125,5092.859375,-0.0859375,0.0,-0.07783100986853242,463.59375,0.0,-27.60304069519043,275.4310144859967,17.291352491990374,1.2775983810424805,30.118419408798218,-2004.248767582444,31.391589403152466,3.8464102745056152,-1468.6875,-1.1496686935424805,463.59375 +162,161,0.6328125,0.71875,0.0,0.0,0.0067531997337937355,0.09612417221069336,5018.9609375,4303.453125,7168.0,7168.0,93.10859608650208,110.80879545211792,862.4700443919759,621.3879477622636,76.98537300831607,64.68800577385029,18.45706605911255,16.12797498703003,94.8957142829895,63.11560082435608,6969.956493794603,9028.417579130502,113.58271908760071,79.47714972496033,226.75879883766174,210.3507137298584,0.0,1463.3359375,0.0,1.0113106966018677,5018.9609375,4303.453125,-0.0859375,0.0,-0.08937097247689962,715.5078125,0.0,-17.700199365615845,241.08209662971228,12.297367234465781,2.3290910720825195,31.780113458633423,-2058.461085335899,34.10556936264038,16.408085107803345,-1463.3359375,-1.0113106966018677,715.5078125 +163,162,0.7421875,0.7734375,0.0,0.0,0.005908493883907795,0.09060955792665482,4709.515625,4097.953125,7168.0,7168.0,95.84719443321228,108.9345543384552,786.1706380201514,601.8957932878234,74.78570491695265,65.80097604043358,17.966858625411987,15.438594579696655,90.39588236808777,59.45525026321411,6968.658123551714,9278.541382935182,108.59301614761353,75.1269416809082,224.74041271209717,204.02146744728088,0.0,1339.2265625,0.0,0.8713599443435669,4709.515625,4097.953125,-0.03125,0.0,-0.08470106404274702,611.5625,0.0,-13.08735990524292,184.274844732328,8.984728876519071,2.528264045715332,30.940632104873657,-2309.8832593834677,33.46607446670532,20.718945264816284,-1339.2265625,-0.8713599443435669,611.5625 +164,163,0.515625,0.609375,0.0,0.0,0.005700146779417992,0.08820351958274841,5266.65625,4890.3203125,7168.0,7168.0,80.43529176712036,122.07605242729187,1047.6309359823288,640.953925395012,89.11511157009407,58.717495016225556,19.188703775405884,18.197428226470947,98.61771512031555,70.30926775932312,7033.74641314424,9180.596250974442,118.0411446094513,88.74155497550964,218.7783088684082,230.8436713218689,0.0,1540.53125,0.0,1.0596423149108887,5266.65625,4890.3203125,-0.09375,0.0,-0.08250337280333042,376.3359375,0.0,-41.64076066017151,406.6770105873168,30.397616553868517,0.9912755489349365,28.30844736099243,-2146.849837830202,29.29958963394165,-12.065362453460693,-1540.53125,-1.0596423149108887,376.3359375 +165,164,0.7109375,0.7890625,0.0,0.0,0.006104565225541592,0.08450207859277725,5653.28125,4922.3671875,7168.0,7168.0,89.65393853187561,118.90577721595764,1008.9071543448197,662.3553274199563,79.95186957069917,60.28302549994203,21.162323713302612,18.45285415649414,106.72725009918213,71.37631058692932,6940.945253546609,9067.86852217216,128.12253856658936,90.06270456314087,237.96326208114624,229.16686987876892,0.0,1467.5390625,0.0,1.1262708902359009,5653.28125,4922.3671875,-0.078125,0.0,-0.07839751336723566,730.9140625,0.0,-29.25183868408203,346.55182692486335,19.668844070757146,2.7094695568084717,35.35093951225281,-2126.923268625552,38.059834003448486,8.79639220237732,-1467.5390625,-1.1262708902359009,730.9140625 +166,165,0.59375,0.6171875,0.0,0.0,0.0064134253188967705,0.10261388123035431,5082.5546875,4637.296875,7168.0,7168.0,92.43616199493408,116.511070728302,879.7517469889841,636.8214585635653,77.54540912670994,61.52205069607006,18.64592719078064,17.231587409973145,95.12589144706726,67.19580078125,7045.221756191613,9125.421423225329,114.00452470779419,84.66015362739563,226.56932711601257,221.02718257904053,0.0,1597.0859375,0.0,1.037354588508606,5082.5546875,4637.296875,-0.0234375,0.0,-0.09620045591145754,445.2578125,0.0,-24.07490873336792,242.93028842541878,16.02335843063988,1.4143397808074951,27.93009066581726,-2080.1996670337157,29.34437108039856,5.542144536972046,-1597.0859375,-1.037354588508606,445.2578125 +167,166,0.84375,0.8984375,0.0,0.0,0.007827508263289928,0.1060924082994461,4514.359375,3935.3671875,7168.0,7168.0,87.42146944999695,105.53861927986145,826.2243869203519,596.6145419529375,81.99358858981351,67.91826583397207,17.3159499168396,15.217621326446533,85.59282064437866,57.71514582633972,6913.6406014546255,8968.997523761936,103.13905715942383,73.16320896148682,211.10106110572815,198.79732584953308,0.0,1404.2109375,0.0,0.9659345149993896,4514.359375,3935.3671875,-0.0546875,0.0,-0.09826490003615618,578.9921875,0.0,-18.117149829864502,229.60984496741435,14.07532275584144,2.0983285903930664,27.87767481803894,-2055.356922307311,29.97584819793701,12.303735256195068,-1404.2109375,-0.9659345149993896,578.9921875 +168,167,0.8046875,0.8515625,0.0,0.0,0.006905543152242899,0.07729309797286987,4515.3515625,4215.890625,7168.0,7168.0,92.97843718528748,108.71351194381714,777.0148346979513,620.4771494720931,77.09314349643893,65.9347662662614,17.255820751190186,15.785837888717651,86.09019660949707,60.18756723403931,6913.83018556539,9252.442416131636,103.57556986808777,76.21735978126526,216.66268944740295,206.56956458091736,0.0,1297.4453125,0.0,0.9473426342010498,4515.3515625,4215.890625,-0.046875,0.0,-0.07038755482062697,299.4609375,0.0,-15.735074758529663,156.53768522585824,11.158377230177521,1.4699828624725342,25.902629375457764,-2338.6122305662466,27.35821008682251,10.093124866485596,-1297.4453125,-0.9473426342010498,299.4609375 +169,168,0.4609375,0.46875,0.0,0.0,0.005556054413318634,0.09081864356994629,5546.828125,5391.8984375,7168.0,7168.0,77.76638555526733,130.0097041130066,1141.2289431521453,663.5687358000011,92.17350078467793,55.134345923666245,20.527775287628174,20.56169366836548,105.12167644500732,78.77190351486206,6974.565330334605,9055.85580860581,125.8807475566864,99.56783628463745,223.72459030151367,249.40150260925293,0.0,1597.65625,0.0,1.2833220958709717,5546.828125,5391.8984375,-0.0078125,0.0,-0.08526258915662766,154.9296875,0.0,-52.24331855773926,477.6602073521442,37.03915486101168,-0.03391838073730469,26.349772930145264,-2081.290478271204,26.31291127204895,-25.676912307739258,-1597.65625,-1.2833220958709717,154.9296875 +170,169,0.5,0.5703125,0.0,0.0,0.0056703053414821625,0.10147568583488464,5948.140625,5337.421875,7168.0,7168.0,81.43247437477112,127.11685180664062,1168.7014392072192,671.8129719724442,88.02385111142765,56.38906170287598,21.550392627716064,20.28438091278076,110.30896735191345,77.22998666763306,7117.481188045779,9153.827813572334,132.09586644172668,97.74755311012268,233.71852469444275,244.89916229248047,0.0,1866.0,0.0,1.2649002075195312,5948.140625,5337.421875,-0.0703125,0.0,-0.09580538049340248,610.71875,0.0,-45.68437743186951,496.888467234775,31.63478940855167,1.2660117149353027,33.078980684280396,-2036.3466255265548,34.348313331604004,-11.18063759803772,-1866.0,-1.2649002075195312,610.71875 +171,170,0.6640625,0.640625,0.0,0.0,0.006801513023674488,0.09428280591964722,4778.8203125,4488.125,7168.0,7168.0,95.79051041603088,113.94930982589722,798.2118966473735,630.1924962048324,74.82995934428604,62.90516380443167,18.103414297103882,16.92394495010376,90.98714470863342,64.99991941452026,6960.554724824842,9170.965216102024,109.32164335250854,82.1548285484314,225.26324653625488,216.7079918384552,0.0,1449.4140625,0.0,1.076395034790039,4778.8203125,4488.125,0.0234375,0.0,-0.08748129289597273,290.6953125,0.0,-18.158799409866333,168.01940044254104,11.924795539854365,1.179469347000122,25.98722529411316,-2210.410491277182,27.16681480407715,8.555254697799683,-1449.4140625,-1.076395034790039,290.6953125 +172,171,0.546875,0.578125,0.0,0.0,0.007133251056075096,0.09226879477500916,4911.5546875,4594.28125,7168.0,7168.0,103.94375228881836,115.5799310207367,756.0326933517263,635.997091803174,68.96037368444212,62.01768712523248,17.911362409591675,16.49495577812195,91.19301080703735,63.20783972740173,7071.3204256903055,9559.63694702968,109.33687973022461,79.9341025352478,233.47854661941528,215.44447922706604,0.0,1551.453125,0.0,1.090036392211914,4911.5546875,4594.28125,-0.03125,0.0,-0.08513554371893406,317.2734375,0.0,-11.636178731918335,120.03560154855222,6.9426865592096405,1.4164066314697266,27.98517107963562,-2488.316521339374,29.402777194976807,18.034067392349243,-1551.453125,-1.090036392211914,317.2734375 +173,172,0.671875,0.6875,0.0,0.0,0.006058076396584511,0.08989371359348297,5361.4609375,5158.7734375,7168.0,7168.0,90.41470980644226,124.28674554824829,948.7767552828858,664.1124492873435,79.27913516888005,57.6730846751263,19.90228581428528,19.466882944107056,100.73955464363098,74.90255212783813,7043.3803535269,9126.5648576737,120.87479138374329,94.60198998451233,231.4064450263977,239.23202919960022,0.0,1572.140625,0.0,1.1827881336212158,5361.4609375,5158.7734375,-0.015625,0.0,-0.08383563719689846,202.6875,0.0,-33.87203574180603,284.66430599554235,21.60605049375375,0.43540287017822266,25.837002515792847,-2083.1845041468005,26.272801399230957,-7.825584173202515,-1572.140625,-1.1827881336212158,202.6875 +174,173,0.8203125,0.8515625,0.0,0.0,0.005761880427598953,0.07521670311689377,4720.078125,4331.921875,7168.0,7168.0,86.34935355186462,110.47644352912903,874.6012204322889,627.3803517373821,83.01162319292446,64.88261000282864,18.24000096321106,16.91819405555725,91.11133456230164,67.22088289260864,6843.078339212176,8536.00808719953,109.58246159553528,84.3700499534607,215.57131266593933,214.82409954071045,0.0,1241.859375,0.0,0.9672491550445557,4720.078125,4331.921875,-0.03125,0.0,-0.06945482268929482,388.15625,0.0,-24.127089977264404,247.2208686949068,18.12901319009582,1.3218069076538086,23.890451669692993,-1692.9297479873549,25.212411642074585,0.7472131252288818,-1241.859375,-0.9672491550445557,388.15625 +175,174,0.59375,0.609375,0.0,0.0,0.005803982261568308,0.07650230079889297,4824.015625,4693.984375,7168.0,7168.0,101.70473146438599,118.01141691207886,758.9052041991542,636.410882651752,70.4785303180317,60.73988591578661,17.90470266342163,17.121637105941772,90.85680556297302,65.97699642181396,6969.758578636063,9345.772518315665,108.99129343032837,83.33178448677063,230.8166480064392,221.45805096626282,0.0,1195.0625,0.0,1.1599977016448975,4824.015625,4693.984375,-0.015625,0.0,-0.07069831853732467,130.03125,0.0,-16.30668544769287,122.49432154740225,9.738644402245093,0.7830655574798584,24.879809141159058,-2376.013939679602,25.65950894355774,9.358597040176392,-1195.0625,-1.1599977016448975,130.03125 +176,175,0.6484375,0.6015625,0.0,0.0,0.005933180917054415,0.09528617560863495,5273.09375,4910.703125,7168.0,7168.0,94.11801552772522,119.56982088088989,896.4224280222578,657.1160634109268,76.15970183613206,59.94823733273332,19.51999592781067,18.30510973930359,98.22197461128235,70.49502158164978,7043.596941906029,9155.965705357185,117.97272968292236,89.03493618965149,231.8062961101532,228.73211812973022,0.0,1537.4921875,0.0,1.0967929363250732,5273.09375,4910.703125,0.046875,0.0,-0.08935299469158053,362.390625,0.0,-25.451805353164673,239.306364611331,16.211464503398737,1.21488618850708,27.72695302963257,-2112.3687634511552,28.937793493270874,3.0741779804229736,-1537.4921875,-1.0967929363250732,362.390625 +177,176,0.34375,0.328125,0.0,0.0,0.0054439580999314785,0.09182043373584747,5786.640625,5588.9765625,7168.0,7168.0,70.86213183403015,132.38761258125305,1306.5687921561691,675.4682198465974,101.15416816401378,54.14403855648225,21.553936004638672,20.441225290298462,107.80754470825195,78.5887839794159,7046.352850867346,9344.195988480264,129.59468054771423,99.2651526927948,220.04703950881958,251.67746472358704,0.0,1675.9375,0.0,1.3023254871368408,5786.640625,5588.9765625,0.015625,0.0,-0.086376475635916,197.6640625,0.0,-61.5254807472229,631.1005723095717,47.01012960753153,1.11271071434021,29.21876072883606,-2297.8431376129174,30.329527854919434,-31.630425214767456,-1675.9375,-1.3023254871368408,197.6640625 +178,177,0.2421875,0.25,0.0,0.0,0.004558555781841278,0.1036556214094162,6568.875,6493.890625,7168.0,7168.0,78.8805480003357,148.1012146472931,1332.4197494108778,701.5624432753366,90.87157964431859,48.3993329634114,22.609651565551758,24.249671459197998,119.97082567214966,94.31101894378662,7182.679582074762,9035.14784956248,142.81506657600403,118.80502963066101,241.50240993499756,286.9455659389496,0.0,2078.0859375,0.0,1.4084525108337402,6568.875,6493.890625,-0.0078125,0.0,-0.09909706562757492,74.984375,0.0,-69.2206666469574,630.8573061355412,42.47224668090718,-1.6400198936462402,25.659806728363037,-1852.4682674877176,24.010036945343018,-45.443156003952026,-2078.0859375,-1.4084525108337402,74.984375 +179,178,0.5859375,0.609375,0.0,0.0,0.007001928985118866,0.0974753275513649,5110.6953125,4830.421875,7168.0,7168.0,83.16172575950623,120.61547374725342,983.2783561571621,640.7697752110345,86.1934974837944,59.42852751231867,18.79154634475708,18.064552783966064,95.96241593360901,70.91940307617188,6989.830273386007,8952.218609596766,114.98618841171265,89.21872639656067,217.7766785621643,229.9150128364563,0.0,1578.515625,0.0,1.153435230255127,5110.6953125,4830.421875,-0.0234375,0.0,-0.09047339856624603,280.2734375,0.0,-37.45374798774719,342.5085809461276,26.764969971475736,0.7269935607910156,25.043012857437134,-1962.388336210759,25.767462015151978,-12.138334274291992,-1578.515625,-1.153435230255127,280.2734375 +180,179,0.734375,0.703125,0.0,0.0,0.0068245758302509785,0.09525999426841736,4927.765625,4788.7734375,7168.0,7168.0,99.10103178024292,118.75850176811218,795.5946429986476,645.1780197565049,72.33022574270548,60.357784017823285,18.415006399154663,17.951902866363525,93.04705572128296,69.0290961265564,6986.250074878102,9159.311587114375,111.69261193275452,87.21474146842957,230.61319875717163,225.90705704689026,0.0,1573.5,0.0,1.1446653604507446,4927.765625,4788.7734375,0.03125,0.0,-0.08843541843816638,138.9921875,0.0,-19.657469987869263,150.4166232421427,11.972441724882195,0.4631035327911377,24.017959594726562,-2173.0615122362724,24.47787046432495,4.706141710281372,-1573.5,-1.1446653604507446,138.9921875 +181,180,0.453125,0.53125,0.0,0.0,0.005283984821289778,0.08748387545347214,5665.5625,5372.515625,7168.0,7168.0,69.83294200897217,128.85311031341553,1298.0836463735607,667.1181610666192,102.6449665987014,55.62923535617366,20.08610773086548,19.88273286819458,104.5425591468811,76.83043599128723,7115.111836461976,9193.257735516414,124.8604576587677,96.95302677154541,214.39603233337402,245.72478532791138,0.0,1622.7109375,0.0,1.3032158613204956,5665.5625,5372.515625,-0.078125,0.0,-0.08219989063218236,293.046875,0.0,-59.02016830444336,630.9654853069414,47.01573124252774,0.20337486267089844,27.712123155593872,-2078.145899054438,27.90743088722229,-31.328752994537354,-1622.7109375,-1.3032158613204956,293.046875 +182,181,0.5390625,0.5703125,0.0,0.0,0.005462354980409145,0.0878094807267189,5408.09375,5245.8125,7168.0,7168.0,100.41619563102722,125.66843914985657,861.7086064278617,667.8924363810386,71.38290745786018,57.03898328403947,20.04917025566101,20.04979395866394,101.69772863388062,77.08969068527222,6999.841683414364,8964.830366507516,121.98018765449524,97.37393641471863,242.2107129096985,242.91812539100647,0.0,1595.421875,0.0,1.3188796043395996,5408.09375,5245.8125,-0.03125,0.0,-0.08234712574630976,162.28125,0.0,-25.252243518829346,193.81617004682312,14.343924173820717,-0.0006237030029296875,24.6080379486084,-1964.9886830931528,24.60625123977661,-0.7074124813079834,-1595.421875,-1.3188796043395996,162.28125 +183,182,0.484375,0.5,0.0,0.0,0.005822153761982918,0.09401750564575195,5265.0625,4972.2890625,7168.0,7168.0,90.86167311668396,123.16102910041809,927.1345894304441,645.9561565950729,78.8891482417994,58.20022820819112,18.9018657207489,18.432043075561523,97.6961898803711,70.537930727005,7117.984855412649,9327.251213907775,116.8300724029541,89.20469379425049,227.7423713207245,233.6659231185913,0.0,1612.6015625,0.0,1.0845253467559814,5265.0625,4972.2890625,-0.015625,0.0,-0.08819535188376904,292.7734375,0.0,-32.29935598373413,281.17843283537115,20.688920033608284,0.46982264518737793,27.15825915336609,-2209.266358495126,27.625378608703613,-5.923551797866821,-1612.6015625,-1.0845253467559814,292.7734375 +184,183,0.65625,0.671875,0.0,0.0,0.005484824068844318,0.0701482743024826,5251.0546875,5049.8359375,7168.0,7168.0,98.7365562915802,122.7438108921051,850.9196406636736,658.260277343213,72.59722507266798,58.39805647146521,19.40987753868103,18.88925838470459,98.64990520477295,72.65618443489075,7023.372182282419,9181.585919885543,118.29294657707214,91.77870059013367,237.0756175518036,234.61360383033752,0.0,1261.5,0.0,1.156051516532898,5251.0546875,5049.8359375,-0.015625,0.0,-0.06466345023363829,201.21875,0.0,-24.007254600524902,192.65936332046056,14.199168601202771,0.5206191539764404,25.993720769882202,-2158.213737603124,26.514245986938477,2.4620137214660645,-1261.5,-1.156051516532898,201.21875 +185,184,0.703125,0.75,0.0,0.0,0.00688868947327137,0.0858069583773613,4572.8359375,4196.4140625,7168.0,7168.0,100.64796948432922,109.86012601852417,726.9433787374296,611.1646457485283,71.2185256863632,65.24660274639919,16.80819296836853,15.110971450805664,85.35909533500671,58.061530351638794,7047.62623876223,9531.216222659903,102.39770197868347,73.40410041809082,223.0101833343506,203.28187441825867,0.0,1190.8125,0.0,0.9265007972717285,4572.8359375,4196.4140625,-0.046875,0.0,-0.07891826890408993,376.421875,0.0,-9.212156534194946,115.77873298890131,5.971922939964003,1.6972215175628662,27.29756498336792,-2483.589983897673,28.99360156059265,19.72830891609192,-1190.8125,-0.9265007972717285,376.421875 +186,185,0.65625,0.671875,0.0,0.0,0.00620269775390625,0.09702479839324951,5768.7265625,5332.2109375,7168.0,7168.0,86.79524326324463,127.22493600845337,1063.4180115154575,670.5868965367865,82.5851709207139,56.34115626140875,22.107658624649048,21.04645085334778,111.22637701034546,80.3336443901062,6829.432194211868,8760.202594352506,133.56717896461487,101.6150631904602,240.42227506637573,248.94965720176697,0.0,1691.625,0.0,1.315413475036621,5768.7265625,5332.2109375,-0.015625,0.0,-0.09082210063934326,436.515625,0.0,-40.42969274520874,392.83111497867094,26.244014659305158,1.0612077713012695,30.892732620239258,-1930.7704001406382,31.952115774154663,-8.527382135391235,-1691.625,-1.315413475036621,436.515625 +187,186,0.53125,0.625,0.0,0.0,0.006593992467969656,0.09728845208883286,5422.796875,5083.6640625,7168.0,7168.0,91.34980463981628,122.77615427970886,949.8077236410661,662.4952986774141,78.4676007602069,58.38267245013921,19.94192409515381,18.448948860168457,101.72507691383362,70.9775002002716,7048.70442720197,9490.627284692991,121.89726567268372,89.65934252738953,234.39667963981628,232.51415181159973,0.0,1745.09375,0.0,1.1388708353042603,5422.796875,5083.6640625,-0.09375,0.0,-0.0906944596208632,339.1328125,0.0,-31.426349639892578,287.312424963652,20.084928310067696,1.4929752349853516,30.74757671356201,-2441.922857491021,32.23792314529419,1.8825278282165527,-1745.09375,-1.1388708353042603,339.1328125 +188,187,0.7890625,0.8515625,0.0,0.0,0.007516460493206978,0.09868721663951874,4810.21875,4077.2265625,7168.0,7168.0,83.40718722343445,107.71118474006653,922.744220996533,605.6532119429337,85.9398361054675,66.54833495052664,19.823633909225464,16.029191970825195,94.91748237609863,61.389888048172,6709.9546264448445,8846.228870361509,114.96934819221497,77.652179479599,218.52576684951782,205.3826642036438,0.0,1319.4921875,0.0,1.0736347436904907,4810.21875,4077.2265625,-0.0625,0.0,-0.09117075614631176,732.9921875,0.0,-24.30399751663208,317.09100905359935,19.39150115494087,3.7944419384002686,33.527594327926636,-2136.274243916664,37.31716871261597,13.143102645874023,-1319.4921875,-1.0736347436904907,732.9921875 +189,188,0.546875,0.6171875,0.0,0.0,0.005758625455200672,0.089149110019207,5729.4921875,5263.5546875,7168.0,7168.0,80.25988721847534,124.95779299736023,1142.187936926202,673.962567518931,89.30986883258377,57.363369086963836,20.667495489120483,19.678616046905518,106.18293237686157,75.61945247650146,7078.793014797095,9151.177075965199,127.08278155326843,95.53341603279114,227.7545416355133,240.5648455619812,0.0,1554.7890625,0.0,1.2373532056808472,5729.4921875,5263.5546875,-0.0703125,0.0,-0.08339048456400633,465.9375,0.0,-44.69790577888489,468.22536940727093,31.94649974561994,0.9888794422149658,30.563479900360107,-2072.3840611681044,31.549365520477295,-12.810303926467896,-1554.7890625,-1.2373532056808472,465.9375 +190,189,0.578125,0.625,0.0,0.0,0.007480373606085777,0.09731416404247284,4935.3671875,4651.9140625,7168.0,7168.0,92.43439769744873,116.74382734298706,854.2910103494895,637.5551212770063,77.5468892377263,61.39939184057075,18.18415641784668,17.67472743988037,91.87903690338135,67.72187399864197,7055.004295284952,9035.854501195152,110.29376649856567,85.63142013549805,222.88795471191406,222.17599606513977,0.0,1529.5390625,0.0,1.0571945905685425,4935.3671875,4651.9140625,-0.046875,0.0,-0.08983379043638706,283.453125,0.0,-24.30942964553833,216.73588907248325,16.147497397155554,0.5094289779663086,24.15716290473938,-1980.8502059102002,24.662346363067627,0.711958646774292,-1529.5390625,-1.0571945905685425,283.453125 +191,190,0.59375,0.59375,0.0,0.0,0.005703936330974102,0.08725302666425705,5423.9296875,5074.9375,7168.0,7168.0,90.87448644638062,122.66072130203247,954.9751354160904,661.9804541998445,78.8780248483648,58.437614942357484,20.64192533493042,18.55211853981018,101.51350927352905,71.03157424926758,7026.2076949593775,9412.49024910769,122.38653779029846,89.82078385353088,233.48044180870056,232.47952270507812,0.0,1528.4296875,0.0,1.1296298503875732,5423.9296875,5074.9375,0.0,0.0,-0.08154909033328295,348.9921875,0.0,-31.786234855651855,292.9946812162459,20.440409906007318,2.0898067951202393,30.481935024261475,-2386.2825541483126,32.56575393676758,1.0009191036224365,-1528.4296875,-1.1296298503875732,348.9921875 +192,191,0.8828125,0.8515625,0.0,0.0,0.006961008533835411,0.09257376194000244,4015.7109375,3694.8515625,7168.0,7168.0,86.77899289131165,101.52577590942383,740.402404536696,582.2917822636663,82.6006359508888,70.60276009508095,15.857937097549438,13.633776903152466,77.67614030838013,52.20889163017273,6837.144558052968,9385.623496301348,93.76181054115295,66.0728325843811,200.44081568717957,187.56453680992126,0.0,1225.84375,0.0,0.8284182548522949,4015.7109375,3694.8515625,0.03125,0.0,-0.08561275340616703,320.859375,0.0,-14.746783018112183,158.11062227302978,11.997875855807848,2.2241601943969727,25.467248678207397,-2548.4789382483805,27.68897795677185,12.8762788772583,-1225.84375,-0.8284182548522949,320.859375 +193,192,0.5546875,0.5703125,0.0,0.0,0.005630863830447197,0.08972897380590439,5877.90625,5509.0,7168.0,7168.0,80.89750719070435,127.89293718338013,1162.5389120866082,689.2014675807636,88.60594410038448,56.04687919335308,21.690247535705566,22.26463222503662,110.56558680534363,85.70121765136719,7011.693442778636,8495.00182088003,132.48674988746643,108.20107650756836,233.42190027236938,256.1653161048889,0.0,1781.4296875,0.0,1.212493896484375,5877.90625,5509.0,-0.015625,0.0,-0.08409810997545719,368.90625,0.0,-46.99542999267578,473.3374445058446,32.5590649070314,-0.5743846893310547,24.86436915397644,-1483.3083781013938,24.28567337989807,-22.74341583251953,-1781.4296875,-1.212493896484375,368.90625 +194,193,0.765625,0.71875,0.0,0.0,0.007334624417126179,0.08738933503627777,4402.953125,4080.7109375,7168.0,7168.0,102.8233597278595,107.43995428085327,685.1288480210267,607.7010683504683,69.71178552200006,66.7163351658034,16.710018157958984,15.108401775360107,83.78855204582214,59.39078092575073,6924.597523569808,9074.724925300978,100.72804236412048,74.7322461605072,223.48343968391418,202.0305461883545,0.0,1243.03125,0.0,0.9733513593673706,4402.953125,4080.7109375,0.046875,0.0,-0.08005471061915159,322.2421875,0.0,-4.616594552993774,77.42777967055838,2.995450356196656,1.601616382598877,24.39777112007141,-2150.12740173117,25.99579620361328,21.452893495559692,-1243.03125,-0.9733513593673706,322.2421875 +195,194,0.5546875,0.6328125,0.0,0.0,0.006384504958987236,0.10102255642414093,5770.7265625,5026.703125,7168.0,7168.0,77.09203386306763,121.59152936935425,1197.6804913981284,661.4543826954344,92.9797754814971,58.95147496850724,21.492231369018555,19.96917200088501,108.7973051071167,76.51310443878174,6944.115015129918,8629.449881075992,130.52270412445068,96.71688199043274,228.44471406936646,238.51559209823608,0.0,1700.21875,0.0,1.2050923109054565,5770.7265625,5026.703125,-0.078125,0.0,-0.0946380514651537,744.0234375,0.0,-44.49949550628662,536.226108702694,34.02830051298986,1.523059368133545,32.28420066833496,-1685.3348659460744,33.805822134017944,-10.070878028869629,-1700.21875,-1.2050923109054565,744.0234375 +196,195,0.6328125,0.6171875,0.0,0.0,0.006373968441039324,0.0860247015953064,5301.7734375,4922.8046875,7168.0,7168.0,88.96092009544373,121.05214977264404,953.5465113106965,650.6689484485279,80.57470620031414,59.21414872402258,19.765584230422974,18.57990074157715,100.15573048591614,71.07862114906311,6956.396769508589,9119.69013355775,120.15293955802917,89.89162278175354,229.14160585403442,231.02637767791748,0.0,1409.9375,0.0,1.070340633392334,5301.7734375,4922.8046875,0.015625,0.0,-0.07965073315426707,378.96875,0.0,-32.09122967720032,302.8775628621686,21.36055747629156,1.1856834888458252,29.077109336853027,-2163.29336404916,30.261316776275635,-1.8847718238830566,-1409.9375,-1.070340633392334,378.96875 +197,196,0.5703125,0.7265625,0.0,0.0,0.0062720924615859985,0.09767310321331024,5858.828125,5461.0703125,7168.0,7168.0,80.3999936580658,128.59631896018982,1165.9360372424085,679.4683215392017,89.15423588818778,55.740320235908406,20.461572408676147,20.98047113418579,106.74586939811707,80.28418755531311,7180.512035939448,8913.050275398147,127.44043827056885,101.49879288673401,228.04050636291504,250.18563795089722,0.0,1685.078125,0.0,1.347341775894165,5858.828125,5461.0703125,-0.15625,0.0,-0.09140101075172424,397.7578125,0.0,-48.19632530212402,486.46771570320675,33.41391565227938,-0.5188987255096436,26.461681842803955,-1732.5382394586995,25.94164538383484,-22.145131587982178,-1685.078125,-1.347341775894165,397.7578125 +198,197,0.5390625,0.515625,0.0,0.0,0.00653450284153223,0.10149289667606354,5699.171875,5458.5703125,7168.0,7168.0,85.4747097492218,131.73052263259888,1066.8272553078803,662.9983944084573,83.86106277553355,54.41411646101039,20.83823013305664,20.570802927017212,106.32076525688171,78.68702626228333,7061.104180224168,9149.475259113811,127.39283514022827,99.4967999458313,232.94275736808777,251.18184208869934,0.0,1854.5859375,0.0,1.2048883438110352,5699.171875,5458.5703125,0.0234375,0.0,-0.09495839383453131,240.6015625,0.0,-46.255812883377075,403.8288608994229,29.446946314523167,0.2674272060394287,27.63373899459839,-2088.3710788896433,27.896035194396973,-18.239084720611572,-1854.5859375,-1.2048883438110352,240.6015625 +199,198,0.6875,0.7578125,0.0,0.0,0.005708303768187761,0.08797280490398407,5785.21875,5399.4765625,7168.0,7168.0,84.42073655128479,125.14704012870789,1096.4545416370368,690.3209609364333,84.90804857696911,57.27662430232506,21.827643871307373,21.831876754760742,109.62449431419373,84.11970949172974,6918.709224109942,8429.475140659086,131.6857123374939,106.18678879737854,236.26299238204956,251.58984184265137,0.0,1618.5546875,0.0,1.3052937984466553,5785.21875,5399.4765625,-0.0703125,0.0,-0.08226450113579631,385.7421875,0.0,-40.726303577423096,406.1335807006035,27.631424274644047,-0.004232883453369141,25.50478482246399,-1510.765916549144,25.498923540115356,-15.326849460601807,-1618.5546875,-1.3052937984466553,385.7421875 +200,199,0.71875,0.6640625,0.0,0.0,0.005850364454090595,0.08289980888366699,5046.34375,4857.8359375,7168.0,7168.0,88.92592549324036,120.7709755897522,907.9635612691769,643.5766095326239,80.60641438637454,59.352008750422215,18.37946653366089,17.374534130096436,93.44398784637451,67.09266662597656,7111.639981509877,9545.171360830205,112.05668258666992,84.69973540306091,221.28256177902222,225.4642369747162,0.0,1475.9765625,0.0,1.1711537837982178,5046.34375,4857.8359375,0.0546875,0.0,-0.0770494444295764,188.5078125,0.0,-31.84505009651184,264.386951736553,21.254405635952324,1.0049324035644531,26.35132122039795,-2433.5313793203286,27.35694718360901,-4.18167519569397,-1475.9765625,-1.1711537837982178,188.5078125 diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/results/scale_reference_s1234_g10_g11_optimized/summary.json b/examples/vime_qwen3_8b_tp4_cp2_200/results/scale_reference_s1234_g10_g11_optimized/summary.json new file mode 100644 index 00000000..a28c88e5 --- /dev/null +++ b/examples/vime_qwen3_8b_tp4_cp2_200/results/scale_reference_s1234_g10_g11_optimized/summary.json @@ -0,0 +1,416 @@ +{ + "bootstrap": { + "reward_g10_minus_g11": { + "ci95": [ + 0.0306640625, + 0.0436328125 + ], + "draws": 20000, + "estimate": 0.037109375, + "method": "paired non-parametric bootstrap over the 200 aligned steps", + "seed": 1234 + }, + "step_time_g11_minus_g10_s": { + "ci95": [ + -23.51596219095588, + -17.82680384796858 + ], + "draws": 20000, + "estimate": -20.718443367481232, + "method": "paired non-parametric bootstrap over the 200 aligned steps", + "seed": 1234 + } + }, + "formulas": { + "longest_sample_tokens_s": "max response tokens / rollout seconds", + "rollout_tokens_per_gpu_s": "mean response tokens * 128 samples / (rollout seconds * 8 GPUs)", + "step_time_s": "VIME perf/step_time wall-clock timer" + }, + "missing_value_policy": "All required values must exist and be finite for all 200 paired steps; no imputation or row deletion.", + "paired": { + "active_tokens_mean": { + "g10": { + "max": 7103.40625, + "mean": 5492.253671875, + "median": 5547.88671875, + "min": 3559.25, + "n": 200, + "p90": 6379.3375, + "p95": 6585.995703124999, + "p99": 6974.569296875 + }, + "g11": { + "max": 7129.2421875, + "mean": 5757.0063671875, + "median": 5840.25390625, + "min": 3733.890625, + "n": 200, + "p90": 6470.08671875, + "p95": 6744.182031249999, + "p99": 7037.579374999999 + }, + "g11_minus_g10_mean": 264.7526953125, + "g11_over_g10": 1.048204746380936 + }, + "actor_train_time_s": { + "g10": { + "max": 100.66892504692078, + "mean": 80.50789266347886, + "median": 80.40358304977417, + "min": 51.76145315170288, + "n": 200, + "p90": 95.12249488830567, + "p95": 97.77730251550673, + "p99": 100.24143914461136 + }, + "g11": { + "max": 128.76038551330566, + "mean": 107.17858437657357, + "median": 108.1092381477356, + "min": 75.00430226325989, + "n": 200, + "p90": 120.00229210853577, + "p95": 122.81895134449002, + "p99": 127.20630695819855 + }, + "g11_minus_g10_mean": 26.670691713094712, + "g11_over_g10": 1.3312804599740002 + }, + "actor_train_tokens_s": { + "g10": { + "max": 9647.806536674922, + "mean": 8993.64551619016, + "median": 9057.04364631704, + "min": 8029.811782659763, + "n": 200, + "p90": 9328.264905670641, + "p95": 9386.301415922178, + "p99": 9554.316318122941 + }, + "g11": { + "max": 7321.05285928952, + "mean": 7050.923564993118, + "median": 7050.615711236377, + "min": 6669.510746786035, + "n": 200, + "p90": 7209.356735330589, + "p95": 7236.05419152076, + "p99": 7277.2575686346845 + }, + "g11_minus_g10_mean": -1942.721951197041, + "g11_over_g10": 0.7839894903907657 + }, + "kl_loss": { + "g10": { + "max": 0.11164486408233643, + "mean": 0.05696508722527142, + "median": 0.06753414869308472, + "min": -3.674428444355726e-05, + "n": 200, + "p90": 0.09705116376280784, + "p95": 0.10061634704470633, + "p99": 0.10367998927831648 + }, + "g11": { + "max": 0.007827508263289928, + "mean": 0.003782993672066368, + "median": 0.003752955002710223, + "min": 0.0, + "n": 200, + "p90": 0.006375022092834115, + "p95": 0.006963054556399583, + "p99": 0.0075022633932530875 + }, + "g11_minus_g10_mean": -0.05318209355320505, + "g11_over_g10": 0.06640898586016944 + }, + "longest_sample_tokens_s": { + "g10": { + "max": 71.14443572197055, + "mean": 55.51318646452719, + "median": 54.42095897200632, + "min": 44.760218224109536, + "n": 200, + "p90": 62.90640664061275, + "p95": 65.22441207296465, + "p99": 68.97004758366674 + }, + "g11": { + "max": 103.82007384579697, + "mean": 87.62363868027555, + "median": 88.46270237665703, + "min": 65.77580916697337, + "n": 200, + "p90": 98.39056173666839, + "p95": 99.7138393178839, + "p99": 102.65503319327365 + }, + "g11_minus_g10_mean": 32.11045221574834, + "g11_over_g10": 1.578429275290599 + }, + "max_abs_diff": { + "g10": { + "max": 1.591546893119812, + "mean": 0.9964957484602928, + "median": 0.978280782699585, + "min": 0.6022037267684937, + "n": 200, + "p90": 1.2850808382034302, + "p95": 1.3155867815017699, + "p99": 1.4539504694938654 + }, + "g11": { + "max": 0.0, + "mean": 0.0, + "median": 0.0, + "min": 0.0, + "n": 200, + "p90": 0.0, + "p95": 0.0, + "p99": 0.0 + }, + "g11_minus_g10_mean": -0.9964957484602928, + "g11_over_g10": 0.0 + }, + "mismatch_count": { + "g10": { + "max": 4758.0078125, + "mean": 2274.6178515625, + "median": 2040.76171875, + "min": 1121.1171875, + "n": 200, + "p90": 3661.65234375, + "p95": 3858.311328124998, + "p99": 4342.794218749999 + }, + "g11": { + "max": 0.0, + "mean": 0.0, + "median": 0.0, + "min": 0.0, + "n": 200, + "p90": 0.0, + "p95": 0.0, + "p99": 0.0 + }, + "g11_minus_g10_mean": -2274.6178515625, + "g11_over_g10": 0.0 + }, + "ref_log_probs_time_s": { + "g10": { + "max": 26.149888277053833, + "mean": 20.904491096735, + "median": 20.985463857650757, + "min": 13.43297791481018, + "n": 200, + "p90": 24.762764906883238, + "p95": 25.649108374118804, + "p99": 26.114913985729217 + }, + "g11": { + "max": 23.863692045211792, + "mean": 20.92132794857025, + "median": 21.15472960472107, + "min": 15.613138914108276, + "n": 200, + "p90": 23.291727018356323, + "p95": 23.698494446277618, + "p99": 23.755542702674866 + }, + "g11_minus_g10_mean": 0.016836851835250854, + "g11_over_g10": 1.0008054179246622 + }, + "reference_kl": { + "g10": { + "max": 0.0, + "mean": 0.0, + "median": 0.0, + "min": 0.0, + "n": 200, + "p90": 0.0, + "p95": 0.0, + "p99": 0.0 + }, + "g11": { + "max": 0.0, + "mean": 0.0, + "median": 0.0, + "min": 0.0, + "n": 200, + "p90": 0.0, + "p95": 0.0, + "p99": 0.0 + }, + "g11_minus_g10_mean": 0.0, + "g11_over_g10": null + }, + "response_len_max_tokens": { + "g10": { + "max": 7168.0, + "mean": 7168.0, + "median": 7168.0, + "min": 7168.0, + "n": 200, + "p90": 7168.0, + "p95": 7168.0, + "p99": 7168.0 + }, + "g11": { + "max": 7168.0, + "mean": 7168.0, + "median": 7168.0, + "min": 7168.0, + "n": 200, + "p90": 7168.0, + "p95": 7168.0, + "p99": 7168.0 + }, + "g11_minus_g10_mean": 0.0, + "g11_over_g10": 1.0 + }, + "response_len_mean_tokens": { + "g10": { + "max": 7103.40625, + "mean": 5492.253671875, + "median": 5547.88671875, + "min": 3559.25, + "n": 200, + "p90": 6379.3375, + "p95": 6585.995703124999, + "p99": 6974.569296875 + }, + "g11": { + "max": 7129.2421875, + "mean": 5757.0063671875, + "median": 5840.25390625, + "min": 3733.890625, + "n": 200, + "p90": 6470.08671875, + "p95": 6744.182031249999, + "p99": 7037.579374999999 + }, + "g11_minus_g10_mean": 264.7526953125, + "g11_over_g10": 1.048204746380936 + }, + "reward": { + "g10": { + "max": 0.921875, + "mean": 0.5285546875, + "median": 0.53125, + "min": 0.0546875, + "n": 200, + "p90": 0.75078125, + "p95": 0.7976562499999997, + "p99": 0.8985937499999999 + }, + "g11": { + "max": 0.8828125, + "mean": 0.4914453125, + "median": 0.484375, + "min": 0.0625, + "n": 200, + "p90": 0.71875, + "p95": 0.7898437499999997, + "p99": 0.8673437499999999 + }, + "g11_minus_g10_mean": -0.037109375, + "g11_over_g10": 0.9297908506392728 + }, + "rollout_time_s": { + "g10": { + "max": 160.1422040462494, + "mean": 130.21705961465835, + "median": 131.7139618396759, + "min": 100.75278449058533, + "n": 200, + "p90": 144.7520049095154, + "p95": 151.20598076581953, + "p99": 155.52637051105498 + }, + "g11": { + "max": 108.97623443603516, + "mean": 82.74613248229026, + "median": 81.02852284908295, + "min": 69.04252457618713, + "n": 200, + "p90": 94.74303011894226, + "p95": 99.63199486732482, + "p99": 106.57158301830292 + }, + "g11_minus_g10_mean": -47.47092713236809, + "g11_over_g10": 0.6354477111305902 + }, + "step_time_s": { + "g10": { + "max": 306.910710811615, + "mean": 251.98820484042167, + "median": 252.34174275398254, + "min": 186.46449184417725, + "n": 200, + "p90": 285.74714694023135, + "p95": 290.72313289642335, + "p99": 301.4975808525085 + }, + "g11": { + "max": 258.1519515514374, + "mean": 231.26976147294044, + "median": 231.45569777488708, + "min": 199.75356698036194, + "n": 200, + "p90": 242.52138652801514, + "p95": 247.6789080619812, + "p99": 253.92121420145034 + }, + "g11_minus_g10_mean": -20.718443367481232, + "g11_over_g10": 0.9177801064910885 + }, + "tokens_per_gpu_s": { + "g10": { + "max": 721.8912111291447, + "mean": 672.3886517288375, + "median": 677.732556671954, + "min": 565.2250733112137, + "n": 200, + "p90": 704.9581159153338, + "p95": 710.8700799072104, + "p99": 716.8867746089913 + }, + "g11": { + "max": 1410.439472415086, + "mean": 1133.9968589355974, + "median": 1213.9946873331705, + "min": 648.025465051748, + "n": 200, + "p90": 1355.7503861958087, + "p95": 1368.1704162827364, + "p99": 1389.6502373664766 + }, + "g11_minus_g10_mean": 461.60820720675997, + "g11_over_g10": 1.6865199256707835 + }, + "train_time_s": { + "g10": { + "max": 126.84098148345947, + "mean": 101.64788591861725, + "median": 101.66368591785431, + "min": 65.4232292175293, + "n": 200, + "p90": 120.45069971084594, + "p95": 123.75672025680542, + "p99": 126.40337255001067 + }, + "g11": { + "max": 152.71362829208374, + "mean": 128.33287456154824, + "median": 129.50647115707397, + "min": 90.84495854377747, + "n": 200, + "p90": 143.10054502487182, + "p95": 146.6778010845184, + "p99": 151.1407914042473 + }, + "g11_minus_g10_mean": 26.684988642930985, + "g11_over_g10": 1.262523793798288 + } + } +} diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/run.py b/examples/vime_qwen3_8b_tp4_cp2_200/run.py new file mode 100644 index 00000000..5bb0af6a --- /dev/null +++ b/examples/vime_qwen3_8b_tp4_cp2_200/run.py @@ -0,0 +1,315 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors + +"""Run and archive the Vime Qwen3-8B TP=4/CP=2 validation entry point. + +This is an integration example, not a synthetic pass generator. A dry run +only records the exact launch contract. ``--run`` executes Vime and records +whether the strict RL-Kernel provider was actually observed in the log. The +report deliberately leaves attention/FFN unclaimed until both framework +readbacks are supplied. +""" + +from __future__ import annotations + +import argparse +import json +import os +import subprocess +from datetime import datetime, timezone +from pathlib import Path +from typing import Any, Mapping + +DEFAULT_CONFIG = Path(__file__).with_name("qwen3_8b_tp4_cp2.json") +PROVIDER_MARKER = "linear_logp provider active" +FALLBACK_MARKERS = ("using native path", "fallback=True", "fallback=true") +RUNTIME_EVIDENCE_SCHEMA = "rlkernel.operator_runtime_evidence.v1" +_OPERATOR_METRICS = { + "attention": ( + "out_max_abs", + "lse_max_abs", + "dq_max_abs", + "dk_max_abs", + "dv_max_abs", + ), + "ffn": ("out_max_abs", "dx_max_abs", "dw_max_abs"), +} + + +def load_config(path: Path) -> dict[str, Any]: + with path.open(encoding="utf-8") as handle: + value = json.load(handle) + if not isinstance(value, dict): + raise ValueError("example config must contain a JSON object") + return value + + +def validate_config(config: Mapping[str, Any]) -> None: + training = config.get("training") + rollout = config.get("rollout") + provider = config.get("linear_logp_provider") + if ( + not isinstance(training, Mapping) + or not isinstance(rollout, Mapping) + or not isinstance(provider, Mapping) + ): + raise ValueError( + "training, rollout, and linear_logp_provider sections are required" + ) + expected = { + "tensor_model_parallel_size": 4, + "context_parallel_size": 2, + "pipeline_model_parallel_size": 1, + "world_size": 8, + } + for name, value in expected.items(): + if training.get(name) != value: + raise ValueError(f"training.{name} must be {value!r}") + if rollout.get("top_p") != 1.0: + raise ValueError( + "rollout.top_p must remain 1.0 for the strict provider contract" + ) + if provider.get("mode") != "strict": + raise ValueError("linear_logp_provider.mode must be strict") + if ( + provider.get("path") + != "rl_engine.integrations.vime.linear_logp_provider.provider" + ): + raise ValueError("example must use the RL-Kernel Vime provider") + if provider.get("backend_id") != "rlkernel.linear_logp.bitwise.v1": + raise ValueError("example must pin the deterministic vocab-parallel backend") + + +def load_runtime_evidence(path: Path | None) -> dict[str, Any] | None: + """Load post-execution readback without treating configuration as evidence.""" + + if path is None: + return None + with path.open(encoding="utf-8") as handle: + value = json.load(handle) + if ( + not isinstance(value, dict) + or value.get("schema_version") != RUNTIME_EVIDENCE_SCHEMA + ): + raise ValueError( + f"runtime evidence must use schema {RUNTIME_EVIDENCE_SCHEMA!r}" + ) + return value + + +def _operator_evidence_status(evidence: Mapping[str, Any] | None, operator: str) -> str: + if evidence is None: + return "unclaimed" + operators = evidence.get("operators") + item = operators.get(operator) if isinstance(operators, Mapping) else None + if not isinstance(item, Mapping): + return "unclaimed" + training = item.get("training") + rollout = item.get("rollout") + comparison = item.get("comparison") + if not isinstance(training, Mapping) or not isinstance(rollout, Mapping): + return "unclaimed" + if not isinstance(comparison, Mapping) or comparison.get("passed") is not True: + return "failed" + required_identity = ("implementation_id", "backend_id", "contract_id") + if any( + not training.get(name) or not rollout.get(name) for name in required_identity + ): + return "failed" + if training["implementation_id"] != rollout["implementation_id"]: + return "failed" + for metric in _OPERATOR_METRICS[operator]: + value = comparison.get(metric) + if ( + not isinstance(value, (int, float)) + or isinstance(value, bool) + or value != 0.0 + ): + return "failed" + return "passed" + + +def validate_runtime_evidence(evidence: Mapping[str, Any] | None) -> None: + """Reject malformed evidence before it can affect a report.""" + + if evidence is None: + return + for operator in _OPERATOR_METRICS: + status = _operator_evidence_status(evidence, operator) + if status == "failed": + raise ValueError( + f"runtime evidence for {operator} is incomplete or non-zero" + ) + + +def _revision(path: Path) -> str | None: + try: + result = subprocess.run( + ["git", "-C", str(path), "rev-parse", "HEAD"], + check=True, + capture_output=True, + text=True, + ) + except (OSError, subprocess.CalledProcessError): + return None + return result.stdout.strip() or None + + +def build_environment(vime_root: Path, rl_kernel_root: Path) -> dict[str, str]: + env = dict(os.environ) + existing = [str(vime_root), str(rl_kernel_root), "/root/Megatron-LM"] + if env.get("PYTHONPATH"): + existing.append(env["PYTHONPATH"]) + env["PYTHONPATH"] = os.pathsep.join(existing) + env["RL_KERNEL_ROOT"] = str(rl_kernel_root) + env["TP_SIZE"] = "4" + env["CP_SIZE"] = "2" + env["ROLLOUT_TOP_P"] = "1.0" + return env + + +def build_command(config: Mapping[str, Any], vime_root: Path) -> list[str]: + script = vime_root / str(config.get("vime_script", "")) + if not script.is_file(): + raise FileNotFoundError(f"Vime entry script does not exist: {script}") + return ["bash", str(script)] + + +def build_report( + config: Mapping[str, Any], + *, + vime_root: Path, + rl_kernel_root: Path, + command: list[str], + status: str, + returncode: int | None, + log_text: str, + log_path: Path | None, + runtime_evidence: Mapping[str, Any] | None = None, + runtime_evidence_path: Path | None = None, +) -> dict[str, Any]: + provider_active = PROVIDER_MARKER in log_text + fallback_observed = any(marker in log_text for marker in FALLBACK_MARKERS) + strict_provider_passed = ( + status == "passed" and provider_active and not fallback_observed + ) + effective_status = ( + "passed" + if strict_provider_passed + else ("failed" if status == "passed" else status) + ) + attention_status = _operator_evidence_status(runtime_evidence, "attention") + ffn_status = _operator_evidence_status(runtime_evidence, "ffn") + return { + "schema_version": "rlkernel.vime_validation_report.v1", + "created_at": datetime.now(timezone.utc).isoformat(), + "status": effective_status, + "claim_boundary": { + "qwen3_8b_tp4_cp2_vime_training": strict_provider_passed, + "attention_train_infer_consistency": attention_status, + "ffn_train_infer_consistency": ffn_status, + "reason": ( + "attention and FFN require executed Megatron/vLLM runtime readbacks; " + "the evidence contract accepts only exact-zero comparison metrics" + ), + }, + "config": dict(config), + "topology": config["training"], + "provider": { + "configured_path": config["linear_logp_provider"]["path"], + "configured_mode": config["linear_logp_provider"]["mode"], + "backend_id": config["linear_logp_provider"]["backend_id"], + "active_observed": provider_active, + "fallback_observed": fallback_observed, + }, + "command": command, + "returncode": returncode, + "artifacts": { + "log": None if log_path is None else str(log_path), + "runtime_evidence": ( + None if runtime_evidence_path is None else str(runtime_evidence_path) + ), + }, + "runtime_evidence": ( + None if runtime_evidence is None else dict(runtime_evidence) + ), + "revisions": { + "vime": _revision(vime_root), + "rl_kernel": _revision(rl_kernel_root), + }, + } + + +def main(argv: list[str] | None = None) -> int: + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--config", type=Path, default=DEFAULT_CONFIG) + parser.add_argument( + "--vime-root", type=Path, default=Path(os.environ.get("VIME_ROOT", ".")) + ) + parser.add_argument( + "--rl-kernel-root", + type=Path, + default=Path(os.environ.get("RL_KERNEL_ROOT", ".")), + ) + parser.add_argument( + "--output", type=Path, default=Path("qwen3_8b_tp4_cp2.validation.json") + ) + parser.add_argument( + "--runtime-evidence", + type=Path, + default=None, + help="post-execution Megatron/vLLM operator readback JSON (strict exact-zero contract)", + ) + parser.add_argument("--run", action="store_true", help="execute the Vime script") + args = parser.parse_args(argv) + + config = load_config(args.config) + validate_config(config) + runtime_evidence = load_runtime_evidence(args.runtime_evidence) + validate_runtime_evidence(runtime_evidence) + vime_root = args.vime_root.resolve() + rl_kernel_root = args.rl_kernel_root.resolve() + command = build_command(config, vime_root) + + status = "not_run" + returncode: int | None = None + log_text = "" + log_path: Path | None = None + if args.run: + args.output.parent.mkdir(parents=True, exist_ok=True) + log_path = args.output.with_suffix(".log") + env = build_environment(vime_root, rl_kernel_root) + with log_path.open("w", encoding="utf-8") as log_handle: + process = subprocess.run( + command, + cwd=vime_root, + env=env, + stdout=log_handle, + stderr=subprocess.STDOUT, + ) + returncode = process.returncode + log_text = log_path.read_text(encoding="utf-8", errors="replace") + status = "passed" if returncode == 0 else "failed" + + report = build_report( + config, + vime_root=vime_root, + rl_kernel_root=rl_kernel_root, + command=command, + status=status, + returncode=returncode, + log_text=log_text, + log_path=log_path, + runtime_evidence=runtime_evidence, + runtime_evidence_path=args.runtime_evidence, + ) + args.output.parent.mkdir(parents=True, exist_ok=True) + args.output.write_text( + json.dumps(report, indent=2, sort_keys=True) + "\n", encoding="utf-8" + ) + print(json.dumps(report, indent=2, sort_keys=True)) + return 0 if report["status"] in {"passed", "not_run"} else 1 + + +if __name__ == "__main__": + raise SystemExit(main()) diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/run_arm.py b/examples/vime_qwen3_8b_tp4_cp2_200/run_arm.py new file mode 100755 index 00000000..cb07b8fb --- /dev/null +++ b/examples/vime_qwen3_8b_tp4_cp2_200/run_arm.py @@ -0,0 +1,595 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors + +"""Submit one append-only VIME × RL-Kernel experiment arm to Ray.""" + +from __future__ import annotations + +import argparse +import hashlib +import json +import os +import re +import subprocess +import sys +from dataclasses import asdict, dataclass +from datetime import datetime, timezone +from pathlib import Path +from typing import Any + + +@dataclass(frozen=True) +class Arm: + group: str + framework_use_rollout_logprobs: bool + attention_case: str + ffn_case: str + logp_case: str + description: str + + +ARMS = { + "G00": Arm("G00", False, "P/P", "P/P", "P/P", "native VIME baseline"), + "G10": Arm("G10", True, "P/P", "P/P", "P/P", "VIME framework-level consistency only"), + "G01": Arm("G01", False, "R/R", "R/R", "R/R", "RL-Kernel operator-level consistency only"), + "G11": Arm( + "G11", + True, + "R/R", + "R/R", + "R/R", + "framework-level and operator-level consistency", + ), + # Short operator-attribution matrix. Framework logp reuse stays disabled so + # that only Attention, FFN, and selected-token logp change across M000-M111. + "M000": Arm("M000", False, "P/P", "P/P", "P/P", "all production"), + "M100": Arm("M100", False, "R/R", "P/P", "P/P", "RL-Kernel attention"), + "M010": Arm("M010", False, "P/P", "R/R", "P/P", "RL-Kernel FFN"), + "M001": Arm("M001", False, "P/P", "P/P", "R/R", "RL-Kernel logp"), + "M110": Arm("M110", False, "R/R", "R/R", "P/P", "RL-Kernel attention and FFN"), + "M101": Arm("M101", False, "R/R", "P/P", "R/R", "RL-Kernel attention and logp"), + "M011": Arm("M011", False, "P/P", "R/R", "R/R", "RL-Kernel FFN and logp"), + "M111": Arm("M111", False, "R/R", "R/R", "R/R", "all RL-Kernel operators"), +} + +TOPOLOGY = { + "gpus": 8, + "actor_gpus": 8, + "rollout_gpus": 8, + "tp": 4, + "cp": 2, + "pp": 1, + "colocate": True, + "offload_train": False, + "offload_rollout": True, + "rollout_gpus_per_engine": 4, + "rollout_engines": 2, +} + +# CP2/P2P requires Transformer Engine's fused attention on this topology. +# Pinning the choice keeps the production arms independent of host-specific +# backend auto-selection. +MEGATRON_ATTENTION_BACKEND = "fused" +RL_KERNEL_LINEAR_LOGP_PROVIDER = "rl_engine.integrations.vime.linear_logp_provider.provider" + +MODEL_ARGS = ( + "--swiglu", + "--num-layers", + "36", + "--hidden-size", + "4096", + "--ffn-hidden-size", + "12288", + "--num-attention-heads", + "32", + "--group-query-attention", + "--num-query-groups", + "8", + "--use-rotary-position-embeddings", + "--disable-bias-linear", + "--normalization", + "RMSNorm", + "--norm-epsilon", + "1e-6", + "--rotary-base", + "1000000", + "--vocab-size", + "151936", + "--kv-channels", + "128", + "--qk-layernorm", + "--untie-embeddings-and-output-weights", +) + + +def _max_engine_decode_batch( + rollout_batch_size: int, n_samples_per_prompt: int, router_policy: str +) -> int: + """Largest decode batch one vLLM engine can hold under the active router. + + round_robin is the only supported policy that guarantees an even split for + this finite request burst. cache_aware can pin a shared prefix and random + can produce an uneven split, so both must graph the full concurrency. + """ + + concurrency = rollout_batch_size * n_samples_per_prompt + engines = TOPOLOGY["rollout_gpus"] // TOPOLOGY["rollout_gpus_per_engine"] + if router_policy == "round_robin" and engines > 1: + return -(-concurrency // engines) # ceil + return concurrency + + +def _linear_logp_provider_args(arm: Arm) -> tuple[str, ...]: + """Install the RL-Kernel provider only on a training-side R route.""" + + training_side, _rollout_side = arm.logp_case.split("/", 1) + if training_side == "P": + return () + if training_side == "R": + return ( + "--linear-logp-provider", + RL_KERNEL_LINEAR_LOGP_PROVIDER, + "--linear-logp-provider-mode", + "strict", + ) + raise ValueError(f"unsupported training logp route: {arm.logp_case!r}") + + +def _path(value: str | None, label: str) -> Path: + if not value: + raise ValueError(f"{label} is required (argument or environment variable)") + path = Path(value).expanduser().resolve() + if not path.exists(): + raise FileNotFoundError(f"{label} does not exist: {path}") + return path + + +def _sha256(path: Path) -> str: + digest = hashlib.sha256() + with path.open("rb") as handle: + for chunk in iter(lambda: handle.read(1024 * 1024), b""): + digest.update(chunk) + return digest.hexdigest() + + +def _revision(path: Path) -> str: + return subprocess.run( + ["git", "-C", str(path), "rev-parse", "HEAD"], + check=True, + capture_output=True, + text=True, + ).stdout.strip() + + +def _repository_state(path: Path) -> dict[str, Any]: + status = subprocess.run( + ["git", "-C", str(path), "status", "--porcelain=v1", "--untracked-files=all"], + check=True, + capture_output=True, + text=True, + ).stdout.splitlines() + tracked_diff = subprocess.run( + ["git", "-C", str(path), "diff", "--binary", "HEAD"], + check=True, + capture_output=True, + ).stdout + return { + "revision": _revision(path), + "dirty": bool(status), + "status": status, + "tracked_diff_sha256": hashlib.sha256(tracked_diff).hexdigest(), + } + + +def _gpu_inventory() -> list[dict[str, str]]: + result = subprocess.run( + [ + "nvidia-smi", + "--query-gpu=index,name,memory.total,driver_version", + "--format=csv,noheader,nounits", + ], + check=True, + capture_output=True, + text=True, + ) + inventory = [] + for line in result.stdout.splitlines(): + index, name, memory, driver = (field.strip() for field in line.split(",", 3)) + inventory.append({"index": index, "name": name, "memory_mib": memory, "driver": driver}) + return inventory + + +def _default_run_id(group: str, num_rollout: int, seed: int) -> str: + stamp = datetime.now(timezone.utc).strftime("%Y%m%dT%H%M%SZ") + return f"{group.lower()}-n{num_rollout}-s{seed}-{stamp}" + + +def _write_json(path: Path, value: Any) -> None: + path.write_text(json.dumps(value, indent=2, sort_keys=True) + "\n", encoding="utf-8") + + +def build_parser() -> argparse.ArgumentParser: + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--group", choices=tuple(ARMS), required=True) + parser.add_argument("--num-rollout", type=int, required=True) + parser.add_argument("--seed", type=int, default=1234) + parser.add_argument("--rollout-seed", type=int, default=42) + parser.add_argument("--run-id", default=None) + parser.add_argument("--output-root", type=Path, required=True) + parser.add_argument("--rl-kernel-root", default=os.environ.get("RL_KERNEL_ROOT")) + parser.add_argument("--vime-root", default=os.environ.get("VIME_ROOT")) + parser.add_argument("--megatron-root", default=os.environ.get("MEGATRON_ROOT")) + parser.add_argument("--model-root", default=os.environ.get("MODEL_ROOT")) + parser.add_argument("--ref-load", default=os.environ.get("TORCH_DIST_ROOT")) + parser.add_argument("--prompt-data", default=os.environ.get("PROMPT_DATA")) + parser.add_argument("--python", default=os.environ.get("RL_KERNEL_REAL_PYTHON", sys.executable)) + parser.add_argument("--ray-bin", default=os.environ.get("RAY_BIN")) + parser.add_argument( + "--ray-address", + default=os.environ.get("RAY_API_SERVER_ADDRESS", "http://127.0.0.1:8265"), + ) + parser.add_argument("--rollout-batch-size", type=int, default=1) + parser.add_argument("--n-samples-per-prompt", type=int, default=8) + parser.add_argument("--global-batch-size", type=int, default=8) + parser.add_argument( + "--use-kl-loss", + action="store_true", + help="Load the reference checkpoint and add a KL term to the policy loss.", + ) + parser.add_argument( + "--kl-loss-coef", + type=float, + default=0.0, + help="Reference-model KL loss coefficient; requires --use-kl-loss.", + ) + parser.add_argument("--max-response-len", type=int, default=7168) + parser.add_argument("--max-tokens-per-gpu", type=int, default=4096) + parser.add_argument("--vllm-gpu-memory-utilization", type=float, default=0.4) + parser.add_argument( + "--router-policy", + choices=("round_robin", "random", "cache_aware"), + default="round_robin", + help=( + "vLLM router policy. cache_aware pins every sample of one prompt to " + "a single engine, which idles the other engine when the concurrent " + "request count stays below router_balance_abs_threshold." + ), + ) + parser.add_argument("--extra-pythonpath", action="append", default=[]) + parser.add_argument("--ld-library-path", default=os.environ.get("LD_LIBRARY_PATH", "")) + parser.add_argument("--wait", action="store_true", help="stream logs until the Ray job exits") + parser.add_argument( + "--allow-dirty", + action="store_true", + help="permit development runs from dirty repository states", + ) + parser.add_argument( + "--dry-run", action="store_true", help="write the manifest without submitting" + ) + return parser + + +def main(argv: list[str] | None = None) -> int: + args = build_parser().parse_args(argv) + if args.num_rollout <= 0: + raise ValueError("--num-rollout must be positive") + trajectories_per_rollout = args.rollout_batch_size * args.n_samples_per_prompt + if args.global_batch_size != trajectories_per_rollout: + raise ValueError( + "--global-batch-size must equal --rollout-batch-size * " + "--n-samples-per-prompt for an exact per-step comparison " + f"({args.global_batch_size} != {trajectories_per_rollout})" + ) + if args.use_kl_loss != (args.kl_loss_coef > 0.0): + raise ValueError( + "--use-kl-loss requires a positive --kl-loss-coef, and a positive " + "coefficient requires --use-kl-loss" + ) + arm = ARMS[args.group] + + script_dir = Path(__file__).resolve().parent + rl_kernel_root = _path(args.rl_kernel_root, "RL-Kernel root") + vime_root = _path(args.vime_root, "VIME root") + megatron_root = _path(args.megatron_root, "Megatron root") + model_root = _path(args.model_root, "HF model root") + ref_load = _path(args.ref_load, "Megatron torch-dist checkpoint") + prompt_data = _path(args.prompt_data, "prompt data") + python = _path(args.python, "Python executable") + ray_bin = _path(args.ray_bin or str(python.parent / "ray"), "Ray executable") + entrypoint = _path( + str(script_dir / "aligned_python_entrypoint.sh"), "aligned Python entrypoint" + ) + repository_state = { + "rl_kernel": _repository_state(rl_kernel_root), + "vime": _repository_state(vime_root), + "megatron": _repository_state(megatron_root), + } + dirty_repositories = [name for name, state in repository_state.items() if state["dirty"]] + if dirty_repositories and not args.allow_dirty: + raise RuntimeError( + "refusing a non-reproducible run from dirty repositories: " + + ", ".join(dirty_repositories) + + "; commit the changes or pass --allow-dirty for a development run" + ) + + run_id = args.run_id or _default_run_id(args.group, args.num_rollout, args.seed) + if not re.fullmatch(r"[A-Za-z0-9][A-Za-z0-9._-]*", run_id): + raise ValueError("--run-id may contain only letters, digits, dot, underscore, and dash") + run_dir = args.output_root.expanduser().resolve() / run_id + run_dir.mkdir(parents=True, exist_ok=False) + + pythonpath = [str(rl_kernel_root), str(vime_root), str(megatron_root)] + pythonpath.extend(str(Path(item).expanduser().resolve()) for item in args.extra_pythonpath) + if os.environ.get("PYTHONPATH"): + pythonpath.extend(item for item in os.environ["PYTHONPATH"].split(os.pathsep) if item) + + max_engine_decode_batch = _max_engine_decode_batch( + args.rollout_batch_size, args.n_samples_per_prompt, args.router_policy + ) + env_vars = { + "RL_KERNEL_ROOT": str(rl_kernel_root), + "RL_KERNEL_REAL_PYTHON": str(python), + "PYTHONPATH": os.pathsep.join(dict.fromkeys(pythonpath)), + "LD_LIBRARY_PATH": args.ld_library_path, + "PYTHONUNBUFFERED": "1", + "PYTHONHASHSEED": str(args.seed), + "CUDA_DEVICE_MAX_CONNECTIONS": "1", + "NCCL_ALGO": "Ring", + "NVTE_ALLOW_NONDETERMINISTIC_ALGO": "0", + "CUBLAS_WORKSPACE_CONFIG": ":16:8", + "CUBLASLT_WORKSPACE_SIZE": "1", + "VLLM_BATCH_INVARIANT": "1", + "RL_KERNEL_VLLM_INTEGRATION": "1", + "RL_KERNEL_CUDA_ONLY": "1", + "VIME_RL_KERNEL_STRICT": "1", + "RL_KERNEL_ATTENTION_CASE": arm.attention_case, + "RL_KERNEL_FFN_CASE": arm.ffn_case, + "RL_KERNEL_LOGP_CASE": arm.logp_case, + "RL_KERNEL_READBACK_DIR": str(run_dir / "readbacks"), + "RL_KERNEL_VLLM_REAL_VOCAB_SIZE": "151936", + "RL_KERNEL_VLLM_PADDED_VOCAB_SIZE": "152064", + "RL_KERNEL_VLLM_TEMPERATURE": "1.0", + "RL_KERNEL_VLLM_CUDAGRAPH_MAX_CAPTURE_SIZE": str(max_engine_decode_batch), + "RL_KERNEL_SEED": str(args.seed), + "RL_KERNEL_ROLLOUT_SEED": str(args.rollout_seed), + "RL_KERNEL_RUN_ID": run_id, + } + if os.environ.get("CUDNN_FRONTEND_CUDART_LIB_NAME"): + env_vars["CUDNN_FRONTEND_CUDART_LIB_NAME"] = os.environ["CUDNN_FRONTEND_CUDART_LIB_NAME"] + + train_command = [ + str(entrypoint), + "train.py", + "--train-backend", + "megatron", + "--actor-num-nodes", + "1", + "--actor-num-gpus-per-node", + str(TOPOLOGY["actor_gpus"]), + "--rollout-num-gpus", + str(TOPOLOGY["rollout_gpus"]), + "--colocate", + "--no-offload-train", + "--offload-rollout", + *MODEL_ARGS, + "--hf-checkpoint", + str(model_root), + "--ref-load", + str(ref_load), + "--load", + str(run_dir / "initial-load"), + "--prompt-data", + str(prompt_data), + "--input-key", + "prompt", + "--label-key", + "label", + "--apply-chat-template", + "--rollout-shuffle", + "--rm-type", + "deepscaler", + "--advantage-estimator", + "grpo", + "--num-rollout", + str(args.num_rollout), + "--rollout-batch-size", + str(args.rollout_batch_size), + "--n-samples-per-prompt", + str(args.n_samples_per_prompt), + "--rollout-max-response-len", + str(args.max_response_len), + "--rollout-temperature", + "1.0", + "--rollout-top-p", + "1.0", + "--global-batch-size", + str(args.global_batch_size), + "--balance-data", + "--tensor-model-parallel-size", + str(TOPOLOGY["tp"]), + "--context-parallel-size", + str(TOPOLOGY["cp"]), + "--cp-comm-type", + "p2p", + "--pipeline-model-parallel-size", + "1", + "--expert-model-parallel-size", + "1", + "--expert-tensor-parallel-size", + "1", + "--use-dynamic-batch-size", + "--max-tokens-per-gpu", + str(args.max_tokens_per_gpu), + "--recompute-granularity", + "full", + "--recompute-method", + "uniform", + "--recompute-num-layers", + "1", + "--custom-megatron-init-path", + "rl_engine.integrations.megatron_runtime.initialize_from_environment", + "--save-debug-train-data", + str(run_dir / "train-data" / "{rollout_id}.rank{rank}.pt"), + "--update-weight-mode", + "full", + "--update-weight-transport", + "disk", + "--update-weight-disk-dir", + str(run_dir / "weight-updates"), + *_linear_logp_provider_args(arm), + "--no-save-optim", + "--attention-dropout", + "0.0", + "--hidden-dropout", + "0.0", + "--transformer-impl", + "transformer_engine", + "--no-persist-layer-norm", + "--no-gradient-accumulation-fusion", + "--no-rope-fusion", + "--attention-softmax-in-fp32", + "--attention-backend", + MEGATRON_ATTENTION_BACKEND, + "--router-policy", + args.router_policy, + "--rollout-num-gpus-per-engine", + str(TOPOLOGY["rollout_gpus_per_engine"]), + "--vllm-gpu-memory-utilization", + str(args.vllm_gpu_memory_utilization), + ] + if arm.framework_use_rollout_logprobs: + train_command.append("--use-rollout-logprobs") + if args.use_kl_loss: + train_command.extend(["--use-kl-loss", "--kl-loss-coef", str(args.kl_loss_coef)]) + if {arm.attention_case, arm.ffn_case, arm.logp_case} == {"R/R"}: + train_command.extend( + [ + "--ci-test", + "--ci-disable-kl-checker", + "--ci-train-rollout-logprob-abs-diff-threshold", + "0", + ] + ) + + submission_id = f"vime200-{run_id}" + runtime_env = {"env_vars": env_vars} + ray_command = [ + str(ray_bin), + "job", + "submit", + f"--address={args.ray_address}", + "--submission-id", + submission_id, + "--runtime-env-json", + json.dumps(runtime_env, separators=(",", ":")), + "--metadata-json", + json.dumps({"group": args.group, "run_id": run_id}), + "--working-dir", + str(vime_root), + ] + if not args.wait: + ray_command.append("--no-wait") + ray_command.extend(["--", *train_command]) + + manifest = { + "schema_version": "rlkernel.vime_qwen3_8b_tp4_cp2_200.run.v1", + "created_at": datetime.now(timezone.utc).isoformat(), + "status": "planned" if args.dry_run else "submitting", + "run_id": run_id, + "submission_id": submission_id, + "arm": asdict(arm), + "num_rollout": args.num_rollout, + "seed": args.seed, + "rollout_seed": args.rollout_seed, + "topology": dict(TOPOLOGY), + "batching": { + "rollout_batch_size": args.rollout_batch_size, + "n_samples_per_prompt": args.n_samples_per_prompt, + "global_batch_size": args.global_batch_size, + "max_response_len": args.max_response_len, + "max_tokens_per_gpu": args.max_tokens_per_gpu, + }, + "algorithm": { + "advantage_estimator": "grpo", + "reward_model": "deepscaler", + "reference_model": { + "enabled": args.use_kl_loss, + "mode": "kl_loss" if args.use_kl_loss else None, + "coefficient": args.kl_loss_coef, + }, + }, + "snapshotting": { + "train_data_every_step": True, + "template": str(run_dir / "train-data" / "{rollout_id}.rank{rank}.pt"), + "full_model_checkpoint_every_step": False, + }, + "rollout_routing": { + "router_policy": args.router_policy, + }, + "training_memory": { + "recompute_granularity": "full", + "recompute_method": "uniform", + "recompute_num_layers": 1, + }, + "vllm_execution": { + "cudagraph_required": True, + "cudagraph_mode": "FULL_DECODE_ONLY", + "capture_sizes": list(range(1, max_engine_decode_batch + 1)), + "enforce_eager": False, + }, + "paths": { + "run_dir": str(run_dir), + "rl_kernel_root": str(rl_kernel_root), + "vime_root": str(vime_root), + "megatron_root": str(megatron_root), + "model_root": str(model_root), + "ref_load": str(ref_load), + "prompt_data": str(prompt_data), + }, + "revisions": {name: state["revision"] for name, state in repository_state.items()}, + "repository_state": repository_state, + "prompt_data_sha256": _sha256(prompt_data), + "gpu_inventory": _gpu_inventory(), + "runtime_env": runtime_env, + "train_command": train_command, + "ray_command": ray_command, + } + manifest_path = run_dir / "manifest.json" + _write_json(manifest_path, manifest) + + if args.dry_run: + print( + json.dumps( + { + "status": "planned", + "run_dir": str(run_dir), + "manifest": str(manifest_path), + } + ) + ) + return 0 + + result = subprocess.run(ray_command, capture_output=True, text=True) + manifest["submission"] = { + "returncode": result.returncode, + "stdout": result.stdout, + "stderr": result.stderr, + } + manifest["status"] = "submitted" if result.returncode == 0 else "submission_failed" + _write_json(manifest_path, manifest) + print(result.stdout, end="") + print(result.stderr, end="", file=sys.stderr) + print( + json.dumps( + { + "status": manifest["status"], + "submission_id": submission_id, + "run_dir": str(run_dir), + } + ) + ) + return result.returncode + + +if __name__ == "__main__": + raise SystemExit(main()) diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/run_supplement_suite.py b/examples/vime_qwen3_8b_tp4_cp2_200/run_supplement_suite.py new file mode 100644 index 00000000..bc18cbbe --- /dev/null +++ b/examples/vime_qwen3_8b_tp4_cp2_200/run_supplement_suite.py @@ -0,0 +1,220 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors + +"""Run the TP4/CP2 supplementary VIME experiment suites serially. + +The controller is deliberately fail-closed: it waits for all eight GPUs to be +idle, refuses existing run IDs, stops on the first failed Ray job, captures the +authoritative Ray log, and seals only runs accepted by ``validate_run.py``. +""" + +from __future__ import annotations + +import argparse +import json +import subprocess +import time +from pathlib import Path + +MODULE_GROUPS = ("M000", "M100", "M010", "M001", "M110", "M101", "M011", "M111") +PRECISION_GROUPS = ("G00", "G10", "G01", "G11") +PRECISION_SEEDS = (1234, 2345, 3456) + + +def parse_args() -> argparse.Namespace: + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--phase", choices=("module", "precision"), required=True) + parser.add_argument("--suite-id", required=True) + parser.add_argument("--output-root", type=Path, required=True) + parser.add_argument("--rl-kernel-root", type=Path, required=True) + parser.add_argument("--vime-root", type=Path, required=True) + parser.add_argument("--megatron-root", type=Path, required=True) + parser.add_argument("--model-root", type=Path, required=True) + parser.add_argument("--ref-load", type=Path, required=True) + parser.add_argument("--prompt-data", type=Path, required=True) + parser.add_argument("--python", type=Path, required=True) + parser.add_argument("--ray-bin", type=Path, required=True) + parser.add_argument("--extra-pythonpath", action="append", default=[]) + parser.add_argument("--ld-library-path", required=True) + parser.add_argument("--idle-memory-mib", type=int, default=1024) + parser.add_argument("--idle-poll-seconds", type=int, default=60) + return parser.parse_args() + + +def run_checked(command: list[str], *, stdout=None) -> None: + result = subprocess.run( + command, + check=False, + text=True, + stdout=stdout, + stderr=subprocess.STDOUT if stdout is not None else None, + ) + if result.returncode: + raise RuntimeError(f"command failed with return code {result.returncode}: {command[0]}") + + +def gpu_memory() -> list[int]: + output = subprocess.run( + [ + "nvidia-smi", + "--query-gpu=memory.used", + "--format=csv,noheader,nounits", + ], + check=True, + capture_output=True, + text=True, + ).stdout + return [int(line.strip()) for line in output.splitlines() if line.strip()] + + +def wait_for_idle_gpus(*, threshold: int, poll_seconds: int) -> None: + while True: + memory = gpu_memory() + if len(memory) != 8: + raise RuntimeError(f"expected exactly 8 GPUs, found {len(memory)}") + if all(value <= threshold for value in memory): + return + print( + json.dumps( + { + "event": "waiting_for_idle_gpus", + "memory_used_mib": memory, + "threshold_mib": threshold, + }, + sort_keys=True, + ), + flush=True, + ) + time.sleep(poll_seconds) + + +def specs(phase: str) -> list[tuple[str, int, int]]: + if phase == "module": + return [(group, 8, 1234) for group in MODULE_GROUPS] + return [(group, 8, seed) for seed in PRECISION_SEEDS for group in PRECISION_GROUPS] + + +def run_one(args: argparse.Namespace, group: str, rounds: int, seed: int) -> None: + run_id = f"{args.suite_id}-{group.lower()}-n{rounds}-b8-s16-" f"refkl001-s{seed}" + run_dir = args.output_root / run_id + submission_id = f"vime200-{run_id}" + if run_dir.exists(): + raise FileExistsError(f"refusing existing run directory: {run_dir}") + status = subprocess.run( + [str(args.ray_bin), "job", "status", submission_id], + check=False, + capture_output=True, + text=True, + ) + if status.returncode == 0: + raise RuntimeError(f"refusing existing Ray submission: {submission_id}") + + wait_for_idle_gpus( + threshold=args.idle_memory_mib, + poll_seconds=args.idle_poll_seconds, + ) + command = [ + str(args.python), + str(args.rl_kernel_root / "examples/vime_qwen3_8b_tp4_cp2_200/run_arm.py"), + "--group", + group, + "--run-id", + run_id, + "--num-rollout", + str(rounds), + "--seed", + str(seed), + "--rollout-seed", + str(seed), + "--output-root", + str(args.output_root), + "--rl-kernel-root", + str(args.rl_kernel_root), + "--vime-root", + str(args.vime_root), + "--megatron-root", + str(args.megatron_root), + "--model-root", + str(args.model_root), + "--ref-load", + str(args.ref_load), + "--prompt-data", + str(args.prompt_data), + "--python", + str(args.python), + "--ray-bin", + str(args.ray_bin), + "--ld-library-path", + args.ld_library_path, + "--rollout-batch-size", + "8", + "--n-samples-per-prompt", + "16", + "--global-batch-size", + "128", + "--max-response-len", + "7168", + "--max-tokens-per-gpu", + "4096", + "--vllm-gpu-memory-utilization", + "0.4", + "--router-policy", + "round_robin", + "--use-kl-loss", + "--kl-loss-coef", + "0.001", + "--wait", + ] + for path in args.extra_pythonpath: + command.extend(["--extra-pythonpath", path]) + + args.output_root.mkdir(parents=True, exist_ok=True) + controller_log = args.output_root / f"{run_id}.controller.log" + started = time.time() + print(json.dumps({"event": "start", "group": group, "run_id": run_id}), flush=True) + with controller_log.open("w", encoding="utf-8") as handle: + run_checked(command, stdout=handle) + + with (run_dir / "run.log").open("w", encoding="utf-8") as handle: + run_checked([str(args.ray_bin), "job", "logs", submission_id], stdout=handle) + run_checked( + [ + str(args.python), + str(args.rl_kernel_root / "examples/vime_qwen3_8b_tp4_cp2_200/validate_run.py"), + "--run-dir", + str(run_dir), + "--seal", + ] + ) + print( + json.dumps( + { + "event": "complete", + "group": group, + "run_id": run_id, + "elapsed_seconds": time.time() - started, + }, + sort_keys=True, + ), + flush=True, + ) + + +def main() -> int: + args = parse_args() + lock = args.output_root / f".{args.suite_id}.lock" + args.output_root.mkdir(parents=True, exist_ok=True) + try: + lock.touch(exist_ok=False) + except FileExistsError as exc: + raise RuntimeError(f"suite lock already exists: {lock}") from exc + try: + for group, rounds, seed in specs(args.phase): + run_one(args, group, rounds, seed) + finally: + lock.unlink(missing_ok=True) + return 0 + + +if __name__ == "__main__": + raise SystemExit(main()) diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/validate_artifacts.py b/examples/vime_qwen3_8b_tp4_cp2_200/validate_artifacts.py new file mode 100644 index 00000000..670087fc --- /dev/null +++ b/examples/vime_qwen3_8b_tp4_cp2_200/validate_artifacts.py @@ -0,0 +1,225 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors + +"""Validate CUDA-only framework readbacks and Vime train/rollout Logp dumps.""" + +from __future__ import annotations + +import argparse +import json +import math +from pathlib import Path +from typing import Any, Mapping + +import torch + +from rl_engine.integrations.runtime import _contains_triton, _runtime_platform + +_FRAMEWORKS = (("megatron", "training"), ("vllm", "rollout")) +_MODULES = ("attention", "ffn", "logp") +_STRICT_LOGP_BACKEND = "rlkernel.linear_logp.bitwise.v1" +_BACKEND_PREFIXES = ("rlkernel.", "pytorch-vocab-parallel-logp") + + +def load_readbacks(directory: Path) -> list[dict[str, Any]]: + values: list[dict[str, Any]] = [] + for path in sorted(directory.glob("*.json")): + value = json.loads(path.read_text(encoding="utf-8")) + if not isinstance(value, dict): + raise ValueError(f"readback must contain an object: {path}") + value["_path"] = str(path) + values.append(value) + if not values: + raise ValueError(f"no framework readbacks found in {directory}") + return values + + +def validate_readbacks(readbacks: list[dict[str, Any]]) -> dict[str, Any]: + errors: list[str] = [] + frameworks: dict[str, Any] = {} + for framework, target in _FRAMEWORKS: + matching = [ + value + for value in readbacks + if value.get("framework") == framework and value.get("target") == target + ] + label = f"{framework}/{target}" + if not matching: + errors.append(f"missing {label} readback") + continue + module_summary: dict[str, Any] = {} + for value in matching: + if value.get("fallbacks"): + errors.append(f"{label} recorded fallback: {value['fallbacks']}") + for module in _MODULES: + hook_count = sum( + module in value.get("installed_hooks", {}) for value in matching + ) + records = [ + value["operators"][module] + for value in matching + if isinstance(value.get("operators"), Mapping) + and module in value["operators"] + ] + call_count = sum(int(record.get("call_count", 0)) for record in records) + if hook_count == 0: + errors.append(f"{label} {module} hook was not installed") + if call_count == 0: + errors.append(f"{label} {module} had zero calls") + backends = sorted({str(record.get("backend_id", "")) for record in records}) + for record in records: + backend = str(record.get("backend_id", "")) + if module == "logp" and backend != _STRICT_LOGP_BACKEND: + errors.append( + f"{label} logp used {backend!r}, expected {_STRICT_LOGP_BACKEND!r}" + ) + elif not backend.startswith(_BACKEND_PREFIXES): + errors.append( + f"{label} {module} used unexpected backend {backend!r}" + ) + if _contains_triton(record): + errors.append(f"{label} {module} used Triton") + if _runtime_platform(record.get("provenance")) != "cuda": + errors.append(f"{label} {module} did not report CUDA execution") + module_summary[module] = { + "installed_processes": hook_count, + "call_count": call_count, + "backend_ids": backends, + } + frameworks[label] = { + "readback_count": len(matching), + "modules": module_summary, + } + return {"passed": not errors, "errors": errors, "frameworks": frameworks} + + +def _load_train_dump(path: Path) -> Mapping[str, Any]: + value = torch.load(path, map_location="cpu", weights_only=False) + if not isinstance(value, Mapping): + raise ValueError(f"train dump must contain a mapping: {path}") + return value + + +def compare_train_rollout_logps(paths: list[Path]) -> dict[str, Any]: + sample_count = 0 + element_count = 0 + mismatch_count = 0 + max_abs_diff = 0.0 + errors: list[str] = [] + for path in paths: + payload = _load_train_dump(path) + samples = payload.get("samples") + if not isinstance(samples, list): + rollout_data = payload.get("rollout_data") + if not isinstance(rollout_data, Mapping): + errors.append(f"{path} has neither samples nor rollout_data") + continue + training_values = rollout_data.get("log_probs") + rollout_values = rollout_data.get("rollout_log_probs") + if not isinstance(training_values, (list, tuple)) or not isinstance( + rollout_values, (list, tuple) + ): + errors.append( + f"{path} rollout_data lacks list log_probs/rollout_log_probs" + ) + continue + if len(training_values) != len(rollout_values): + errors.append( + f"{path} logprob list length mismatch: " + f"{len(training_values)} != {len(rollout_values)}" + ) + samples = [ + {"log_probs": training, "rollout_log_probs": rollout} + for training, rollout in zip( + training_values, rollout_values, strict=False + ) + ] + for sample_index, sample in enumerate(samples): + if not isinstance(sample, Mapping): + errors.append(f"{path} sample {sample_index} is not a mapping") + continue + training = sample.get("log_probs") + rollout = sample.get("rollout_log_probs") + if not isinstance(training, torch.Tensor) or not isinstance( + rollout, torch.Tensor + ): + errors.append( + f"{path} sample {sample_index} lacks tensor log_probs/rollout_log_probs" + ) + continue + sample_count += 1 + if training.shape != rollout.shape: + errors.append( + f"{path} sample {sample_index} shape mismatch: " + f"{tuple(training.shape)} != {tuple(rollout.shape)}" + ) + continue + if training.dtype != rollout.dtype: + errors.append( + f"{path} sample {sample_index} dtype mismatch: " + f"{training.dtype} != {rollout.dtype}" + ) + element_count += training.numel() + mismatch_count += int(torch.ne(training, rollout).sum().item()) + if training.numel(): + diff = (training.float() - rollout.float()).abs() + if not bool(torch.isfinite(diff).all().item()): + errors.append(f"{path} sample {sample_index} has non-finite drift") + else: + max_abs_diff = max(max_abs_diff, float(diff.max().item())) + if not paths: + errors.append("no Vime train dump was found") + if sample_count == 0: + errors.append("no comparable train/rollout samples were found") + torch_equal = not errors and mismatch_count == 0 + return { + "passed": torch_equal and max_abs_diff == 0.0, + "torch_equal": torch_equal, + "mismatch_count": mismatch_count, + "max_abs_diff": max_abs_diff if math.isfinite(max_abs_diff) else None, + "sample_count": sample_count, + "element_count": element_count, + "errors": errors, + "artifacts": [str(path) for path in paths], + } + + +def validate_artifacts(readback_dir: Path, train_data_dir: Path) -> dict[str, Any]: + readbacks = validate_readbacks(load_readbacks(readback_dir)) + train_paths = sorted(train_data_dir.glob("*.pt")) + bitwise = compare_train_rollout_logps(train_paths) + return { + "schema_version": "rlkernel.vime_cuda_bitwise_validation.v1", + "passed": bool(readbacks["passed"] and bitwise["passed"]), + "runtime_policy": {"platform": "cuda", "triton_allowed": False}, + "readbacks": readbacks, + "train_rollout_logp": bitwise, + } + + +def main(argv: list[str] | None = None) -> int: + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--readback-dir", type=Path, required=True) + parser.add_argument("--train-data-dir", type=Path, required=True) + parser.add_argument("--output", type=Path, required=True) + args = parser.parse_args(argv) + + try: + report = validate_artifacts(args.readback_dir, args.train_data_dir) + except Exception as exc: + report = { + "schema_version": "rlkernel.vime_cuda_bitwise_validation.v1", + "passed": False, + "runtime_policy": {"platform": "cuda", "triton_allowed": False}, + "error": str(exc), + } + args.output.parent.mkdir(parents=True, exist_ok=True) + args.output.write_text( + json.dumps(report, indent=2, sort_keys=True) + "\n", encoding="utf-8" + ) + print(json.dumps(report, indent=2, sort_keys=True)) + return 0 if report["passed"] else 1 + + +if __name__ == "__main__": + raise SystemExit(main()) diff --git a/examples/vime_qwen3_8b_tp4_cp2_200/validate_run.py b/examples/vime_qwen3_8b_tp4_cp2_200/validate_run.py new file mode 100755 index 00000000..70e56ca9 --- /dev/null +++ b/examples/vime_qwen3_8b_tp4_cp2_200/validate_run.py @@ -0,0 +1,501 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors + +"""Validate and optionally seal one append-only 200-rollout experiment arm.""" + +from __future__ import annotations + +import argparse +import ast +import json +import math +import re +from collections.abc import Mapping +from pathlib import Path +from typing import Any + +import torch + +from rl_engine.integrations.runtime import _contains_triton, _runtime_platform + + +RECORD_RE = re.compile(r"\b(rollout|step|perf)\s+(\d+):\s+(\{.*\})\s*$") +FRAMEWORKS = (("megatron", "training"), ("vllm", "rollout")) +MODULES = ("attention", "ffn", "logp") +EXPECTED_TOPOLOGY = { + "gpus": 8, + "actor_gpus": 8, + "rollout_gpus": 8, + "tp": 4, + "cp": 2, + "pp": 1, + "colocate": True, + "offload_train": False, + "offload_rollout": True, + "rollout_gpus_per_engine": 4, + "rollout_engines": 2, +} +CASE_FIELDS = { + "attention": "attention_case", + "ffn": "ffn_case", + "logp": "logp_case", +} +RL_KERNEL_LINEAR_LOGP_PROVIDER = ( + "rl_engine.integrations.vime.linear_logp_provider.provider" +) +VIME_NATIVE_LINEAR_LOGP_MARKER = ( + "linear_logp native active: " + "backend_id=vime.utils.ppo_utils.calculate_log_probs_and_entropy " + "contract_id=vime.native.linear_logp.v1 route=unconfigured device=cuda" +) +CUDA_GRAPH_LAUNCHER_MARKERS = ( + "required vLLM full-decode CUDA Graph capture sizes", + "strict vLLM full-decode CUDA Graph capture sizes", +) + + +def _load_json(path: Path) -> dict[str, Any]: + value = json.loads(path.read_text(encoding="utf-8")) + if not isinstance(value, dict): + raise ValueError(f"expected a JSON object: {path}") + return value + + +def _parse_runtime_records(log_text: str) -> dict[str, dict[int, dict[str, Any]]]: + records: dict[str, dict[int, dict[str, Any]]] = { + "rollout": {}, + "step": {}, + "perf": {}, + } + for line in log_text.splitlines(): + match = RECORD_RE.search(line) + if not match: + continue + try: + value = ast.literal_eval(match.group(3)) + except (SyntaxError, ValueError): + continue + if isinstance(value, dict): + records[match.group(1)][int(match.group(2))] = value + return records + + +def _validate_cudagraph(log_text: str, manifest: Mapping[str, Any]) -> dict[str, Any]: + execution = manifest.get("vllm_execution", {}) + capture_sizes = ( + execution.get("capture_sizes", []) if isinstance(execution, Mapping) else [] + ) + compact_sizes = "[" + ",".join(str(value) for value in capture_sizes) + "]" + checks = { + "launcher_marker": any( + f"{marker}: {compact_sizes}" in log_text + for marker in CUDA_GRAPH_LAUNCHER_MARKERS + ), + "engine_mode": bool(re.search(r"cudagraph_mode.*FULL_DECODE_ONLY", log_text)), + "not_eager": "enforce_eager=False" in log_text, + "capture_sizes": ( + f"'cudagraph_capture_sizes': {capture_sizes}" in log_text + or f'"cudagraph_capture_sizes":{compact_sizes}' in log_text + ), + } + return { + "passed": all(checks.values()), + "checks": checks, + "expected_capture_sizes": capture_sizes, + } + + +def _side(case_id: str, target: str) -> str: + training, rollout = case_id.split("/", 1) + selected = training if target == "training" else rollout + return "rl_kernel" if selected == "R" else "production" + + +def _load_readbacks(directory: Path) -> list[dict[str, Any]]: + values = [] + for path in sorted(directory.glob("*.json")): + value = _load_json(path) + value["_path"] = str(path) + values.append(value) + if not values: + raise ValueError(f"no framework readbacks found in {directory}") + return values + + +def _reported_backend_ids(value: Any) -> set[str]: + backend_ids: set[str] = set() + if isinstance(value, Mapping): + for key, nested in value.items(): + if key in {"actual_backend", "backend_id"} and isinstance(nested, str): + backend_ids.add(nested) + backend_ids.update(_reported_backend_ids(nested)) + elif isinstance(value, (list, tuple)): + for nested in value: + backend_ids.update(_reported_backend_ids(nested)) + return backend_ids + + +def _validate_readbacks( + readbacks: list[dict[str, Any]], arm: Mapping[str, Any], log_text: str +) -> dict[str, Any]: + errors: list[str] = [] + frameworks: dict[str, Any] = {} + for framework, target in FRAMEWORKS: + matching = [ + value + for value in readbacks + if value.get("framework") == framework and value.get("target") == target + ] + label = f"{framework}/{target}" + if not matching: + errors.append(f"missing {label} readback") + continue + for value in matching: + if value.get("fallbacks"): + errors.append(f"{label} recorded fallback in {value['_path']}") + + module_summary: dict[str, Any] = {} + for module in MODULES: + case_id = str(arm[CASE_FIELDS[module]]) + expected = _side(case_id, target) + records = [ + value["operators"][module] + for value in matching + if isinstance(value.get("operators"), Mapping) + and module in value["operators"] + ] + installed_count = sum( + module in value.get("installed_hooks", {}) for value in matching + ) + call_count = sum(int(record.get("call_count", 0)) for record in records) + implementations = sorted( + {str(record.get("implementation", "")) for record in records} + ) + backend_ids = sorted( + {str(record.get("backend_id", "")) for record in records} + ) + case_ids = sorted({str(record.get("case_id", "")) for record in records}) + native_megatron_logp = ( + framework == "megatron" + and target == "training" + and module == "logp" + and expected == "production" + ) + if native_megatron_logp: + marker_present = VIME_NATIVE_LINEAR_LOGP_MARKER in log_text + if installed_count: + errors.append( + f"{label} production logp unexpectedly installed an RL-Kernel hook" + ) + if records: + errors.append( + f"{label} production logp unexpectedly entered provider readback" + ) + if not marker_present: + errors.append( + f"{label} production logp did not report Vime's native backend marker" + ) + module_summary[module] = { + "case_id": case_id, + "expected_implementation": expected, + "installed_processes": installed_count, + "call_count": call_count, + "implementations": implementations, + "backend_ids": backend_ids, + "native_marker_present": marker_present, + "native_backend_id": ( + "vime.utils.ppo_utils.calculate_log_probs_and_entropy" + ), + } + continue + if installed_count == 0: + errors.append(f"{label} {module} hook was not installed") + if call_count == 0: + errors.append(f"{label} {module} had zero calls") + if any(value != expected for value in implementations): + errors.append( + f"{label} {module} implementation {implementations!r} != {expected!r}" + ) + if any(value != case_id for value in case_ids): + errors.append(f"{label} {module} case IDs {case_ids!r} != {case_id!r}") + for record in records: + if _contains_triton(record): + errors.append(f"{label} {module} used Triton") + if _runtime_platform(record.get("provenance")) != "cuda": + errors.append(f"{label} {module} did not report CUDA execution") + if record.get("provenance", {}).get("fallback") is True: + errors.append(f"{label} {module} provenance recorded fallback") + if expected == "rl_kernel" and not str( + record.get("backend_id", "") + ).startswith("rlkernel."): + errors.append(f"{label} {module} did not use an RL-Kernel backend") + provenance = record.get("provenance", {}) + reported_backend_ids = _reported_backend_ids(record) + strict_execution = ( + isinstance(provenance, Mapping) + and ( + provenance.get("deterministic_linear_logp") is True + or ( + isinstance(provenance.get("execution"), Mapping) + and provenance["execution"].get("strict_backend") is True + ) + ) + ) + if expected == "production" and ( + any(value.startswith("rlkernel.") for value in reported_backend_ids) + or strict_execution + ): + errors.append( + f"{label} {module} production route executed an RL-Kernel backend: " + f"{sorted(reported_backend_ids)!r}" + ) + module_summary[module] = { + "case_id": case_id, + "expected_implementation": expected, + "installed_processes": installed_count, + "call_count": call_count, + "implementations": implementations, + "backend_ids": backend_ids, + } + frameworks[label] = {"readback_count": len(matching), "modules": module_summary} + return {"passed": not errors, "errors": errors, "frameworks": frameworks} + + +def _numeric(value: Any, label: str, errors: list[str]) -> float | None: + if not isinstance(value, (int, float)) or isinstance(value, bool): + errors.append(f"{label} is missing or non-numeric") + return None + result = float(value) + if not math.isfinite(result): + errors.append(f"{label} is non-finite") + return None + return result + + +def _validate_runtime_logprobs( + records: Mapping[int, Mapping[str, Any]], + expected_rounds: int, + global_batch_size: int, + require_zero: bool, +) -> dict[str, Any]: + errors: list[str] = [] + rows: list[dict[str, Any]] = [] + for step in sorted(records): + record = records[step] + mismatch = _numeric( + record.get("train/train_current_rollout_logprob_mismatch_count"), + f"step {step} mismatch_count", + errors, + ) + maximum = _numeric( + record.get("train/train_current_rollout_logprob_max_abs_diff"), + f"step {step} max_abs_diff", + errors, + ) + mean_abs = _numeric( + record.get("train/train_rollout_logprob_abs_diff"), + f"step {step} mean_abs_diff", + errors, + ) + numel = _numeric( + record.get("train/train_current_rollout_logprob_numel"), + f"step {step} active_token_count", + errors, + ) + if numel is not None and numel <= 0: + errors.append(f"step {step} has no active tokens") + total_mismatches = None if mismatch is None else mismatch * global_batch_size + total_active_tokens = None if numel is None else numel * global_batch_size + rows.append( + { + "step": step, + "bitwise_mismatch_count": total_mismatches, + "max_abs_dlogp": maximum, + "mean_abs_dlogp": mean_abs, + "active_token_count": total_active_tokens, + "vime_mean_mismatch_count_per_sample": mismatch, + "vime_mean_active_tokens_per_sample": numel, + } + ) + if len(rows) != expected_rounds: + errors.append(f"observed {len(rows)} train steps, expected {expected_rounds}") + bitwise_zero = bool(rows) and all( + row["bitwise_mismatch_count"] == 0.0 and row["max_abs_dlogp"] == 0.0 + for row in rows + ) + if require_zero and not bitwise_zero: + errors.append("R/R arm did not achieve bitwise-zero runtime metrics") + return { + "passed": not errors, + "errors": errors, + "evidence_source": ( + "VIME runtime torch.ne/max metrics; VIME reports sample means, converted " + "to counts with global_batch_size" + ), + "bitwise_zero": bitwise_zero, + "rows": rows, + "total_active_token_exposure": sum( + row["active_token_count"] or 0.0 for row in rows + ), + } + + +def _inspect_offline_dumps(directory: Path) -> dict[str, Any]: + paths = sorted(directory.glob("*.pt")) + comparable = 0 + for path in paths: + payload = torch.load(path, map_location="cpu", weights_only=False) + rollout_data = ( + payload.get("rollout_data", {}) if isinstance(payload, Mapping) else {} + ) + if isinstance(rollout_data, Mapping) and "log_probs" in rollout_data: + comparable += 1 + return { + "status": "available" if paths and comparable == len(paths) else "unavailable", + "artifact_count": len(paths), + "comparable_artifact_count": comparable, + "reason": ( + None + if paths and comparable == len(paths) + else "current VIME dump lacks captured training log_probs; runtime exact metrics are used" + ), + } + + +def validate_run(run_dir: Path) -> dict[str, Any]: + manifest = _load_json(run_dir / "manifest.json") + arm = manifest.get("arm") + if not isinstance(arm, Mapping): + raise ValueError("manifest.arm is missing") + log_text = (run_dir / "run.log").read_text(encoding="utf-8", errors="replace") + records = _parse_runtime_records(log_text) + require_zero = all(str(arm[CASE_FIELDS[module]]) == "R/R" for module in MODULES) + cudagraph = _validate_cudagraph(log_text, manifest) + readbacks = _validate_readbacks( + _load_readbacks(run_dir / "readbacks"), arm, log_text + ) + logprobs = _validate_runtime_logprobs( + records["step"], + int(manifest["num_rollout"]), + int(manifest["batching"]["global_batch_size"]), + require_zero, + ) + global_errors = [] + algorithm = manifest.get("algorithm", {}) + if ( + not isinstance(algorithm, Mapping) + or algorithm.get("advantage_estimator") != "grpo" + ): + global_errors.append("manifest does not explicitly select GRPO") + train_command = manifest.get("train_command", []) + expected_algorithm_pair = ["--advantage-estimator", "grpo"] + if not isinstance(train_command, list) or not any( + train_command[index : index + 2] == expected_algorithm_pair + for index in range(max(0, len(train_command) - 1)) + ): + global_errors.append("train command does not explicitly select GRPO") + if manifest.get("topology") != EXPECTED_TOPOLOGY: + global_errors.append("manifest does not contain the required TP4/CP2 colocated topology") + required_command_pairs = ( + ("--actor-num-gpus-per-node", "8"), + ("--rollout-num-gpus", "8"), + ("--tensor-model-parallel-size", "4"), + ("--context-parallel-size", "2"), + ("--rollout-num-gpus-per-engine", "4"), + ) + if isinstance(train_command, list): + for flag, value in required_command_pairs: + if not any( + train_command[index : index + 2] == [flag, value] + for index in range(max(0, len(train_command) - 1)) + ): + global_errors.append(f"train command does not contain {flag} {value}") + if "--colocate" not in train_command: + global_errors.append("train command does not enable colocated execution") + if "--no-offload-train" not in train_command: + global_errors.append("train command does not keep the TP4 actor resident") + if "--offload-rollout" not in train_command: + global_errors.append("train command does not offload rollout during training") + training_logp_implementation = _side(str(arm["logp_case"]), "training") + provider_pair = ["--linear-logp-provider", RL_KERNEL_LINEAR_LOGP_PROVIDER] + strict_mode_pair = ["--linear-logp-provider-mode", "strict"] + has_provider = any( + train_command[index : index + 2] == provider_pair + for index in range(max(0, len(train_command) - 1)) + ) + has_strict_mode = any( + train_command[index : index + 2] == strict_mode_pair + for index in range(max(0, len(train_command) - 1)) + ) + if training_logp_implementation == "production": + if "--linear-logp-provider" in train_command: + global_errors.append( + "production Megatron logp must not configure a linear_logp provider" + ) + if "--linear-logp-provider-mode" in train_command: + global_errors.append( + "production Megatron logp must not configure provider mode" + ) + elif not has_provider or not has_strict_mode: + global_errors.append( + "RL-Kernel Megatron logp must configure the strict RL-Kernel provider" + ) + expected_recompute = { + "recompute_granularity": "full", + "recompute_method": "uniform", + "recompute_num_layers": 1, + } + if manifest.get("training_memory") != expected_recompute: + global_errors.append( + "manifest does not contain the required recompute configuration" + ) + if re.search(r"fallback=true", log_text, re.IGNORECASE): + global_errors.append("run log contains fallback=true") + if "Traceback (most recent call last)" in log_text: + global_errors.append("run log contains a traceback") + report = { + "schema_version": "rlkernel.vime_qwen3_8b_tp4_cp2_200.validation.v1", + "run_id": manifest.get("run_id"), + "group": arm.get("group"), + "passed": bool( + cudagraph["passed"] + and readbacks["passed"] + and logprobs["passed"] + and not global_errors + ), + "errors": global_errors, + "cudagraph": cudagraph, + "runtime_readbacks": readbacks, + "train_rollout_logprob": logprobs, + "offline_tensor_comparison": _inspect_offline_dumps(run_dir / "train-data"), + } + return report + + +def main(argv: list[str] | None = None) -> int: + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--run-dir", type=Path, required=True) + parser.add_argument("--output", type=Path, default=None) + parser.add_argument("--seal", action="store_true") + args = parser.parse_args(argv) + run_dir = args.run_dir.resolve() + output = args.output.resolve() if args.output else run_dir / "run-validation.json" + try: + report = validate_run(run_dir) + except Exception as exc: + report = { + "schema_version": "rlkernel.vime_qwen3_8b_tp4_cp2_200.validation.v1", + "passed": False, + "errors": [f"{type(exc).__name__}: {exc}"], + } + output.write_text( + json.dumps(report, indent=2, sort_keys=True) + "\n", encoding="utf-8" + ) + print(json.dumps(report, indent=2, sort_keys=True)) + if args.seal and report["passed"]: + (run_dir / "COMPLETE").touch(exist_ok=False) + return 0 if report["passed"] else 1 + + +if __name__ == "__main__": + raise SystemExit(main()) diff --git a/examples/vime_rocm_attention_ablation/README.md b/examples/vime_rocm_attention_ablation/README.md new file mode 100644 index 00000000..77207679 --- /dev/null +++ b/examples/vime_rocm_attention_ablation/README.md @@ -0,0 +1,190 @@ +# Vime ROCm Attention operator ablation + +This example runs one real Vime rollout/training step for each Attention +implementation pairing: + +| Case | Megatron training | vLLM rollout | Purpose | +| --- | --- | --- | --- | +| `P/P` | framework production | framework production | native baseline | +| `P/R` | framework production | RL-Kernel | rollout-side attribution | +| `R/P` | RL-Kernel | framework production | training-side attribution | +| `R/R` | RL-Kernel | RL-Kernel | strict ROCm control | + +Only `RL_KERNEL_ATTENTION_CASE` changes. FFN and Logp stay on strict `R/R`; model, +reference checkpoint, prompt data, topology, batch shape, optimizer, seeds, +dropout, and vLLM execution mode are frozen. Megatron sequence parallelism is +explicitly disabled in every arm: the strict projection hook does not yet +implement the complete SP all-gather/reduce-scatter contract. Context +parallelism remains supported and is validated using Vime's zigzag shards. + +Every arm explicitly selects `VLLM_ATTENTION_BACKEND=ROCM_AITER_FA`. The +Megatron actor loads the same Hugging Face weights used to initialize vLLM, +pins `--start-rollout-id 0`, and never reuses another arm's output. The +Attention-only matrix does not instantiate a zero-coefficient reference model; +the reference checkpoint remains sealed as an input for later full-path runs. + +`--get-mismatch-metrics` uses the included `metrics_only_tis` hook. The hook +returns the policy loss and masks unchanged—no TIS weighting or rejection is +introduced—and writes the rank/call evidence consumed by validation. + +This is the executable Attention **operator cross-configuration matrix**. It +does not claim to execute PR #230's compact `A0`-`A7` diagnostic taxonomy. +Those rows describe one-at-a-time root-cause probes; most still need concrete +runtime mutation and restoration hooks. No row label is treated as execution +evidence here. + +## Why every arm uses eager vLLM + +All four arms pass `--vllm-enforce-eager`. The current strict ROCm QKV/O +projection uses a fixed-tree collective with Python/lock bookkeeping, so it is +not a valid vLLM fullgraph capture target. Freezing eager mode across `P` and +`R` arms makes this a correctness-first, one-factor experiment. A future graph +path should be introduced as a separate controlled variable, not enabled only +for some arms. + +## Prerequisites + +Use a ROCm environment with Vime, Megatron-LM, vLLM, AITER, and this RL-Kernel +checkout installed (normally `pip install -e /work/RL-Kernel`). A source-only +`PYTHONPATH` entry is insufficient because vLLM discovers RL-Kernel through the +installed `vllm.general_plugins` entry point. The launcher is based on Vime's +Qwen3-8B AMD launcher and the existing `vime_qwen3_8b_tp2_cp2` example, but is +parameterized and avoids their CUDA-only flags. + +The default formal topology reuses PR #377's colocated eight-GPU schedule: + +- all eight GPUs host Megatron training with TP=4, CP=2, PP=1 and sequence parallelism disabled; +- the same eight GPUs host two vLLM TP=4 engines while rollout is active; +- the actor remains resident and rollout is offloaded between generation phases; +- the vLLM router is pinned to `round_robin`, with two independent prompt requests so both engines receive work; +- one rollout, 2 prompts, 1 sample per prompt, global batch 2; +- response length 32 and at most 256 packed tokens per training GPU; +- Qwen3-8B BF16 with deterministic seeds 1234/42. +- strict Logp vocabulary layout: 151936 real rows, padded to 152064 rows for TP4. + +The batch defaults are deliberately small correctness settings for the strict +path and may be overridden consistently on the CLI. On the rollout side, the current strict route logically gathers the +vLLM paged KV layout before invoking dense AITER Attention; it does not claim a +native paged Attention kernel or publishable performance numbers. + +The command refuses to reuse a non-empty run directory or an already-running +Ray cluster. Each arm receives its own dump, readback, and log directory. A +final checkpoint is written only when `RLK_ABLATION_SAVE_CHECKPOINT=1`; the +matrix never loads another arm's output checkpoint. Before creating the +run directory, it also requires the reference checkpoint's +`latest_checkpointed_iteration.txt` marker and verifies that the installed +RL-Kernel distribution exposes the expected vLLM plugin entry point. + +## Inspect the plan + +No configuration JSON is checked into the repository. Supply paths on the CLI +or through the matching environment variables: + +```bash +python examples/vime_rocm_attention_ablation/run.py \ + --vime-root /work/vime \ + --rl-kernel-root /work/RL-Kernel \ + --megatron-root /work/Megatron-LM-vime \ + --model-root /app/model/Qwen3-8B \ + --reference-checkpoint /app/model/Qwen3-8B_torch_dist \ + --prompt-data /app/model/dapo-math-17k/dapo-math-17k.jsonl +``` + +Without `--run`, this prints the exact four-arm plan and does not start Ray or +write results. + +## Execute all four arms + +```bash +python examples/vime_rocm_attention_ablation/run.py \ + --vime-root /work/vime \ + --rl-kernel-root /work/RL-Kernel \ + --megatron-root /work/Megatron-LM-vime \ + --model-root /app/model/Qwen3-8B \ + --reference-checkpoint /app/model/Qwen3-8B_torch_dist \ + --prompt-data /app/model/dapo-math-17k/dapo-math-17k.jsonl \ + --run-dir /work/RL-Kernel/runs/vime-rocm-attention-$(date -u +%Y%m%dT%H%M%SZ) \ + --run +``` + +The runner content-hashes the prompt dataset, launcher, and small checkpoint +index/config manifests. For the large model/checkpoint trees it seals every +relative file name, size, nanosecond mtime, and symlink target without rereading +all 8B weight shards. It also records each source revision, tracked dirty state, +and tracked-diff digest. Dirty checkouts are allowed, but the complete seal must +remain identical before and after the four arms. + +The default resource arguments are equivalent to: + +```bash +python examples/vime_rocm_attention_ablation/run.py \ + ... \ + --visible-gpus 0,1,2,3,4,5,6,7 \ + --num-gpus 8 \ + --tp-size 4 \ + --cp-size 2 \ + --rollout-tp-size 4 \ + --run +``` + +CP remains an Attention matrix dimension here; sequence parallelism remains +off even when CP is greater than one. + +## Runtime artifacts and validation + +Generated JSON is runtime evidence and lives only under `--run-dir` (the +repository's top-level `runs/` directory is ignored): + +```text +/ + matrix-plan.json + frozen-inputs.before.json + frozen-inputs.after.json + matrix-validation.json + arms/ + p-p/ + p-r/ + r-p/ + r-r/ + launch.json + launcher.log + validation.json + checkpoint/ # only when RLK_ABLATION_SAVE_CHECKPOINT=1 + dump/rollout_data/*.pt + mismatch_sidecars/*.pt + readbacks/ +``` + +An arm passes only after execution proves the requested route on both sides: + +- a Megatron/training and vLLM/rollout Attention hook was installed and called; +- `P` records production and does not resolve to RL-Kernel; +- `R` records ROCm execution, the strict AITER/CK Attention runtime/core and + fixed no-Split-KV schedule, no fallback/reference path, and the approved + deterministic `rlkernel.rocm.triton_det_gemm` QKV/O projection; +- the no-correction TIS hook atomically records rank/call-local training and + rollout logprobs, full response masks, and sequence lengths as `.pt` + sidecars; validation computes finite, non-empty mismatch count, max/mean + absolute drift, forward mismatch KL, and K3 KL from that evidence; +- CP sidecars are interpreted with Vime's two-ended zigzag response slice and + TP replicas are deduplicated. This observation path never depends on the + transient `partition` field that Vime removes before saving train debug data; +- the `R/R` control is bitwise equal for training versus rollout logprobs. + +The matrix additionally requires identical content fingerprints before/after +and identical rollout sample/token identity across all four arms. If Attention +changes generated tokens, the result remains useful operational evidence, but +the cross-arm metric comparison is marked invalid instead of comparing +different samples. + +To revalidate a completed arm without launching Vime: + +```bash +python examples/vime_rocm_attention_ablation/validate_artifacts.py \ + --arm-dir /path/to/run/arms/r-r \ + --case R/R +``` + +Do not add `matrix-plan.json`, validation JSON, rollout dumps, mismatch +sidecars, checkpoints, or MI300X result files to the PR. Publish them as CI/job +artifacts when needed. diff --git a/examples/vime_rocm_attention_ablation/__init__.py b/examples/vime_rocm_attention_ablation/__init__.py new file mode 100644 index 00000000..1d8ab8cb --- /dev/null +++ b/examples/vime_rocm_attention_ablation/__init__.py @@ -0,0 +1 @@ +"""ROCm Vime Attention operator cross-configuration example.""" diff --git a/examples/vime_rocm_attention_ablation/launch_arm.sh b/examples/vime_rocm_attention_ablation/launch_arm.sh new file mode 100644 index 00000000..c040d1e0 --- /dev/null +++ b/examples/vime_rocm_attention_ablation/launch_arm.sh @@ -0,0 +1,304 @@ +#!/usr/bin/env bash +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors + +# Launch one correctness-first Qwen3-8B/Vime ROCm Attention ablation arm. +# `run.py` supplies every RLK_ABLATION_* value and invokes this script once for +# each of P/P, P/R, R/P, and R/R. This launcher never reuses another arm's +# output checkpoint and never stops a Ray cluster it did not start. + +set -euo pipefail + +: "${RLK_ABLATION_CASE_ID:?}" +: "${RLK_ABLATION_ARM_DIR:?}" +: "${RLK_ABLATION_VIME_ROOT:?}" +: "${RLK_ABLATION_RL_KERNEL_ROOT:?}" +: "${RLK_ABLATION_MEGATRON_ROOT:?}" +: "${RLK_ABLATION_MODEL_ROOT:?}" +: "${RLK_ABLATION_REFERENCE_CHECKPOINT:?}" +: "${RLK_ABLATION_PROMPT_DATA:?}" +: "${RLK_ABLATION_NUM_GPUS:?}" +: "${RLK_ABLATION_TP_SIZE:?}" +: "${RLK_ABLATION_CP_SIZE:?}" +: "${RLK_ABLATION_ROLLOUT_TP_SIZE:?}" +: "${RLK_ABLATION_COLOCATE:?}" +: "${RLK_ABLATION_OFFLOAD_TRAIN:?}" +: "${RLK_ABLATION_OFFLOAD_ROLLOUT:?}" +: "${RLK_ABLATION_ROUTER_POLICY:?}" +: "${RLK_ABLATION_NUM_ROLLOUT:?}" +: "${RLK_ABLATION_ROLLOUT_BATCH_SIZE:?}" +: "${RLK_ABLATION_SAMPLES_PER_PROMPT:?}" +: "${RLK_ABLATION_GLOBAL_BATCH_SIZE:?}" +: "${RLK_ABLATION_MAX_RESPONSE_LENGTH:?}" +: "${RLK_ABLATION_MAX_TOKENS_PER_GPU:?}" +: "${RLK_ABLATION_SEED:?}" +: "${RLK_ABLATION_ROLLOUT_SEED:?}" +: "${RLK_ABLATION_RAY_PORT:?}" +: "${RLK_ABLATION_RAY_DASHBOARD_PORT:?}" +: "${RL_KERNEL_READBACK_DIR:?}" +: "${RL_KERNEL_MISMATCH_SIDECAR_DIR:?}" +: "${RL_KERNEL_VLLM_REAL_VOCAB_SIZE:?}" +: "${RL_KERNEL_VLLM_PADDED_VOCAB_SIZE:?}" + +case "${RLK_ABLATION_CASE_ID}" in + P/P|P/R|R/P|R/R) ;; + *) + echo "RLK_ABLATION_CASE_ID must be P/P, P/R, R/P, or R/R" >&2 + exit 2 + ;; +esac + +if [[ "${RL_KERNEL_ATTENTION_CASE:-}" != "${RLK_ABLATION_CASE_ID}" ]]; then + echo "RL_KERNEL_ATTENTION_CASE disagrees with the arm ID" >&2 + exit 2 +fi +if [[ "${RL_KERNEL_FFN_CASE:-}" != "R/R" || "${RL_KERNEL_LOGP_CASE:-}" != "R/R" ]]; then + echo "the strict dense matrix requires FFN=R/R and Logp=R/R" >&2 + exit 2 +fi +if [[ "${RL_KERNEL_VLLM_INTEGRATION:-}" != "1" ]]; then + echo "RL_KERNEL_VLLM_INTEGRATION=1 is required for rollout route readback" >&2 + exit 2 +fi + +TRAIN_GPUS=$((RLK_ABLATION_TP_SIZE * RLK_ABLATION_CP_SIZE)) +if [[ "${RLK_ABLATION_COLOCATE}" == "1" ]]; then + ROLLOUT_GPUS="${RLK_ABLATION_NUM_GPUS}" + if (( TRAIN_GPUS != RLK_ABLATION_NUM_GPUS )); then + echo "colocated training TP*CP must use all visible GPUs" >&2 + exit 2 + fi +else + ROLLOUT_GPUS=$((RLK_ABLATION_NUM_GPUS - TRAIN_GPUS)) +fi +if (( TRAIN_GPUS <= 0 || ROLLOUT_GPUS <= 0 )); then + echo "the requested TP/CP topology does not leave a valid rollout allocation" >&2 + exit 2 +fi +if (( ROLLOUT_GPUS % RLK_ABLATION_ROLLOUT_TP_SIZE != 0 )); then + echo "rollout GPU count must be divisible by rollout TP size" >&2 + exit 2 +fi +if [[ "${RLK_ABLATION_ROUTER_POLICY}" != "round_robin" ]]; then + echo "the two-engine strict matrix requires round_robin routing" >&2 + exit 2 +fi + +COLOCATE_ARGS=() +if [[ "${RLK_ABLATION_COLOCATE}" == "1" ]]; then + COLOCATE_ARGS+=(--colocate) +fi +if [[ "${RLK_ABLATION_OFFLOAD_TRAIN}" == "1" ]]; then + COLOCATE_ARGS+=(--offload-train) +else + COLOCATE_ARGS+=(--no-offload-train) +fi +if [[ "${RLK_ABLATION_OFFLOAD_ROLLOUT}" == "1" ]]; then + COLOCATE_ARGS+=(--offload-rollout) +else + COLOCATE_ARGS+=(--no-offload-rollout) +fi + +for required in \ + "${RLK_ABLATION_VIME_ROOT}/train.py" \ + "${RLK_ABLATION_VIME_ROOT}/scripts/models/qwen3-8B.sh" \ + "${RLK_ABLATION_RL_KERNEL_ROOT}/rl_engine" \ + "${RLK_ABLATION_RL_KERNEL_ROOT}/examples/vime_rocm_attention_ablation/tis_metrics.py" \ + "${RLK_ABLATION_MEGATRON_ROOT}" \ + "${RLK_ABLATION_MODEL_ROOT}" \ + "${RLK_ABLATION_REFERENCE_CHECKPOINT}" \ + "${RLK_ABLATION_PROMPT_DATA}" +do + if [[ ! -e "${required}" ]]; then + echo "required ROCm matrix path does not exist: ${required}" >&2 + exit 3 + fi +done + +unset CUBLASLT_WORKSPACE_SIZE CUBLAS_WORKSPACE_CONFIG NCCL_ALGO +unset RL_KERNEL_CUDA_ONLY RL_KERNEL_DET_GEMM_SM90_ONLY RL_KERNEL_PRECOMPILE_FA4 +unset VLLM_BATCH_INVARIANT NVTE_FUSED_ATTN NVTE_FLASH_ATTN NVTE_UNFUSED_ATTN + +export PYTHONUNBUFFERED=1 +export PYTHONPATH="${RLK_ABLATION_RL_KERNEL_ROOT}/examples:${RLK_ABLATION_RL_KERNEL_ROOT}:${RLK_ABLATION_VIME_ROOT}:${RLK_ABLATION_MEGATRON_ROOT}:${PYTHONPATH:-}" +export HIP_VISIBLE_DEVICES="${HIP_VISIBLE_DEVICES:?}" +export CUDA_VISIBLE_DEVICES="${HIP_VISIBLE_DEVICES}" +export RAY_EXPERIMENTAL_NOSET_HIP_VISIBLE_DEVICES=1 +export RAY_EXPERIMENTAL_NOSET_CUDA_VISIBLE_DEVICES=1 +export PYTORCH_ROCM_ARCH="${PYTORCH_ROCM_ARCH:-gfx942}" +export PYTORCH_ALLOC_CONF="${PYTORCH_ALLOC_CONF:-expandable_segments:True}" +export CUDA_DEVICE_MAX_CONNECTIONS=1 +export NCCL_NVLS_ENABLE=0 +export HSA_NO_SCRATCH_RECLAIM="${HSA_NO_SCRATCH_RECLAIM:-1}" +export VLLM_ROCM_USE_AITER=1 +# The strict paged materializer consumes AITER's packed NHD cache and rejects +# the optional shuffled physical layout. +export VLLM_ROCM_SHUFFLE_KV_CACHE_LAYOUT=0 +export VLLM_ATTENTION_BACKEND=ROCM_AITER_FA +export MIOPEN_DEBUG_CONV_DIRECT="${MIOPEN_DEBUG_CONV_DIRECT:-0}" + +mkdir -p \ + "${RLK_ABLATION_ARM_DIR}/dump" \ + "${RL_KERNEL_READBACK_DIR}" \ + "${RL_KERNEL_MISMATCH_SIDECAR_DIR}" + +SAVE_ARGS=() +if [[ "${RLK_ABLATION_SAVE_CHECKPOINT:-0}" == "1" ]]; then + mkdir -p "${RLK_ABLATION_ARM_DIR}/checkpoint" + SAVE_ARGS=( + --save "${RLK_ABLATION_ARM_DIR}/checkpoint" + --save-interval 1 + ) +fi + +python3 - <<'PY' +import os +import torch + +expected = int(os.environ["RLK_ABLATION_NUM_GPUS"]) +if torch.version.hip is None: + raise SystemExit("PyTorch is not a ROCm build") +if not torch.cuda.is_available() or torch.cuda.device_count() != expected: + raise SystemExit( + f"expected {expected} visible ROCm devices, got {torch.cuda.device_count()}" + ) +print(f"ROCm gate: HIP={torch.version.hip}, devices={torch.cuda.device_count()}") +PY + +cd "${RLK_ABLATION_VIME_ROOT}" +# shellcheck source=/dev/null +source "${RLK_ABLATION_VIME_ROOT}/scripts/models/qwen3-8B.sh" + +RUNTIME_ENV_JSON="$(python3 - <<'PY' +import json +import os + +names = [ + "PYTHONPATH", + "PYTHONUNBUFFERED", + "HIP_VISIBLE_DEVICES", + "CUDA_VISIBLE_DEVICES", + "RAY_EXPERIMENTAL_NOSET_HIP_VISIBLE_DEVICES", + "RAY_EXPERIMENTAL_NOSET_CUDA_VISIBLE_DEVICES", + "PYTORCH_ROCM_ARCH", + "PYTORCH_ALLOC_CONF", + "CUDA_DEVICE_MAX_CONNECTIONS", + "NCCL_NVLS_ENABLE", + "HSA_NO_SCRATCH_RECLAIM", + "VLLM_ROCM_USE_AITER", + "VLLM_ROCM_SHUFFLE_KV_CACHE_LAYOUT", + "VLLM_ATTENTION_BACKEND", + "MIOPEN_DEBUG_CONV_DIRECT", + "RL_KERNEL_ATTENTION_CASE", + "RL_KERNEL_FFN_CASE", + "RL_KERNEL_LOGP_CASE", + "RL_KERNEL_VLLM_REAL_VOCAB_SIZE", + "RL_KERNEL_VLLM_PADDED_VOCAB_SIZE", + "RL_KERNEL_VLLM_INTEGRATION", + "RL_KERNEL_READBACK_DIR", + "RL_KERNEL_MISMATCH_SIDECAR_DIR", +] +print(json.dumps({"env_vars": {name: os.environ[name] for name in names}})) +PY +)" + +if ray status >/dev/null 2>&1; then + echo "an existing Ray cluster is active; refusing to stop or reuse it" >&2 + exit 4 +fi + +ray_started=0 +cleanup_ray() { + if [[ "${ray_started}" == "1" ]]; then + ray stop --force >/dev/null 2>&1 || true + fi +} +trap cleanup_ray EXIT + +ray start --head \ + --node-ip-address=127.0.0.1 \ + --port="${RLK_ABLATION_RAY_PORT}" \ + --num-gpus="${RLK_ABLATION_NUM_GPUS}" \ + --disable-usage-stats \ + --dashboard-host=127.0.0.1 \ + --dashboard-port="${RLK_ABLATION_RAY_DASHBOARD_PORT}" +ray_started=1 + +# Correctness first: eager mode is frozen across all four arms. The strict +# ROCm projection collective performs Python/lock bookkeeping and is not a +# valid vLLM fullgraph capture target yet. +ray job submit \ + --address="http://127.0.0.1:${RLK_ABLATION_RAY_DASHBOARD_PORT}" \ + --runtime-env-json="${RUNTIME_ENV_JSON}" \ + --working-dir="${RLK_ABLATION_VIME_ROOT}" \ + -- python3 "${RLK_ABLATION_VIME_ROOT}/train.py" \ + --train-backend megatron \ + --actor-num-nodes 1 \ + --actor-num-gpus-per-node "${TRAIN_GPUS}" \ + --rollout-num-gpus "${ROLLOUT_GPUS}" \ + "${COLOCATE_ARGS[@]}" \ + --seed "${RLK_ABLATION_SEED}" \ + --rollout-seed "${RLK_ABLATION_ROLLOUT_SEED}" \ + "${MODEL_ARGS[@]}" \ + --hf-checkpoint "${RLK_ABLATION_MODEL_ROOT}" \ + --ref-load "${RLK_ABLATION_REFERENCE_CHECKPOINT}" \ + --load "${RLK_ABLATION_REFERENCE_CHECKPOINT}" \ + --start-rollout-id 0 \ + "${SAVE_ARGS[@]}" \ + --prompt-data "${RLK_ABLATION_PROMPT_DATA}" \ + --input-key prompt \ + --label-key label \ + --apply-chat-template \ + --rollout-shuffle \ + --rm-type deepscaler \ + --num-rollout "${RLK_ABLATION_NUM_ROLLOUT}" \ + --rollout-batch-size "${RLK_ABLATION_ROLLOUT_BATCH_SIZE}" \ + --n-samples-per-prompt "${RLK_ABLATION_SAMPLES_PER_PROMPT}" \ + --rollout-max-response-len "${RLK_ABLATION_MAX_RESPONSE_LENGTH}" \ + --rollout-temperature 1.0 \ + --rollout-top-p 1.0 \ + --global-batch-size "${RLK_ABLATION_GLOBAL_BATCH_SIZE}" \ + --balance-data \ + --optimizer adam \ + --lr 1e-6 \ + --lr-decay-style constant \ + --weight-decay 0.1 \ + --adam-beta1 0.9 \ + --adam-beta2 0.98 \ + --advantage-estimator grpo \ + --entropy-coef 0 \ + --eps-clip 0.2 \ + --eps-clip-high 0.28 \ + --tensor-model-parallel-size "${RLK_ABLATION_TP_SIZE}" \ + --context-parallel-size "${RLK_ABLATION_CP_SIZE}" \ + --pipeline-model-parallel-size 1 \ + --expert-model-parallel-size 1 \ + --expert-tensor-parallel-size 1 \ + --recompute-granularity full \ + --recompute-method uniform \ + --recompute-num-layers 1 \ + --use-dynamic-batch-size \ + --max-tokens-per-gpu "${RLK_ABLATION_MAX_TOKENS_PER_GPU}" \ + --router-policy "${RLK_ABLATION_ROUTER_POLICY}" \ + --rollout-num-gpus-per-engine "${RLK_ABLATION_ROLLOUT_TP_SIZE}" \ + --vllm-gpu-memory-utilization "${VLLM_GPU_MEMORY_UTILIZATION:-0.4}" \ + --vllm-attention-backend ROCM_AITER_FA \ + --vllm-enforce-eager \ + --vllm-disable-custom-all-reduce \ + --attention-dropout 0 \ + --hidden-dropout 0 \ + --accumulate-allreduce-grads-in-fp32 \ + --attention-softmax-in-fp32 \ + --attention-backend flash \ + --train-memory-margin-bytes 2147483648 \ + --no-gradient-accumulation-fusion \ + --linear-logp-provider \ + rl_engine.integrations.vime.linear_logp_provider.provider \ + --linear-logp-provider-mode strict \ + --get-mismatch-metrics \ + --custom-tis-function-path \ + vime_rocm_attention_ablation.tis_metrics.metrics_only_tis \ + --save-debug-rollout-data \ + "${RLK_ABLATION_ARM_DIR}/dump/rollout_data/{rollout_id}.pt" \ + --custom-megatron-init-path \ + rl_engine.integrations.megatron_runtime.initialize_from_environment diff --git a/examples/vime_rocm_attention_ablation/run.py b/examples/vime_rocm_attention_ablation/run.py new file mode 100644 index 00000000..3314b265 --- /dev/null +++ b/examples/vime_rocm_attention_ablation/run.py @@ -0,0 +1,795 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors + +"""Plan or run the four-arm ROCm Vime Attention implementation matrix. + +Every arm starts from the same model and Megatron checkpoint, uses the same +prompt data and seeds, and changes only ``RL_KERNEL_ATTENTION_CASE``. FFN and +Logp remain on the strict ``R/R`` path. Runtime JSON is written below ``--run-dir``; this example +does not require or ship a checked-in JSON configuration or result file. +""" + +from __future__ import annotations + +import argparse +import hashlib +import json +import os +import subprocess +from dataclasses import dataclass +from datetime import datetime, timezone +from importlib import metadata as importlib_metadata +from pathlib import Path +from typing import Any, Mapping + +try: + from .validate_artifacts import ( + CASE_IMPLEMENTATIONS, + validate_arm, + validate_matrix, + write_report, + ) +except ImportError: # Direct ``python examples/.../run.py`` execution. + from validate_artifacts import ( # type: ignore[no-redef] + CASE_IMPLEMENTATIONS, + validate_arm, + validate_matrix, + write_report, + ) + +PLAN_SCHEMA_VERSION = "rlkernel.vime_rocm_attention_operator_plan.v1" +FROZEN_SCHEMA_VERSION = "rlkernel.vime_rocm_attention_frozen_inputs.v1" +CASE_ORDER = ("P/P", "P/R", "R/P", "R/R") +RL_KERNEL_PLUGIN_ENTRY_POINT = "rl_engine.integrations.vllm_runtime:register_vllm_plugin" + +_CUDA_ONLY_ENVIRONMENT = ( + "CUBLASLT_WORKSPACE_SIZE", + "CUBLAS_WORKSPACE_CONFIG", + "NCCL_ALGO", + "RL_KERNEL_CUDA_ONLY", + "RL_KERNEL_DET_GEMM_SM90_ONLY", + "RL_KERNEL_PRECOMPILE_FA4", + "VLLM_BATCH_INVARIANT", +) + + +@dataclass(frozen=True) +class MatrixConfig: + vime_root: Path + rl_kernel_root: Path + megatron_root: Path + model_root: Path + reference_checkpoint: Path + prompt_data: Path + run_dir: Path + launcher: Path + visible_gpus: str = "0,1,2,3,4,5,6,7" + num_gpus: int = 8 + tensor_parallel_size: int = 4 + context_parallel_size: int = 2 + rollout_tensor_parallel_size: int = 4 + colocate: bool = True + offload_train: bool = False + offload_rollout: bool = True + router_policy: str = "round_robin" + num_rollout: int = 1 + rollout_batch_size: int = 2 + samples_per_prompt: int = 1 + global_batch_size: int = 2 + real_vocab_size: int = 151936 + padded_vocab_size: int = 152064 + max_response_length: int = 32 + max_tokens_per_gpu: int = 256 + seed: int = 1234 + rollout_seed: int = 42 + ray_port: int = 6385 + ray_dashboard_port: int = 28265 + + @property + def training_gpus(self) -> int: + return self.tensor_parallel_size * self.context_parallel_size + + @property + def rollout_gpus(self) -> int: + return self.num_gpus if self.colocate else self.num_gpus - self.training_gpus + + @property + def rollout_engines(self) -> int: + return self.rollout_gpus // self.rollout_tensor_parallel_size + + def validate(self, *, require_paths: bool) -> None: + visible = [item.strip() for item in self.visible_gpus.split(",") if item.strip()] + if len(visible) != self.num_gpus or len(set(visible)) != len(visible): + raise ValueError("visible_gpus must contain exactly num_gpus unique device IDs") + for name in ( + "num_gpus", + "tensor_parallel_size", + "context_parallel_size", + "rollout_tensor_parallel_size", + "num_rollout", + "rollout_batch_size", + "samples_per_prompt", + "global_batch_size", + "real_vocab_size", + "padded_vocab_size", + "max_response_length", + "max_tokens_per_gpu", + ): + if getattr(self, name) <= 0: + raise ValueError(f"{name} must be positive") + if self.training_gpus > self.num_gpus: + raise ValueError("training TP*CP cannot exceed num_gpus") + if self.colocate and self.training_gpus != self.num_gpus: + raise ValueError("colocated training TP*CP must use all visible GPUs") + if self.rollout_gpus <= 0: + raise ValueError("non-colocated training TP*CP must leave GPUs for rollout") + if self.rollout_gpus % self.rollout_tensor_parallel_size: + raise ValueError("rollout GPU count must be divisible by rollout TP") + if self.router_policy != "round_robin": + raise ValueError("the two-engine strict matrix requires round_robin routing") + if self.rollout_batch_size < self.rollout_engines: + raise ValueError( + "rollout_batch_size must issue at least one request per rollout engine" + ) + generated = self.rollout_batch_size * self.samples_per_prompt + if self.global_batch_size != generated: + raise ValueError( + "global_batch_size must equal rollout_batch_size*samples_per_prompt " + "for the one-step frozen matrix" + ) + if self.real_vocab_size > self.padded_vocab_size: + raise ValueError("real_vocab_size cannot exceed padded_vocab_size") + if self.padded_vocab_size % self.tensor_parallel_size: + raise ValueError("padded_vocab_size must be divisible by training TP") + if self.padded_vocab_size % 64: + raise ValueError("padded_vocab_size must be divisible by 64 vocab tiles") + if not 1024 <= self.ray_port <= 65535 or not 1024 <= self.ray_dashboard_port <= 65535: + raise ValueError("Ray ports must be between 1024 and 65535") + if abs(self.ray_port - self.ray_dashboard_port) < len(CASE_ORDER): + raise ValueError("Ray GCS and dashboard port ranges overlap") + dashboard_ports = range( + self.ray_dashboard_port, self.ray_dashboard_port + len(CASE_ORDER) + ) + if any(10001 <= port <= 19999 for port in dashboard_ports): + raise ValueError("Ray dashboard ports overlap the default client/worker range") + if not require_paths: + return + required = { + "vime_root": self.vime_root, + "rl_kernel_root": self.rl_kernel_root, + "megatron_root": self.megatron_root, + "model_root": self.model_root, + "reference_checkpoint": self.reference_checkpoint, + "prompt_data": self.prompt_data, + "launcher": self.launcher, + } + for name, path in required.items(): + if not path.exists(): + raise FileNotFoundError(f"{name} does not exist: {path}") + if not (self.vime_root / "train.py").is_file(): + raise FileNotFoundError(f"Vime train.py is missing below {self.vime_root}") + if not (self.vime_root / "scripts" / "models" / "qwen3-8B.sh").is_file(): + raise FileNotFoundError("Vime Qwen3-8B model argument script is missing") + if not (self.rl_kernel_root / "rl_engine").is_dir(): + raise FileNotFoundError("rl_kernel_root does not contain rl_engine") + metrics_hook = ( + self.rl_kernel_root + / "examples" + / "vime_rocm_attention_ablation" + / "tis_metrics.py" + ) + if not metrics_hook.is_file(): + raise FileNotFoundError(f"Attention mismatch metrics hook is missing: {metrics_hook}") + _validate_checkpoint_marker(self.reference_checkpoint) + _validate_rl_kernel_plugin_installation() + + def frozen_parameters(self) -> dict[str, Any]: + return { + "model": "Qwen/Qwen3-8B", + "visible_gpus": self.visible_gpus, + "num_gpus": self.num_gpus, + "training": { + "num_gpus": self.training_gpus, + "tensor_parallel_size": self.tensor_parallel_size, + "context_parallel_size": self.context_parallel_size, + "pipeline_parallel_size": 1, + "sequence_parallel": False, + "dtype": "bf16", + "attention_backend": "flash", + "attention_dropout": 0.0, + "hidden_dropout": 0.0, + }, + "rollout": { + "num_gpus": self.rollout_gpus, + "engine_count": self.rollout_engines, + "tensor_parallel_size": self.rollout_tensor_parallel_size, + "router_policy": self.router_policy, + "temperature": 1.0, + "top_p": 1.0, + # Vime's deterministic-inference flag exports + # VLLM_BATCH_INVARIANT=1, which native ROCM_AITER_FA correctly + # declares unsupported. The R route owns its deterministic + # per-row schedule; the P route remains the native baseline. + "vllm_batch_invariant": False, + "enforce_eager": True, + "custom_all_reduce": False, + "attention_backend": "ROCM_AITER_FA", + "shuffle_kv_cache_layout": False, + }, + "placement": { + "colocate": self.colocate, + "offload_train": self.offload_train, + "offload_rollout": self.offload_rollout, + }, + "batch": { + "start_rollout_id": 0, + "num_rollout": self.num_rollout, + "rollout_batch_size": self.rollout_batch_size, + "samples_per_prompt": self.samples_per_prompt, + "global_batch_size": self.global_batch_size, + "max_response_length": self.max_response_length, + "max_tokens_per_gpu": self.max_tokens_per_gpu, + }, + "optimizer": { + "name": "adam", + "lr": 1e-6, + "weight_decay": 0.1, + "beta1": 0.9, + "beta2": 0.98, + }, + "seed": self.seed, + "rollout_seed": self.rollout_seed, + "mismatch_metrics_hook": ( + "vime_rocm_attention_ablation.tis_metrics.metrics_only_tis" + ), + "ffn_case": "R/R", + "logp_case": "R/R", + "real_vocab_size": self.real_vocab_size, + "padded_vocab_size": self.padded_vocab_size, + "platform": "rocm", + } + + def paths(self) -> dict[str, str]: + return { + "vime_root": str(self.vime_root.resolve()), + "rl_kernel_root": str(self.rl_kernel_root.resolve()), + "megatron_root": str(self.megatron_root.resolve()), + "model_root": str(self.model_root.resolve()), + "reference_checkpoint": str(self.reference_checkpoint.resolve()), + "prompt_data": str(self.prompt_data.resolve()), + "launcher": str(self.launcher.resolve()), + "run_dir": str(self.run_dir.resolve()), + } + + +def _validate_checkpoint_marker(checkpoint: Path) -> None: + marker = checkpoint / "latest_checkpointed_iteration.txt" + if not marker.is_file(): + raise FileNotFoundError( + "reference_checkpoint is not a Megatron checkpoint: missing " + f"{marker}" + ) + value = marker.read_text(encoding="utf-8").strip() + if value == "release": + return + try: + iteration = int(value) + except ValueError as exc: + raise ValueError(f"invalid Megatron checkpoint marker {marker}: {value!r}") from exc + if iteration < 0: + raise ValueError(f"invalid Megatron checkpoint iteration in {marker}: {iteration}") + + +def _normalize_distribution_name(value: str) -> str: + return value.lower().replace("_", "-").replace(".", "-") + + +def _validate_rl_kernel_plugin_installation() -> None: + """Fail before Ray starts unless vLLM can discover the installed plugin.""" + + try: + distribution = importlib_metadata.distribution("RL-Kernel") + except importlib_metadata.PackageNotFoundError as exc: + raise RuntimeError( + "RL-Kernel must be installed (for example, `pip install -e `) " + "so vLLM can discover its plugin entry point" + ) from exc + + candidates = [ + entry_point + for entry_point in importlib_metadata.entry_points(group="vllm.general_plugins") + if entry_point.name == "rl_kernel" + ] + if not candidates: + raise RuntimeError( + "installed RL-Kernel does not expose the `rl_kernel` entry point in " + "vllm.general_plugins" + ) + matching = [ + entry_point + for entry_point in candidates + if entry_point.value == RL_KERNEL_PLUGIN_ENTRY_POINT + ] + if not matching: + values = sorted({entry_point.value for entry_point in candidates}) + raise RuntimeError( + "the visible rl_kernel vLLM plugin entry point has an unexpected target: " + f"{values!r}" + ) + + installed_name = distribution.metadata.get("Name", "RL-Kernel") + owners = { + entry_point.dist.metadata.get("Name", entry_point.dist.name) + for entry_point in matching + if getattr(entry_point, "dist", None) is not None + } + if owners and _normalize_distribution_name(installed_name) not in { + _normalize_distribution_name(owner) for owner in owners + }: + raise RuntimeError( + "the visible rl_kernel vLLM plugin entry point is not owned by the " + "installed RL-Kernel distribution" + ) + + +def _canonical_fingerprint(value: Mapping[str, Any]) -> str: + encoded = json.dumps(value, sort_keys=True, separators=(",", ":")).encode("utf-8") + return hashlib.sha256(encoded).hexdigest() + + +def _hash_file(path: Path) -> tuple[str, int]: + digest = hashlib.sha256() + size = 0 + with path.open("rb") as handle: + while chunk := handle.read(1024 * 1024): + digest.update(chunk) + size += len(chunk) + return digest.hexdigest(), size + + +def fingerprint_path(path: Path) -> dict[str, Any]: + """Seal a file or tree without rereading every multi-GB weight shard. + + Individual files (the prompt dataset and launcher) are content-hashed. For + checkpoint trees, every relative path, size, mtime, and symlink target is + sealed, while small checkpoint/config manifests receive an additional + content hash. + """ + + resolved = path.resolve() + if not resolved.exists(): + raise FileNotFoundError(resolved) + single_file = resolved.is_file() + files = ( + [resolved] + if single_file + else sorted(item for item in resolved.rglob("*") if item.is_file()) + ) + aggregate = hashlib.sha256() + byte_count = 0 + content_hashed = 0 + for item in files: + relative = item.name if single_file else item.relative_to(resolved).as_posix() + stat = item.stat() + size = stat.st_size + link_target = os.readlink(item) if item.is_symlink() else None + manifest = ( + single_file + or item.name == "latest_checkpointed_iteration.txt" + or item.name.endswith(".index.json") + or item.name + in { + "config.json", + "generation_config.json", + "metadata.json", + "tokenizer.json", + "tokenizer_config.json", + } + ) + record = { + "path": relative, + "size": size, + "mtime_ns": stat.st_mtime_ns, + "symlink": link_target, + "content_sha256": _hash_file(item)[0] if manifest else None, + } + content_hashed += int(manifest) + aggregate.update(json.dumps(record, sort_keys=True, separators=(",", ":")).encode()) + aggregate.update(b"\n") + byte_count += size + return { + "path": str(resolved), + "kind": "file" if single_file else "directory", + "seal_mode": "content" if single_file else "metadata_plus_checkpoint_manifests", + "file_count": len(files), + "content_hashed_file_count": content_hashed, + "byte_count": byte_count, + "sha256": aggregate.hexdigest(), + } + + +def _git_identity(path: Path) -> dict[str, Any]: + def run(*args: str) -> str: + result = subprocess.run( + ["git", "-C", str(path), *args], + check=True, + capture_output=True, + text=True, + ) + return result.stdout + + try: + revision = run("rev-parse", "HEAD").strip() + status = run("status", "--porcelain=v1", "--untracked-files=no") + diff = subprocess.run( + ["git", "-C", str(path), "diff", "--binary", "HEAD"], + check=True, + capture_output=True, + ).stdout + except (OSError, subprocess.CalledProcessError) as exc: + raise ValueError(f"source root is not a readable Git checkout: {path}") from exc + return { + "path": str(path.resolve()), + "revision": revision, + "tracked_dirty": bool(status.strip()), + "tracked_status_sha256": hashlib.sha256(status.encode()).hexdigest(), + "tracked_diff_sha256": hashlib.sha256(diff).hexdigest(), + } + + +def frozen_input_manifest(config: MatrixConfig) -> dict[str, Any]: + """Seal inputs and tracked source state before/after the four executions.""" + + payload = { + "schema_version": FROZEN_SCHEMA_VERSION, + "parameters": config.frozen_parameters(), + "inputs": { + "model_root": fingerprint_path(config.model_root), + "reference_checkpoint": fingerprint_path(config.reference_checkpoint), + "prompt_data": fingerprint_path(config.prompt_data), + }, + "sources": { + "vime": _git_identity(config.vime_root), + "rl_kernel": _git_identity(config.rl_kernel_root), + "megatron": _git_identity(config.megatron_root), + }, + "launcher": fingerprint_path(config.launcher), + } + payload["fingerprint"] = _canonical_fingerprint(payload) + return payload + + +def case_slug(case_id: str) -> str: + if case_id not in CASE_IMPLEMENTATIONS: + raise ValueError(f"unknown case {case_id!r}") + return case_id.lower().replace("/", "-") + + +def build_arm_environment( + config: MatrixConfig, + case_id: str, + arm_dir: Path, + *, + arm_index: int, + base_environment: Mapping[str, str] | None = None, +) -> dict[str, str]: + """Build the exact inherited environment for one Vime/Ray arm.""" + + if case_id not in CASE_IMPLEMENTATIONS: + raise ValueError(f"unknown case {case_id!r}") + env = dict(os.environ if base_environment is None else base_environment) + for name in _CUDA_ONLY_ENVIRONMENT: + env.pop(name, None) + existing_pythonpath = env.get("PYTHONPATH", "") + python_paths = [ + str((config.rl_kernel_root / "examples").resolve()), + str(config.rl_kernel_root.resolve()), + str(config.vime_root.resolve()), + str(config.megatron_root.resolve()), + ] + if existing_pythonpath: + python_paths.append(existing_pythonpath) + env.update( + { + "PYTHONPATH": os.pathsep.join(python_paths), + "PYTHONUNBUFFERED": "1", + "HIP_VISIBLE_DEVICES": config.visible_gpus, + # PyTorch on ROCm and Ray still consume CUDA_VISIBLE_DEVICES. + "CUDA_VISIBLE_DEVICES": config.visible_gpus, + "RAY_EXPERIMENTAL_NOSET_HIP_VISIBLE_DEVICES": "1", + "RAY_EXPERIMENTAL_NOSET_CUDA_VISIBLE_DEVICES": "1", + "VLLM_ROCM_USE_AITER": "1", + "VLLM_ROCM_SHUFFLE_KV_CACHE_LAYOUT": "0", + "VLLM_ATTENTION_BACKEND": "ROCM_AITER_FA", + "RL_KERNEL_ATTENTION_CASE": case_id, + "RL_KERNEL_FFN_CASE": "R/R", + "RL_KERNEL_LOGP_CASE": "R/R", + "RL_KERNEL_VLLM_REAL_VOCAB_SIZE": str(config.real_vocab_size), + "RL_KERNEL_VLLM_PADDED_VOCAB_SIZE": str(config.padded_vocab_size), + "RL_KERNEL_VLLM_INTEGRATION": "1", + "RL_KERNEL_READBACK_DIR": str((arm_dir / "readbacks").resolve()), + "RL_KERNEL_MISMATCH_SIDECAR_DIR": str( + (arm_dir / "mismatch_sidecars").resolve() + ), + "RLK_ABLATION_CASE_ID": case_id, + "RLK_ABLATION_ARM_DIR": str(arm_dir.resolve()), + "RLK_ABLATION_VIME_ROOT": str(config.vime_root.resolve()), + "RLK_ABLATION_RL_KERNEL_ROOT": str(config.rl_kernel_root.resolve()), + "RLK_ABLATION_MEGATRON_ROOT": str(config.megatron_root.resolve()), + "RLK_ABLATION_MODEL_ROOT": str(config.model_root.resolve()), + "RLK_ABLATION_REFERENCE_CHECKPOINT": str( + config.reference_checkpoint.resolve() + ), + "RLK_ABLATION_PROMPT_DATA": str(config.prompt_data.resolve()), + "RLK_ABLATION_NUM_GPUS": str(config.num_gpus), + "RLK_ABLATION_TP_SIZE": str(config.tensor_parallel_size), + "RLK_ABLATION_CP_SIZE": str(config.context_parallel_size), + "RLK_ABLATION_ROLLOUT_TP_SIZE": str(config.rollout_tensor_parallel_size), + "RLK_ABLATION_COLOCATE": "1" if config.colocate else "0", + "RLK_ABLATION_OFFLOAD_TRAIN": "1" if config.offload_train else "0", + "RLK_ABLATION_OFFLOAD_ROLLOUT": "1" if config.offload_rollout else "0", + "RLK_ABLATION_ROUTER_POLICY": config.router_policy, + "RLK_ABLATION_NUM_ROLLOUT": str(config.num_rollout), + "RLK_ABLATION_ROLLOUT_BATCH_SIZE": str(config.rollout_batch_size), + "RLK_ABLATION_SAMPLES_PER_PROMPT": str(config.samples_per_prompt), + "RLK_ABLATION_GLOBAL_BATCH_SIZE": str(config.global_batch_size), + "RLK_ABLATION_MAX_RESPONSE_LENGTH": str(config.max_response_length), + "RLK_ABLATION_MAX_TOKENS_PER_GPU": str(config.max_tokens_per_gpu), + "RLK_ABLATION_SEED": str(config.seed), + "RLK_ABLATION_ROLLOUT_SEED": str(config.rollout_seed), + "RLK_ABLATION_RAY_PORT": str(config.ray_port + arm_index), + "RLK_ABLATION_RAY_DASHBOARD_PORT": str( + config.ray_dashboard_port + arm_index + ), + } + ) + return env + + +def public_arm_environment(environment: Mapping[str, str]) -> dict[str, str]: + """Return only experiment variables; never serialize arbitrary host secrets.""" + + names = { + "CUDA_VISIBLE_DEVICES", + "HIP_VISIBLE_DEVICES", + "RL_KERNEL_ATTENTION_CASE", + "RL_KERNEL_FFN_CASE", + "RL_KERNEL_LOGP_CASE", + "RL_KERNEL_VLLM_REAL_VOCAB_SIZE", + "RL_KERNEL_VLLM_PADDED_VOCAB_SIZE", + "RL_KERNEL_MISMATCH_SIDECAR_DIR", + "RL_KERNEL_READBACK_DIR", + "RL_KERNEL_VLLM_INTEGRATION", + "VLLM_ROCM_USE_AITER", + "VLLM_ROCM_SHUFFLE_KV_CACHE_LAYOUT", + "VLLM_ATTENTION_BACKEND", + } + return { + name: value + for name, value in sorted(environment.items()) + if name in names or name.startswith("RLK_ABLATION_") + } + + +def build_plan(config: MatrixConfig) -> dict[str, Any]: + arms: list[dict[str, Any]] = [] + for index, case_id in enumerate(CASE_ORDER): + arm_dir = config.run_dir / "arms" / case_slug(case_id) + environment = build_arm_environment( + config, + case_id, + arm_dir, + arm_index=index, + base_environment={}, + ) + arms.append( + { + "case_id": case_id, + "training_implementation": CASE_IMPLEMENTATIONS[case_id]["training"], + "rollout_implementation": CASE_IMPLEMENTATIONS[case_id]["rollout"], + "arm_dir": str(arm_dir.resolve()), + "command": ["bash", str(config.launcher.resolve())], + "environment": public_arm_environment(environment), + } + ) + payload = { + "schema_version": PLAN_SCHEMA_VERSION, + "matrix_kind": "attention_operator_implementation_cross_config", + "parameters": config.frozen_parameters(), + "paths": config.paths(), + "arms": arms, + "claim_boundary": { + "implemented": ["P/P", "P/R", "R/P", "R/R"], + "not_implemented": "A0-A7 runtime mutation matrix", + }, + } + payload["plan_fingerprint"] = _canonical_fingerprint(payload) + return payload + + +def _prepare_run_dir(path: Path) -> None: + if path.exists() and any(path.iterdir()): + raise FileExistsError(f"run directory is not empty: {path}") + path.mkdir(parents=True, exist_ok=True) + + +def execute_matrix(config: MatrixConfig, *, fail_fast: bool = False) -> dict[str, Any]: + config.validate(require_paths=True) + _prepare_run_dir(config.run_dir) + plan = build_plan(config) + write_report(config.run_dir / "matrix-plan.json", plan) + + frozen_before = frozen_input_manifest(config) + write_report(config.run_dir / "frozen-inputs.before.json", frozen_before) + reports: dict[str, Mapping[str, Any]] = {} + + for index, case_id in enumerate(CASE_ORDER): + arm_dir = config.run_dir / "arms" / case_slug(case_id) + readback_dir = arm_dir / "readbacks" + dump_dir = arm_dir / "dump" + checkpoint_dir = arm_dir / "checkpoint" + sidecar_dir = arm_dir / "mismatch_sidecars" + for directory in (readback_dir, dump_dir, checkpoint_dir, sidecar_dir): + directory.mkdir(parents=True, exist_ok=False) + environment = build_arm_environment( + config, + case_id, + arm_dir, + arm_index=index, + ) + arm_manifest = { + "schema_version": "rlkernel.vime_rocm_attention_arm_launch.v1", + "case_id": case_id, + "expected_implementations": CASE_IMPLEMENTATIONS[case_id], + "frozen_input_fingerprint": frozen_before["fingerprint"], + "command": ["bash", str(config.launcher.resolve())], + "environment": public_arm_environment(environment), + "started_at": datetime.now(timezone.utc).isoformat(), + } + write_report(arm_dir / "launch.json", arm_manifest) + + log_path = arm_dir / "launcher.log" + returncode = 127 + try: + with log_path.open("w", encoding="utf-8") as log_handle: + process = subprocess.run( + ["bash", str(config.launcher.resolve())], + cwd=config.rl_kernel_root, + env=environment, + stdout=log_handle, + stderr=subprocess.STDOUT, + check=False, + ) + returncode = process.returncode + except OSError as exc: + log_path.write_text(f"failed to start launcher: {exc}\n", encoding="utf-8") + + report = validate_arm(arm_dir, case_id, launcher_returncode=returncode) + write_report(arm_dir / "validation.json", report) + reports[case_id] = report + if fail_fast and report["passed"] is not True: + break + + frozen_after = frozen_input_manifest(config) + write_report(config.run_dir / "frozen-inputs.after.json", frozen_after) + summary = validate_matrix( + reports, + frozen_before=frozen_before, + frozen_after=frozen_after, + ) + summary = { + **summary, + "created_at": datetime.now(timezone.utc).isoformat(), + "run_dir": str(config.run_dir.resolve()), + "plan_fingerprint": plan["plan_fingerprint"], + } + write_report(config.run_dir / "matrix-validation.json", summary) + return summary + + +def _path_argument(value: str | None, environment_name: str) -> Path | None: + selected = value or os.getenv(environment_name) + return None if not selected else Path(selected) + + +def build_parser() -> argparse.ArgumentParser: + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--vime-root", default=None) + parser.add_argument("--rl-kernel-root", default=None) + parser.add_argument("--megatron-root", default=None) + parser.add_argument("--model-root", default=None) + parser.add_argument("--reference-checkpoint", default=None) + parser.add_argument("--prompt-data", default=None) + parser.add_argument( + "--run-dir", + type=Path, + default=Path("runs") + / "vime_rocm_attention_ablation" + / datetime.now(timezone.utc).strftime("%Y%m%dT%H%M%SZ"), + ) + parser.add_argument( + "--launcher", + type=Path, + default=Path(__file__).with_name("launch_arm.sh"), + ) + parser.add_argument("--visible-gpus", default="0,1,2,3,4,5,6,7") + parser.add_argument("--num-gpus", type=int, default=8) + parser.add_argument("--tp-size", type=int, default=4) + parser.add_argument("--cp-size", type=int, default=2) + parser.add_argument("--rollout-tp-size", type=int, default=4) + parser.add_argument("--num-rollout", type=int, default=1) + parser.add_argument("--rollout-batch-size", type=int, default=1) + parser.add_argument("--samples-per-prompt", type=int, default=2) + parser.add_argument("--global-batch-size", type=int, default=2) + parser.add_argument("--max-response-length", type=int, default=32) + parser.add_argument("--max-tokens-per-gpu", type=int, default=256) + parser.add_argument("--seed", type=int, default=1234) + parser.add_argument("--rollout-seed", type=int, default=42) + parser.add_argument("--ray-port", type=int, default=6385) + parser.add_argument("--ray-dashboard-port", type=int, default=28265) + parser.add_argument("--run", action="store_true", help="execute all four Vime arms") + parser.add_argument("--fail-fast", action="store_true") + return parser + + +def config_from_args(args: argparse.Namespace) -> MatrixConfig: + values = { + "vime_root": _path_argument(args.vime_root, "VIME_ROOT"), + "rl_kernel_root": _path_argument(args.rl_kernel_root, "RL_KERNEL_ROOT"), + "megatron_root": _path_argument(args.megatron_root, "MEGATRON_ROOT"), + "model_root": _path_argument(args.model_root, "MODEL_ROOT"), + "reference_checkpoint": _path_argument( + args.reference_checkpoint, "TORCH_DIST_ROOT" + ), + "prompt_data": _path_argument(args.prompt_data, "PROMPT_DATA"), + } + missing = [name for name, value in values.items() if value is None] + if missing: + options = ", ".join("--" + name.replace("_", "-") for name in missing) + raise ValueError(f"missing required paths: {options}") + return MatrixConfig( + **values, # type: ignore[arg-type] + run_dir=args.run_dir, + launcher=args.launcher, + visible_gpus=args.visible_gpus, + num_gpus=args.num_gpus, + tensor_parallel_size=args.tp_size, + context_parallel_size=args.cp_size, + rollout_tensor_parallel_size=args.rollout_tp_size, + num_rollout=args.num_rollout, + rollout_batch_size=args.rollout_batch_size, + samples_per_prompt=args.samples_per_prompt, + global_batch_size=args.global_batch_size, + max_response_length=args.max_response_length, + max_tokens_per_gpu=args.max_tokens_per_gpu, + seed=args.seed, + rollout_seed=args.rollout_seed, + ray_port=args.ray_port, + ray_dashboard_port=args.ray_dashboard_port, + ) + + +def main(argv: list[str] | None = None) -> int: + args = build_parser().parse_args(argv) + try: + config = config_from_args(args) + config.validate(require_paths=args.run) + if args.run: + report = execute_matrix(config, fail_fast=args.fail_fast) + else: + report = {"status": "planned", **build_plan(config)} + except Exception as exc: + report = { + "schema_version": PLAN_SCHEMA_VERSION, + "status": "error", + "error_type": f"{type(exc).__module__}.{type(exc).__qualname__}", + "error": str(exc), + } + print(json.dumps(report, indent=2, sort_keys=True)) + return 2 + print(json.dumps(report, indent=2, sort_keys=True)) + return 0 if report.get("passed", True) else 1 + + +if __name__ == "__main__": + raise SystemExit(main()) diff --git a/examples/vime_rocm_attention_ablation/tis_metrics.py b/examples/vime_rocm_attention_ablation/tis_metrics.py new file mode 100644 index 00000000..c71e4e30 --- /dev/null +++ b/examples/vime_rocm_attention_ablation/tis_metrics.py @@ -0,0 +1,138 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors + +"""Mismatch metrics hook that deliberately leaves Vime's policy loss unchanged.""" + +from __future__ import annotations + +import itertools +import os +import threading +from pathlib import Path +from typing import Any + +import torch +import torch.distributed as dist + +SIDECAR_SCHEMA_VERSION = "rlkernel.vime_rocm_attention_mismatch_sidecar.v1" +SIDECAR_DIRECTORY_ENV = "RL_KERNEL_MISMATCH_SIDECAR_DIR" + +_CALL_COUNTER = itertools.count() +_CALL_COUNTER_LOCK = threading.Lock() + + +def _cpu_vector(value: Any, *, label: str) -> torch.Tensor: + try: + tensor = value if isinstance(value, torch.Tensor) else torch.as_tensor(value) + except Exception as exc: + raise ValueError(f"{label} is not tensor-like") from exc + return tensor.detach().to(device="cpu").reshape(-1).contiguous() + + +def _global_rank() -> int: + if dist.is_available() and dist.is_initialized(): + return dist.get_rank() + try: + return int(os.environ.get("RANK", "0")) + except ValueError as exc: + raise ValueError("RANK must be an integer when torch.distributed is unavailable") from exc + + +def _write_sidecar( + args: Any, + *, + train_log_probs: list[torch.Tensor], + rollout_log_probs: list[torch.Tensor], + loss_masks: list[torch.Tensor], + total_lengths: list[int], + response_lengths: list[int], +) -> None: + directory_value = os.environ.get(SIDECAR_DIRECTORY_ENV) + if not directory_value: + raise RuntimeError( + f"{SIDECAR_DIRECTORY_ENV} must name an arm-local directory for mismatch evidence" + ) + count = len(train_log_probs) + fields = { + "rollout_log_probs": rollout_log_probs, + "loss_masks": loss_masks, + "total_lengths": total_lengths, + "response_lengths": response_lengths, + } + if any(len(value) != count for value in fields.values()): + lengths = {"train_log_probs": count, **{key: len(value) for key, value in fields.items()}} + raise ValueError(f"mismatch sidecar fields have different sample counts: {lengths}") + + rank = _global_rank() + with _CALL_COUNTER_LOCK: + call_index = next(_CALL_COUNTER) + payload = { + "schema_version": SIDECAR_SCHEMA_VERSION, + "rank": rank, + "call_index": call_index, + "tensor_parallel_size": int(args.tensor_model_parallel_size), + "context_parallel_size": int(args.context_parallel_size), + "train_log_probs": [ + _cpu_vector(value, label="train_log_probs") for value in train_log_probs + ], + "rollout_log_probs": [ + _cpu_vector(value, label="rollout_log_probs") + for value in rollout_log_probs + ], + "loss_masks": [_cpu_vector(value, label="loss_masks") for value in loss_masks], + "total_lengths": [int(value) for value in total_lengths], + "response_lengths": [int(value) for value in response_lengths], + } + + directory = Path(directory_value) + directory.mkdir(parents=True, exist_ok=True) + path = directory / f"rank{rank:05d}.call{call_index:08d}.pt" + temporary = path.with_name(f".{path.name}.{os.getpid()}.tmp") + torch.save(payload, temporary) + os.replace(temporary, path) + + +def metrics_only_tis( + args: Any, + *, + pg_loss: torch.Tensor, + train_log_probs: list[torch.Tensor], + rollout_log_probs: list[torch.Tensor], + loss_masks: list[torch.Tensor], + total_lengths: list[int], + response_lengths: list[int], + **_: Any, +) -> tuple[torch.Tensor, list[torch.Tensor], dict[str, torch.Tensor]]: + """Return TIS diagnostics without applying TIS weights or rejection masks. + + Current Vime requires ``--custom-tis-function-path`` whenever + ``--get-mismatch-metrics`` is enabled. Its built-in fallback multiplies + ``pg_loss`` by clipped importance weights, which would introduce a second + experimental variable into this Attention-only matrix. This hook reports + the same ratio diagnostics while returning both loss inputs verbatim. + """ + + training = torch.cat(train_log_probs, dim=0).detach() + rollout = torch.cat(rollout_log_probs, dim=0).detach() + if training.shape != rollout.shape: + raise ValueError( + "training and rollout log probabilities must have identical shapes: " + f"{tuple(training.shape)} != {tuple(rollout.shape)}" + ) + + ratio = torch.exp(training - rollout) + clipped = torch.clamp(ratio, min=args.tis_clip_low, max=args.tis_clip) + metrics = { + "tis": ratio, + "tis_clipfrac": (clipped != ratio).to(dtype=ratio.dtype), + "tis_abs": (ratio - 1).abs(), + } + _write_sidecar( + args, + train_log_probs=train_log_probs, + rollout_log_probs=rollout_log_probs, + loss_masks=loss_masks, + total_lengths=total_lengths, + response_lengths=response_lengths, + ) + return pg_loss, loss_masks, metrics diff --git a/examples/vime_rocm_attention_ablation/validate_artifacts.py b/examples/vime_rocm_attention_ablation/validate_artifacts.py new file mode 100644 index 00000000..e7ae89d9 --- /dev/null +++ b/examples/vime_rocm_attention_ablation/validate_artifacts.py @@ -0,0 +1,1157 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors + +"""Validate one ROCm Vime Attention operator arm and the four-arm matrix. + +The executable matrix in this directory is the Attention implementation matrix +(``P/P``, ``P/R``, ``R/P``, and ``R/R``). It is intentionally not the compact +``A0``-``A7`` Attention root-cause manifest: those rows need separate, real +mutation hooks before they can be claimed as executed experiments. +""" + +from __future__ import annotations + +import hashlib +import json +import math +from collections.abc import Iterable, Mapping, Sequence +from pathlib import Path +from typing import Any + +import torch + +try: + from .tis_metrics import SIDECAR_SCHEMA_VERSION +except ImportError: # Direct ``python examples/.../validate_artifacts.py`` execution. + from tis_metrics import SIDECAR_SCHEMA_VERSION # type: ignore[no-redef] + +SCHEMA_VERSION = "rlkernel.vime_rocm_attention_operator_arm.v1" +MATRIX_SCHEMA_VERSION = "rlkernel.vime_rocm_attention_operator_matrix.v1" +LAUNCH_SCHEMA_VERSION = "rlkernel.vime_rocm_attention_arm_launch.v1" + +CASE_IMPLEMENTATIONS = { + "P/P": {"training": "production", "rollout": "production"}, + "P/R": {"training": "production", "rollout": "rl_kernel"}, + "R/P": {"training": "rl_kernel", "rollout": "production"}, + "R/R": {"training": "rl_kernel", "rollout": "rl_kernel"}, +} +FRAMEWORK_TARGETS = (("megatron", "training"), ("vllm", "rollout")) + +STRICT_ROCM_BACKEND_ID = "rlkernel.rocm.attention.aiter_ck_ag_rs.v1" +STRICT_ROCM_CORE_BACKEND_ID = "aiter.rocm.ck_dense_mha" +STRICT_ROCM_CORE_ID = "rlkernel.attention.rocm.aiter_ck_dense_mha.v1" +STRICT_ROCM_SCHEDULE_ID = "single_batch_aiter_ck_dense_mha_no_splitkv" +ROCM_DETERMINISTIC_PROJECTION_BACKEND_ID = "rlkernel.rocm.triton_det_gemm" +ROCM_DETERMINISTIC_COLLECTIVE_BACKEND_ID = "rocm_ipc_fixed_tree" +ROCM_FFN_BACKEND_ID = "rlkernel.rocm.det_gemm_swiglu" +STRICT_FFN_BACKEND_ID = "rlkernel.ffn.qwen3.deterministic.v1" +STRICT_LINEAR_LOGP_BACKEND_ID = "rlkernel.linear_logp.bitwise.v1" +ROCM_LOGP_KERNEL_BACKEND_ID = "rocm-vocab-parallel-logp-ws2" + +_FALLBACK_KEYS = { + "attention_fallback", + "fallback", + "fallback_used", + "split_kv_fallback", + "used_fallback", +} +_TRITON_KEYS = {"triton_used", "uses_triton"} +_REFERENCE_KEYS = {"reference_only"} + + +def _case_id(value: str) -> str: + normalized = value.strip().upper() + if normalized not in CASE_IMPLEMENTATIONS: + raise ValueError(f"unknown Attention operator case {value!r}") + return normalized + + +def _walk_key_values(value: Any) -> Iterable[tuple[str, Any]]: + if isinstance(value, Mapping): + for key, item in value.items(): + normalized = str(key).strip().lower() + yield normalized, item + yield from _walk_key_values(item) + elif isinstance(value, (list, tuple)): + for item in value: + yield from _walk_key_values(item) + + +def _values_for_keys(value: Any, keys: set[str]) -> list[Any]: + return [item for key, item in _walk_key_values(value) if key in keys] + + +def _truthy_flag(value: Any, keys: set[str]) -> bool: + return any(item not in (False, None, "", 0) for item in _values_for_keys(value, keys)) + + +def _contains_string(value: Any, needle: str) -> bool: + normalized = needle.lower() + if isinstance(value, str): + return normalized in value.lower() + if isinstance(value, Mapping): + return any(_contains_string(item, normalized) for item in value.values()) + if isinstance(value, (list, tuple)): + return any(_contains_string(item, normalized) for item in value) + return False + + +def _runtime_platform(provenance: Any) -> str | None: + values = _values_for_keys(provenance, {"runtime_platform", "platform"}) + normalized = { + str(value).strip().lower() + for value in values + if isinstance(value, str) and value.strip() + } + if normalized & {"rocm", "hip"}: + return "rocm" + if "cuda" in normalized: + return "cuda" + return None + + +def _has_exact_value(value: Any, keys: set[str], expected: str) -> bool: + return any(str(item).strip() == expected for item in _values_for_keys(value, keys)) + + +def load_readbacks(directory: Path) -> list[dict[str, Any]]: + """Load framework readbacks emitted into one arm-local directory.""" + + values: list[dict[str, Any]] = [] + for path in sorted(directory.glob("*.json")): + value = json.loads(path.read_text(encoding="utf-8")) + if not isinstance(value, dict): + raise ValueError(f"readback must contain an object: {path}") + value = dict(value) + value["_path"] = str(path) + values.append(value) + if not values: + raise ValueError(f"no framework readbacks found in {directory}") + return values + + +def validate_launch_manifest(path: Path, case_id: str) -> dict[str, Any]: + """Validate the arm-local frozen configuration emitted before execution.""" + + normalized_case = _case_id(case_id) + errors: list[str] = [] + try: + value = json.loads(path.read_text(encoding="utf-8")) + except Exception as exc: + return { + "passed": False, + "errors": [f"cannot load launch manifest {path}: {exc}"], + "frozen_input_fingerprint": None, + } + if not isinstance(value, Mapping): + return { + "passed": False, + "errors": [f"launch manifest must contain an object: {path}"], + "frozen_input_fingerprint": None, + } + if value.get("schema_version") != LAUNCH_SCHEMA_VERSION: + errors.append(f"launch manifest does not use {LAUNCH_SCHEMA_VERSION}") + if value.get("case_id") != normalized_case: + errors.append("launch manifest carries the wrong case_id") + if value.get("expected_implementations") != CASE_IMPLEMENTATIONS[normalized_case]: + errors.append("launch manifest carries the wrong implementation mapping") + fingerprint = value.get("frozen_input_fingerprint") + if not isinstance(fingerprint, str) or not fingerprint: + errors.append("launch manifest lacks a frozen input fingerprint") + + environment = value.get("environment") + if not isinstance(environment, Mapping): + errors.append("launch manifest lacks its public environment") + environment = {} + expected_environment = { + "RL_KERNEL_ATTENTION_CASE": normalized_case, + "RL_KERNEL_FFN_CASE": "R/R", + "RL_KERNEL_LOGP_CASE": "R/R", + "RL_KERNEL_VLLM_REAL_VOCAB_SIZE": "151936", + "RL_KERNEL_VLLM_PADDED_VOCAB_SIZE": "152064", + "RL_KERNEL_VLLM_INTEGRATION": "1", + "VLLM_ATTENTION_BACKEND": "ROCM_AITER_FA", + "VLLM_ROCM_USE_AITER": "1", + "VLLM_ROCM_SHUFFLE_KV_CACHE_LAYOUT": "0", + } + for name, expected in expected_environment.items(): + if environment.get(name) != expected: + errors.append(f"launch environment {name} is not frozen to {expected!r}") + hip_visible = environment.get("HIP_VISIBLE_DEVICES") + if not isinstance(hip_visible, str) or not hip_visible: + errors.append("launch environment lacks HIP_VISIBLE_DEVICES") + if environment.get("CUDA_VISIBLE_DEVICES") != hip_visible: + errors.append("HIP_VISIBLE_DEVICES and CUDA_VISIBLE_DEVICES differ") + sidecar_directory = environment.get("RL_KERNEL_MISMATCH_SIDECAR_DIR") + if not isinstance(sidecar_directory, str) or not sidecar_directory: + errors.append("launch environment lacks RL_KERNEL_MISMATCH_SIDECAR_DIR") + for name in ("RLK_ABLATION_TP_SIZE", "RLK_ABLATION_CP_SIZE"): + try: + size = int(environment.get(name)) + except (TypeError, ValueError): + size = 0 + if size <= 0: + errors.append(f"launch environment {name} is not a positive integer") + return { + "passed": not errors, + "errors": errors, + "frozen_input_fingerprint": fingerprint, + "environment": dict(environment), + } + + +def _readback_plan_error(readback: Mapping[str, Any], case_id: str) -> str | None: + plan = readback.get("plan") + cases = plan.get("cases") if isinstance(plan, Mapping) else None + if not isinstance(cases, Mapping): + return "readback does not contain an integration plan" + attention = cases.get("attention") + ffn = cases.get("ffn") + logp = cases.get("logp") + if not isinstance(attention, Mapping) or attention.get("case_id") != case_id: + return f"readback Attention plan is not {case_id}" + for module, item in (("ffn", ffn), ("logp", logp)): + if not isinstance(item, Mapping) or item.get("case_id") != "R/R": + return f"readback changed frozen {module} case away from R/R" + return None + + +def _validate_rlkernel_record( + record: Mapping[str, Any], + *, + label: str, + framework: str, + errors: list[str], +) -> None: + backend_id = str(record.get("backend_id", "")) + provenance = record.get("provenance") + if not backend_id.startswith("rlkernel."): + errors.append(f"{label} selected RL-Kernel but reported backend {backend_id!r}") + if _runtime_platform(provenance) != "rocm": + errors.append(f"{label} RL-Kernel route did not prove ROCm execution") + if _truthy_flag(provenance, _FALLBACK_KEYS) or _truthy_flag( + provenance, _REFERENCE_KEYS + ): + errors.append(f"{label} RL-Kernel route reported a fallback/reference path") + fallback_values = _values_for_keys(provenance, {"fallback"}) + reference_values = _values_for_keys(provenance, _REFERENCE_KEYS) + if not fallback_values or not any(value is False for value in fallback_values): + errors.append(f"{label} did not explicitly prove fallback=false") + if not reference_values or not any(value is False for value in reference_values): + errors.append(f"{label} did not explicitly prove reference_only=false") + if not _has_exact_value( + provenance, + {"actual_backend", "backend_id"}, + STRICT_ROCM_BACKEND_ID, + ): + errors.append( + f"{label} did not prove strict ROCm backend {STRICT_ROCM_BACKEND_ID!r}" + ) + if not _has_exact_value( + provenance, + {"strict_core_id", "core_id"}, + STRICT_ROCM_CORE_ID, + ): + errors.append(f"{label} did not prove strict ROCm core {STRICT_ROCM_CORE_ID!r}") + if not _has_exact_value( + provenance, + {"strict_schedule", "schedule_id"}, + STRICT_ROCM_SCHEDULE_ID, + ): + errors.append( + f"{label} did not prove strict ROCm schedule {STRICT_ROCM_SCHEDULE_ID!r}" + ) + production_ready = _values_for_keys(provenance, {"production_ready"}) + if not production_ready or not any(value is True for value in production_ready): + errors.append(f"{label} strict ROCm provenance is not production-ready") + native_arithmetic = _values_for_keys(provenance, {"native_attention_arithmetic"}) + if not native_arithmetic or not any(value is True for value in native_arithmetic): + errors.append(f"{label} did not prove native Attention arithmetic") + if not _has_exact_value(provenance, {"split_kv"}, "disabled"): + errors.append(f"{label} did not prove split_kv=disabled") + core_backends = _values_for_keys(provenance, {"core_actual_backends"}) + if not any( + isinstance(value, (list, tuple)) and STRICT_ROCM_CORE_BACKEND_ID in value + for value in core_backends + ): + errors.append( + f"{label} did not prove AITER/CK core backend {STRICT_ROCM_CORE_BACKEND_ID!r}" + ) + + projections = _values_for_keys(provenance, {"deterministic_projection"}) + projection_ok = any( + isinstance(value, Mapping) + and value.get("backend_id") == ROCM_DETERMINISTIC_PROJECTION_BACKEND_ID + and value.get("deterministic") is True + and value.get("split_k") is False + and value.get("accumulation_dtype") == "fp32" + and value.get("reduction_order") == "k_ascending" + and value.get("triton_used") is True + and isinstance(value.get("roles"), (list, tuple)) + and set(value["roles"]) == {"qkv", "o_proj"} + for value in projections + ) + if not projection_ok: + errors.append( + f"{label} did not prove deterministic ROCm QKV/O projection " + f"{ROCM_DETERMINISTIC_PROJECTION_BACKEND_ID!r}" + ) + + # Triton is permitted only for the declared deterministic QKV/O projection. + # The Attention core itself must remain the native AITER/CK implementation. + for key, item in _walk_key_values(provenance): + if key in _TRITON_KEYS and item is True: + continue + if isinstance(item, str) and "triton" in item.lower(): + if item != ROCM_DETERMINISTIC_PROJECTION_BACKEND_ID: + errors.append(f"{label} reported an unapproved Triton backend {item!r}") + + if framework == "vllm": + layouts = _values_for_keys(provenance, {"framework_layout"}) + if "vllm_paged_kv" not in layouts: + errors.append(f"{label} did not prove the vLLM paged-KV execution boundary") + tp_values = _values_for_keys(provenance, {"tp_world_size"}) + try: + tp_world_size = max(int(value) for value in tp_values) + except (TypeError, ValueError): + tp_world_size = 0 + collective_values = _values_for_keys( + provenance, {"deterministic_all_reduce_backend"} + ) + expected_collective = ( + "none" if tp_world_size == 1 else ROCM_DETERMINISTIC_COLLECTIVE_BACKEND_ID + ) + if tp_world_size <= 0 or expected_collective not in collective_values: + errors.append( + f"{label} did not prove the deterministic ROCm O-projection collective" + ) + elif framework == "megatron": + cp_values = _values_for_keys(provenance, {"cp_world_size"}) + try: + cp_world_size = max(int(value) for value in cp_values) + except (TypeError, ValueError): + cp_world_size = 0 + communication = _values_for_keys(provenance, {"communication_backend"}) + expected_communication = "none" if cp_world_size == 1 else "rccl_ag_rs" + if cp_world_size <= 0 or expected_communication not in communication: + errors.append(f"{label} did not prove the expected ROCm CP communication path") + tp_values = _values_for_keys(provenance, {"tp_world_size"}) + try: + tp_world_size = max(int(value) for value in tp_values) + except (TypeError, ValueError): + tp_world_size = 0 + expected_tp_collective = ( + "none" if tp_world_size == 1 else ROCM_DETERMINISTIC_COLLECTIVE_BACKEND_ID + ) + qkv_collectives = _values_for_keys( + provenance, {"tp_qkv_dgrad_collective"} + ) + output_collectives = _values_for_keys( + provenance, {"tp_output_projection_collective"} + ) + if ( + tp_world_size <= 0 + or expected_tp_collective not in qkv_collectives + or expected_tp_collective not in output_collectives + ): + errors.append( + f"{label} did not prove deterministic ROCm TP projection collectives" + ) + + +def _validate_production_record( + record: Mapping[str, Any], + *, + label: str, + errors: list[str], +) -> None: + backend_id = str(record.get("backend_id", "")) + provenance = record.get("provenance") + if backend_id.startswith("rlkernel.") or _contains_string(provenance, "rlkernel."): + errors.append(f"{label} selected production but executed an RL-Kernel backend") + if _truthy_flag(provenance, _FALLBACK_KEYS): + errors.append(f"{label} production route reported fallback") + + +def _validate_strict_dense_record( + record: Mapping[str, Any], + *, + module: str, + framework: str, + label: str, + errors: list[str], +) -> None: + provenance = record.get("provenance") + expected_backend = ( + STRICT_FFN_BACKEND_ID if module == "ffn" else STRICT_LINEAR_LOGP_BACKEND_ID + ) + if record.get("case_id") != "R/R" or record.get("implementation") != "rl_kernel": + errors.append(f"{label} did not execute the fixed R/R route") + if record.get("backend_id") != expected_backend: + errors.append(f"{label} reported backend {record.get('backend_id')!r}") + if record.get("execution_mode", "eager") != "eager": + errors.append(f"{label} did not execute in frozen eager mode") + if int(record.get("call_count", 0)) <= 0: + errors.append(f"{label} had zero executed calls") + if _runtime_platform(provenance) != "rocm": + errors.append(f"{label} did not prove ROCm execution") + if _truthy_flag(provenance, _FALLBACK_KEYS): + errors.append(f"{label} reported a fallback") + + if module == "ffn": + if not _has_exact_value(provenance, {"actual_backend"}, ROCM_FFN_BACKEND_ID): + errors.append(f"{label} did not prove the strict ROCm FFN backend") + if not _has_exact_value( + provenance, + {"deterministic_all_reduce_backend"}, + ROCM_DETERMINISTIC_COLLECTIVE_BACKEND_ID, + ): + errors.append(f"{label} did not prove the ROCm fixed-tree FFN reduction") + return + + expected_entrypoint = ( + "rocm_deterministic_linear_logp_tp" + if framework == "megatron" + else "rocm_vocab_parallel_logp_from_local_logits_tp" + ) + if not _has_exact_value( + provenance, {"logprob_kernel_backend"}, ROCM_LOGP_KERNEL_BACKEND_ID + ): + errors.append(f"{label} did not prove the ROCm WS2 logp kernel") + if not _has_exact_value(provenance, {"strict_entrypoint"}, expected_entrypoint): + errors.append(f"{label} did not prove strict entrypoint {expected_entrypoint!r}") + deterministic_values = _values_for_keys(provenance, {"deterministic_linear_logp"}) + if not deterministic_values or not any(value is True for value in deterministic_values): + errors.append(f"{label} did not prove deterministic linear logp") + + +def validate_attention_readbacks( + readbacks: Sequence[Mapping[str, Any]], + case_id: str, +) -> dict[str, Any]: + """Validate exact training/rollout routing and strict ROCm provenance.""" + + normalized_case = _case_id(case_id) + expected = CASE_IMPLEMENTATIONS[normalized_case] + errors: list[str] = [] + frameworks: dict[str, Any] = {} + + for framework, target in FRAMEWORK_TARGETS: + label = f"{framework}/{target}" + matching = [ + value + for value in readbacks + if value.get("framework") == framework and value.get("target") == target + ] + if not matching: + errors.append(f"missing {label} readback") + continue + + for value in matching: + plan_error = _readback_plan_error(value, normalized_case) + if plan_error: + errors.append(f"{label}: {plan_error}") + if value.get("fallbacks"): + errors.append(f"{label} recorded fallback: {value['fallbacks']}") + + hook_count = sum( + isinstance(value.get("installed_hooks"), Mapping) + and bool(value["installed_hooks"].get("attention")) + for value in matching + ) + if hook_count == 0: + errors.append(f"{label} Attention hook was not installed") + + records = [ + value["operators"]["attention"] + for value in matching + if isinstance(value.get("operators"), Mapping) + and isinstance(value["operators"].get("attention"), Mapping) + ] + call_count = sum(int(record.get("call_count", 0)) for record in records) + if call_count <= 0: + errors.append(f"{label} Attention had zero executed calls") + + expected_implementation = expected[target] + backend_ids: set[str] = set() + for record in records: + record_label = f"{label} Attention" + backend_ids.add(str(record.get("backend_id", ""))) + if record.get("case_id") != normalized_case: + errors.append(f"{record_label} record has the wrong case_id") + if record.get("execution_mode", "eager") != "eager": + errors.append(f"{record_label} did not execute in frozen eager mode") + if record.get("implementation") != expected_implementation: + errors.append( + f"{record_label} implementation={record.get('implementation')!r}, " + f"expected {expected_implementation!r}" + ) + continue + if expected_implementation == "rl_kernel": + _validate_rlkernel_record( + record, + label=record_label, + framework=framework, + errors=errors, + ) + else: + _validate_production_record(record, label=record_label, errors=errors) + + frameworks[label] = { + "readback_count": len(matching), + "installed_processes": hook_count, + "call_count": call_count, + "expected_implementation": expected_implementation, + "backend_ids": sorted(backend_ids), + } + + for module in ("ffn", "logp"): + dense_label = f"{label} {module.upper()}" + dense_records = [ + value["operators"][module] + for value in matching + if isinstance(value.get("operators"), Mapping) + and isinstance(value["operators"].get(module), Mapping) + ] + hook_count = sum( + isinstance(value.get("installed_hooks"), Mapping) + and bool(value["installed_hooks"].get(module)) + for value in matching + ) + if hook_count == 0: + errors.append(f"{dense_label} hook was not installed") + if not dense_records: + errors.append(f"missing {dense_label} execution record") + for record in dense_records: + _validate_strict_dense_record( + record, + module=module, + framework=framework, + label=dense_label, + errors=errors, + ) + frameworks[f"{label}/{module}"] = { + "readback_count": len(dense_records), + "installed_processes": hook_count, + "call_count": sum(int(record.get("call_count", 0)) for record in dense_records), + "expected_implementation": "rl_kernel", + "backend_ids": sorted( + {str(record.get("backend_id", "")) for record in dense_records} + ), + } + + return { + "passed": not errors, + "errors": errors, + "frameworks": frameworks, + } + + +def _tensor(value: Any, *, label: str) -> torch.Tensor: + if isinstance(value, torch.Tensor): + result = value.detach().cpu() + else: + try: + result = torch.as_tensor(value) + except Exception as exc: # pragma: no cover - message is exercised through callers + raise ValueError(f"{label} is not tensor-like") from exc + return result.reshape(-1).contiguous() + + +def _scalar_key(value: Any) -> Any: + if isinstance(value, torch.Tensor): + if value.numel() != 1: + raise ValueError("sample identity key tensors must be scalar") + return value.detach().cpu().item() + return value + + +def _positive_int(value: Any, *, label: str) -> int: + value = _scalar_key(value) + if isinstance(value, bool) or not isinstance(value, int) or value <= 0: + raise ValueError(f"{label} must be a positive integer") + return value + + +def slice_response_mask_for_cp( + loss_mask: Any, + *, + total_length: int, + response_length: int, + context_parallel_size: int, + context_parallel_rank: int, +) -> torch.Tensor: + """Apply Vime's two-ended CP zigzag response slice to a full loss mask.""" + + mask = _tensor(loss_mask, label="loss_masks") + if response_length < 0 or total_length < response_length: + raise ValueError( + f"invalid total/response lengths: {total_length}/{response_length}" + ) + if mask.numel() != response_length: + raise ValueError( + "full loss mask length does not match response_length: " + f"{mask.numel()} != {response_length}" + ) + if context_parallel_size <= 0: + raise ValueError("context_parallel_size must be positive") + if not 0 <= context_parallel_rank < context_parallel_size: + raise ValueError("context_parallel_rank is outside the CP world") + if context_parallel_size == 1: + return mask + + prompt_length = total_length - response_length + chunk_size = (total_length + 2 * context_parallel_size - 1) // ( + 2 * context_parallel_size + ) + chunks = ( + ( + context_parallel_rank * chunk_size, + (context_parallel_rank + 1) * chunk_size, + ), + ( + (2 * context_parallel_size - context_parallel_rank - 1) * chunk_size, + (2 * context_parallel_size - context_parallel_rank) * chunk_size, + ), + ) + response_chunks: list[torch.Tensor] = [] + for start, end in chunks: + logit_start = max(start, prompt_length - 1) + logit_end = min(end, total_length - 1) + if logit_start >= logit_end: + continue + response_start = logit_start - (prompt_length - 1) + response_end = logit_end - (prompt_length - 1) + response_chunks.append(mask[response_start:response_end]) + if not response_chunks: + return mask.new_empty((0,)) + return torch.cat(response_chunks, dim=0) + + +def _hash_value(digest: Any, value: Any) -> None: + if value is None: + digest.update(b"N") + elif isinstance(value, bool): + digest.update(b"B1" if value else b"B0") + elif isinstance(value, int): + digest.update(f"I{value};".encode()) + elif isinstance(value, float): + digest.update(f"F{value.hex()};".encode()) + elif isinstance(value, str): + encoded = value.encode("utf-8") + digest.update(f"S{len(encoded)}:".encode()) + digest.update(encoded) + elif isinstance(value, torch.Tensor): + tensor = value.detach().cpu().contiguous() + digest.update(f"T{tensor.dtype}:{tuple(tensor.shape)}:".encode()) + digest.update(tensor.view(torch.uint8).numpy().tobytes()) + elif isinstance(value, Mapping): + digest.update(b"{") + for key in sorted(value, key=lambda item: str(item)): + _hash_value(digest, str(key)) + _hash_value(digest, value[key]) + digest.update(b"}") + elif isinstance(value, (list, tuple)): + digest.update(b"[") + for item in value: + _hash_value(digest, item) + digest.update(b"]") + else: + _hash_value(digest, str(value)) + + +def _value_fingerprint(value: Any) -> str: + digest = hashlib.sha256() + _hash_value(digest, value) + return digest.hexdigest() + + +def load_rollout_identity(directory: Path) -> dict[str, Any]: + """Hash the frozen sample/token identity emitted by Vime rollout dumps.""" + + errors: list[str] = [] + identities: list[dict[str, Any]] = [] + paths = sorted(directory.glob("*.pt")) + for path in paths: + try: + payload = torch.load(path, map_location="cpu", weights_only=False) + samples = payload.get("samples") if isinstance(payload, Mapping) else None + if not isinstance(samples, list): + raise ValueError("payload does not contain a samples list") + outer_rollout_id = payload.get("rollout_id") + for ordinal, sample in enumerate(samples): + if not isinstance(sample, Mapping): + raise ValueError(f"sample {ordinal} is not a mapping") + tokens = _tensor(sample.get("tokens"), label="tokens") + response_length = sample.get("response_length") + if isinstance(response_length, bool) or not isinstance(response_length, int): + raise ValueError(f"sample {ordinal} has an invalid response_length") + identity = { + "dump_file": path.name, + "ordinal": ordinal, + "outer_rollout_id": outer_rollout_id, + "rollout_id": sample.get("rollout_id"), + "group_index": sample.get("group_index"), + "index": sample.get("index"), + "prompt": sample.get("prompt"), + "tokens": tokens, + "response_length": response_length, + "loss_mask": sample.get("loss_mask"), + } + identities.append( + { + "sort_key": ( + path.name, + ordinal, + ), + "identity": identity, + } + ) + except Exception as exc: + errors.append(f"{path}: {exc}") + + identities.sort(key=lambda item: item["sort_key"]) + digest = hashlib.sha256() + token_count = 0 + for item in identities: + identity = item["identity"] + token_count += int(identity["tokens"].numel()) + _hash_value(digest, identity) + if not paths: + errors.append(f"no rollout dumps found in {directory}") + if not identities: + errors.append("no rollout sample identity was found") + return { + "passed": not errors, + "errors": errors, + "fingerprint": digest.hexdigest() if identities else None, + "sample_count": len(identities), + "token_count": token_count, + "artifacts": [str(path) for path in paths], + } + + +def _sidecar_samples( + path: Path, + *, + tensor_parallel_size: int, + context_parallel_size: int, +) -> list[dict[str, Any]]: + payload = torch.load(path, map_location="cpu", weights_only=False) + if not isinstance(payload, Mapping): + raise ValueError("mismatch sidecar must contain a mapping") + if payload.get("schema_version") != SIDECAR_SCHEMA_VERSION: + raise ValueError(f"mismatch sidecar does not use {SIDECAR_SCHEMA_VERSION}") + if payload.get("tensor_parallel_size") != tensor_parallel_size: + raise ValueError("mismatch sidecar tensor_parallel_size differs from launch") + if payload.get("context_parallel_size") != context_parallel_size: + raise ValueError("mismatch sidecar context_parallel_size differs from launch") + + training_values = payload.get("train_log_probs") + rollout_values = payload.get("rollout_log_probs") + loss_masks = payload.get("loss_masks") + total_lengths = payload.get("total_lengths") + response_lengths = payload.get("response_lengths") + required_lists = { + "train_log_probs": training_values, + "rollout_log_probs": rollout_values, + "loss_masks": loss_masks, + "total_lengths": total_lengths, + "response_lengths": response_lengths, + } + if any(not isinstance(value, (list, tuple)) for value in required_lists.values()): + missing = [ + name + for name, value in required_lists.items() + if not isinstance(value, (list, tuple)) + ] + raise ValueError(f"mismatch sidecar lacks list fields: {', '.join(missing)}") + assert isinstance(training_values, (list, tuple)) + assert isinstance(rollout_values, (list, tuple)) + assert isinstance(loss_masks, (list, tuple)) + assert isinstance(total_lengths, (list, tuple)) + assert isinstance(response_lengths, (list, tuple)) + count = len(training_values) + if not all( + len(value) == count + for value in (rollout_values, loss_masks, total_lengths, response_lengths) + ): + raise ValueError("mismatch sidecar sample lists have different lengths") + + rank = _scalar_key(payload.get("rank")) + if isinstance(rank, bool) or not isinstance(rank, int) or rank < 0: + raise ValueError("mismatch sidecar rank must be a non-negative integer") + call_index = _scalar_key(payload.get("call_index")) + if isinstance(call_index, bool) or not isinstance(call_index, int) or call_index < 0: + raise ValueError("mismatch sidecar call_index must be a non-negative integer") + # The launcher fixes Megatron's parallel order to its default + # tp-cp-ep-dp-pp, with EP=PP=1. TP ranks are replicas for these values; + # CP ranks own distinct zigzag response shards and must remain distinct. + context_parallel_rank = (rank // tensor_parallel_size) % context_parallel_size + data_parallel_rank = rank // (tensor_parallel_size * context_parallel_size) + + values: list[dict[str, Any]] = [] + for index in range(count): + logical_key = ( + data_parallel_rank, + call_index, + index, + ) + key = (*logical_key, context_parallel_rank) + training = _tensor(training_values[index], label="train_log_probs") + rollout = _tensor(rollout_values[index], label="rollout_log_probs") + mask = _tensor(loss_masks[index], label="loss_masks").to(torch.bool) + total_length = _positive_int( + total_lengths[index], label="mismatch sidecar total_lengths" + ) + response_length = _positive_int( + response_lengths[index], label="mismatch sidecar response_lengths" + ) + mask = slice_response_mask_for_cp( + mask, + total_length=total_length, + response_length=response_length, + context_parallel_size=context_parallel_size, + context_parallel_rank=context_parallel_rank, + ).to(torch.bool) + values.append( + { + "key": key, + "logical_key": logical_key, + "training": training, + "rollout": rollout, + "mask": mask, + "total_length": total_length, + "response_length": response_length, + } + ) + return values + + +def compare_train_rollout_logps( + directory: Path, + *, + require_exact: bool, + tensor_parallel_size: int = 1, + context_parallel_size: int = 1, +) -> dict[str, Any]: + """Compute selected-token metrics from the custom hook's rank/call sidecars.""" + + if tensor_parallel_size <= 0 or context_parallel_size <= 0: + raise ValueError("tensor/context parallel sizes must be positive") + errors: list[str] = [] + unique: dict[tuple[Any, ...], dict[str, Any]] = {} + paths = sorted(directory.glob("*.pt")) + for path in paths: + try: + for sample in _sidecar_samples( + path, + tensor_parallel_size=tensor_parallel_size, + context_parallel_size=context_parallel_size, + ): + key = sample["key"] + fingerprint = _value_fingerprint( + { + "training": sample["training"], + "rollout": sample["rollout"], + "mask": sample["mask"], + "total_length": sample["total_length"], + "response_length": sample["response_length"], + } + ) + previous = unique.get(key) + if previous is not None: + if previous["fingerprint"] != fingerprint: + errors.append(f"replicated model-parallel sample {key!r} is inconsistent") + continue + unique[key] = {**sample, "fingerprint": fingerprint} + except Exception as exc: + errors.append(f"{path}: {exc}") + + if context_parallel_size > 1: + shards_by_sample: dict[tuple[Any, ...], set[int]] = {} + lengths_by_sample: dict[tuple[Any, ...], set[tuple[int, int]]] = {} + for key, sample in unique.items(): + shards_by_sample.setdefault(sample["logical_key"], set()).add(int(key[-1])) + lengths_by_sample.setdefault(sample["logical_key"], set()).add( + (sample["total_length"], sample["response_length"]) + ) + expected_shards = set(range(context_parallel_size)) + for logical_key, actual_shards in shards_by_sample.items(): + if actual_shards != expected_shards: + errors.append( + f"sample {logical_key!r} is missing context-parallel shards: " + f"expected {sorted(expected_shards)}, got {sorted(actual_shards)}" + ) + if len(lengths_by_sample[logical_key]) != 1: + errors.append( + f"sample {logical_key!r} has inconsistent lengths across " + "context-parallel shards" + ) + + mismatch_count = 0 + element_count = 0 + sum_abs_diff = 0.0 + sum_mismatch_kl = 0.0 + sum_mismatch_k3_kl = 0.0 + max_abs_diff = 0.0 + for key, sample in unique.items(): + training = sample["training"] + rollout = sample["rollout"] + mask = sample["mask"] + if training.shape != rollout.shape: + errors.append( + f"sample {key!r} train/rollout shape mismatch: " + f"{tuple(training.shape)} != {tuple(rollout.shape)}" + ) + continue + if training.dtype != rollout.dtype: + errors.append( + f"sample {key!r} train/rollout dtype mismatch: " + f"{training.dtype} != {rollout.dtype}" + ) + continue + if mask.numel() != training.numel(): + errors.append( + f"sample {key!r} mask/logprob length mismatch: " + f"{mask.numel()} != {training.numel()}" + ) + continue + active_training = training[mask] + active_rollout = rollout[mask] + if active_training.numel() == 0: + continue + if not bool(torch.isfinite(active_training).all()) or not bool( + torch.isfinite(active_rollout).all() + ): + errors.append(f"sample {key!r} contains non-finite log probabilities") + continue + mismatch_count += int(torch.ne(active_training, active_rollout).sum().item()) + delta = active_training.to(torch.float64) - active_rollout.to(torch.float64) + absolute = delta.abs() + k3 = torch.exp(delta) - delta - 1.0 + if not bool(torch.isfinite(k3).all()): + errors.append(f"sample {key!r} produced a non-finite mismatch_k3_kl") + continue + element_count += int(delta.numel()) + sum_abs_diff += float(absolute.sum().item()) + sum_mismatch_kl += float((-delta).sum().item()) + sum_mismatch_k3_kl += float(k3.sum().item()) + max_abs_diff = max(max_abs_diff, float(absolute.max().item())) + + if not paths: + errors.append(f"no mismatch sidecars found in {directory}") + if element_count == 0: + errors.append("no active train/rollout logprob elements were found") + mean_abs_diff = sum_abs_diff / element_count if element_count else None + mismatch_kl = sum_mismatch_kl / element_count if element_count else None + mismatch_k3_kl = sum_mismatch_k3_kl / element_count if element_count else None + exact = element_count > 0 and mismatch_count == 0 and max_abs_diff == 0.0 + if require_exact and not exact: + errors.append("R/R requires bitwise-equal training and rollout log probabilities") + return { + "passed": not errors, + "errors": errors, + "require_exact": require_exact, + "torch_equal": exact, + "mismatch_count": mismatch_count, + "max_abs_diff": max_abs_diff if element_count else None, + "train_rollout_logprob_abs_diff": mean_abs_diff, + "mismatch_kl": mismatch_kl, + "mismatch_k3_kl": mismatch_k3_kl, + "sample_count": len( + {sample["logical_key"] for sample in unique.values()} + ), + "element_count": element_count, + "artifacts": [str(path) for path in paths], + } + + +def validate_arm( + arm_dir: Path, + case_id: str, + *, + launcher_returncode: int = 0, +) -> dict[str, Any]: + """Validate one completed Vime arm without accepting configured-only evidence.""" + + normalized_case = _case_id(case_id) + errors: list[str] = [] + launch = validate_launch_manifest(arm_dir / "launch.json", normalized_case) + try: + readbacks = validate_attention_readbacks( + load_readbacks(arm_dir / "readbacks"), + normalized_case, + ) + except Exception as exc: + readbacks = {"passed": False, "errors": [str(exc)], "frameworks": {}} + rollout_identity = load_rollout_identity(arm_dir / "dump" / "rollout_data") + environment = launch.get("environment", {}) + try: + tensor_parallel_size = int(environment.get("RLK_ABLATION_TP_SIZE", 1)) + context_parallel_size = int(environment.get("RLK_ABLATION_CP_SIZE", 1)) + except (TypeError, ValueError): + tensor_parallel_size = 0 + context_parallel_size = 0 + try: + metrics = compare_train_rollout_logps( + arm_dir / "mismatch_sidecars", + require_exact=normalized_case == "R/R", + tensor_parallel_size=tensor_parallel_size, + context_parallel_size=context_parallel_size, + ) + except Exception as exc: + metrics = { + "passed": False, + "errors": [f"cannot validate mismatch sidecars: {exc}"], + } + if launcher_returncode != 0: + errors.append(f"Vime launcher exited with status {launcher_returncode}") + errors.extend(launch["errors"]) + errors.extend(readbacks["errors"]) + errors.extend(rollout_identity["errors"]) + errors.extend(metrics["errors"]) + return { + "schema_version": SCHEMA_VERSION, + "case_id": normalized_case, + "passed": not errors, + "launcher_returncode": launcher_returncode, + "errors": errors, + "expected_implementations": CASE_IMPLEMENTATIONS[normalized_case], + "launch": launch, + "attention_readbacks": readbacks, + "rollout_identity": rollout_identity, + "metrics": metrics, + "claim_boundary": { + "matrix_kind": "attention_operator_implementation_cross_config", + "executed_case": normalized_case, + "a0_a7_mutation_matrix_executed": False, + }, + } + + +def validate_matrix( + arm_reports: Mapping[str, Mapping[str, Any]], + *, + frozen_before: Mapping[str, Any], + frozen_after: Mapping[str, Any], +) -> dict[str, Any]: + """Validate all four arms, frozen sources, and cross-arm sample identity.""" + + errors: list[str] = [] + normalized_reports = {_case_id(case): report for case, report in arm_reports.items()} + if set(normalized_reports) != set(CASE_IMPLEMENTATIONS): + missing = sorted(set(CASE_IMPLEMENTATIONS) - set(normalized_reports)) + extra = sorted(set(normalized_reports) - set(CASE_IMPLEMENTATIONS)) + errors.append( + f"matrix cases differ from the four-arm contract; missing={missing}, extra={extra}" + ) + + before_fingerprint = frozen_before.get("fingerprint") + after_fingerprint = frozen_after.get("fingerprint") + frozen_sources_match = ( + isinstance(before_fingerprint, str) + and before_fingerprint + and before_fingerprint == after_fingerprint + ) + if not frozen_sources_match: + errors.append("frozen model/checkpoint/data/revision fingerprint changed during the matrix") + + rollout_fingerprints: dict[str, str | None] = {} + metrics: dict[str, Any] = {} + for case_id, report in sorted(normalized_reports.items()): + if report.get("case_id") != case_id: + errors.append(f"{case_id} report carries the wrong case_id") + if report.get("passed") is not True: + errors.append(f"{case_id} arm did not pass strict validation") + launch = report.get("launch") + arm_fingerprint = ( + launch.get("frozen_input_fingerprint") if isinstance(launch, Mapping) else None + ) + if arm_fingerprint != before_fingerprint: + errors.append(f"{case_id} arm was not launched with the sealed frozen inputs") + identity = report.get("rollout_identity") + rollout_fingerprints[case_id] = ( + identity.get("fingerprint") if isinstance(identity, Mapping) else None + ) + arm_metrics = report.get("metrics") + if isinstance(arm_metrics, Mapping): + metrics[case_id] = { + key: arm_metrics.get(key) + for key in ( + "mismatch_count", + "max_abs_diff", + "train_rollout_logprob_abs_diff", + "mismatch_kl", + "mismatch_k3_kl", + "sample_count", + "element_count", + ) + } + + identity_values = list(rollout_fingerprints.values()) + rollout_identity_match = ( + len(identity_values) == len(CASE_IMPLEMENTATIONS) + and all(isinstance(value, str) and value for value in identity_values) + and len(set(identity_values)) == 1 + ) + if not rollout_identity_match: + errors.append("rollout sample/token identity is not frozen across all four arms") + + return { + "schema_version": MATRIX_SCHEMA_VERSION, + "passed": not errors, + "errors": errors, + "cases": list(CASE_IMPLEMENTATIONS), + "frozen_sources": { + "matched": frozen_sources_match, + "before": before_fingerprint, + "after": after_fingerprint, + }, + "rollout_identity": { + "matched": rollout_identity_match, + "fingerprints": rollout_fingerprints, + }, + "metrics": metrics, + "claim_boundary": { + "matrix_kind": "attention_operator_implementation_cross_config", + "implemented_cases": list(CASE_IMPLEMENTATIONS), + "a0_a7_mutation_matrix_executed": False, + "note": ( + "A0-A7 remain a diagnostic taxonomy until each row has a real " + "runtime mutation and restoration hook" + ), + }, + } + + +def write_report(path: Path, report: Mapping[str, Any]) -> None: + path.parent.mkdir(parents=True, exist_ok=True) + temporary = path.with_suffix(path.suffix + ".tmp") + temporary.write_text(json.dumps(report, indent=2, sort_keys=True) + "\n", encoding="utf-8") + temporary.replace(path) + + +def main(argv: list[str] | None = None) -> int: + import argparse + + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--arm-dir", type=Path, required=True) + parser.add_argument("--case", choices=tuple(CASE_IMPLEMENTATIONS), required=True) + parser.add_argument("--launcher-returncode", type=int, default=0) + parser.add_argument("--output", type=Path, default=None) + args = parser.parse_args(argv) + + report = validate_arm( + args.arm_dir, + args.case, + launcher_returncode=args.launcher_returncode, + ) + output = args.output or args.arm_dir / "validation.json" + write_report(output, report) + print(json.dumps(report, indent=2, sort_keys=True)) + return 0 if report["passed"] else 1 + + +if __name__ == "__main__": + raise SystemExit(main()) diff --git a/rl_engine/_C.pyi b/rl_engine/_C.pyi index a50cadd1..4e60e52b 100644 --- a/rl_engine/_C.pyi +++ b/rl_engine/_C.pyi @@ -14,6 +14,10 @@ def deterministic_collective_create( def deterministic_collective_destroy(handle: int) -> None: ... def deterministic_collective_stage(handle: int, input: torch.Tensor) -> None: ... def deterministic_collective_all_reduce(handle: int, output: torch.Tensor) -> None: ... +def deterministic_collective_prepare_staged(handle: int, input: torch.Tensor) -> None: ... +def deterministic_collective_all_reduce_staged( + handle: int, input: torch.Tensor, output: torch.Tensor +) -> None: ... def deterministic_collective_all_reduce_fused( handle: int, input: torch.Tensor, output: torch.Tensor ) -> None: ... @@ -22,6 +26,9 @@ def deterministic_collective_all_gather(handle: int, output: torch.Tensor) -> No def deterministic_collective_all_gather_fused( handle: int, input: torch.Tensor, output: torch.Tensor ) -> None: ... +def deterministic_collective_all_gather_many( + handle: int, inputs: list[torch.Tensor], outputs: list[torch.Tensor] +) -> None: ... def deterministic_collective_rocm_all_reduce( rank_inputs: torch.Tensor, output: torch.Tensor, diff --git a/rl_engine/distributed/collectives.py b/rl_engine/distributed/collectives.py index 673828e4..fe0f6eee 100644 --- a/rl_engine/distributed/collectives.py +++ b/rl_engine/distributed/collectives.py @@ -1,15 +1,11 @@ # SPDX-License-Identifier: Apache-2.0 # Copyright (c) 2026 RL-Kernel Contributors -"""Deterministic collectives for CUDA IPC and ROCm rank-ordered transport. - -ROCm uses HIP IPC where it wins and RCCL otherwise. Reduction arithmetic stays -outside RCCL and follows the same fixed balanced rank tree on every rank. -""" from __future__ import annotations import socket import threading +from collections.abc import Iterable from types import TracebackType from typing import Any @@ -18,19 +14,15 @@ _SUPPORTED_WORLD_SIZES = (1, 2, 4, 8) _DEFAULT_MAX_SIZE_BYTES = 64 * 1024 * 1024 -# Packing two independent lanes saves a collective launch for small tensors, -# but doubles the message size seen by RCCL. On MI300X, separate AllGather -# transports win once the packed payload reaches the multi-megabyte regime. -# Keep the crossover explicit and easy to retune with new RCCL releases. -_PACKED_REDUCE_SCATTER_MAX_BYTES = 8 * 1024 * 1024 -_ROCM_IPC_DIRECT_ALL_REDUCE_MAX_BYTES = 768 * 1024 -_ROCM_IPC_SHARDED_ALL_REDUCE_MIN_BYTES = 2176 * 1024 -_ROCM_IPC_ALL_GATHER_MAX_BYTES = 256 * 1024 _COLLECTIVE_STAGING_FRAMES = 3 -_COLLECTIVE_FRAME_METADATA_BYTES = 3 * 8 +_COLLECTIVE_FRAME_METADATA_BYTES = 4 * 8 +_DIRECT_STAGING_MAX_BYTES = 256 * 1024 _REDUCTION_DTYPES = (torch.float32, torch.float16, torch.bfloat16) +_REDUCTION_DTYPE_BYTES = {torch.float32: 4, torch.float16: 2, torch.bfloat16: 2} _COLLECTIVES: dict[tuple[int, int, int, int], Any] = {} DETERMINISTIC_ALL_REDUCE_OP = "rl_kernel::deterministic_all_reduce_" +DETERMINISTIC_STAGING_RESERVE_OP = "rl_kernel::deterministic_staging_reserve_" +DETERMINISTIC_STAGED_ALL_REDUCE_OP = "rl_kernel::deterministic_staged_all_reduce" @torch.library.custom_op(DETERMINISTIC_ALL_REDUCE_OP, mutates_args={"input"}) @@ -61,6 +53,63 @@ def deterministic_all_reduce_inplace( return input +@torch.library.custom_op(DETERMINISTIC_STAGING_RESERVE_OP, mutates_args={"staging"}) +def _deterministic_staging_reserve_(staging: torch.Tensor, collective_handle: int) -> None: + """Wait until every peer has finished reading the shared direct-output slot.""" + + from rl_engine import _C + + _C.deterministic_collective_prepare_staged(collective_handle, staging) + + +@_deterministic_staging_reserve_.register_fake +def _deterministic_staging_reserve_fake( + staging: torch.Tensor, + collective_handle: int, +) -> None: + del staging, collective_handle + + +@torch.library.custom_op(DETERMINISTIC_STAGED_ALL_REDUCE_OP, mutates_args=()) +def _deterministic_staged_all_reduce( + staging: torch.Tensor, + collective_handle: int, +) -> torch.Tensor: + """Reduce a GEMM result already resident in the local CUDA IPC payload.""" + + from rl_engine import _C + + output = torch.empty_like(staging) + _C.deterministic_collective_all_reduce_staged(collective_handle, staging, output) + return output + + +@_deterministic_staged_all_reduce.register_fake +def _deterministic_staged_all_reduce_fake( + staging: torch.Tensor, + collective_handle: int, +) -> torch.Tensor: + del collective_handle + return torch.empty_like(staging) + + +def deterministic_staging_reserve( + staging: torch.Tensor, + *, + collective_handle: int, +) -> torch.Tensor: + _deterministic_staging_reserve_(staging, collective_handle) + return staging + + +def deterministic_all_reduce_staged( + staging: torch.Tensor, + *, + collective_handle: int, +) -> torch.Tensor: + return _deterministic_staged_all_reduce(staging, collective_handle) + + class DeterministicCollective: """Correctness-first TP-invariant CUDA collectives for one eight-GPU node. @@ -128,10 +177,11 @@ def __init__( "deterministic_collective_destroy", "deterministic_collective_stage", "deterministic_collective_all_reduce", - "deterministic_collective_all_reduce_fused", + "deterministic_collective_prepare_staged", + "deterministic_collective_all_reduce_staged", "deterministic_collective_reduce_scatter", "deterministic_collective_all_gather", - "deterministic_collective_all_gather_fused", + "deterministic_collective_all_gather_many", ) missing = [name for name in required_symbols if not hasattr(_C, name)] if missing: @@ -146,6 +196,7 @@ def __init__( self._lock = threading.Lock() self._handle = 0 self._validated_signatures: set[tuple[Any, ...]] = set() + self._direct_staging_views: dict[tuple[tuple[int, ...], torch.dtype], torch.Tensor] = {} self._staging = torch.zeros( _COLLECTIVE_STAGING_FRAMES * (self.max_size_bytes + _COLLECTIVE_FRAME_METADATA_BYTES), dtype=torch.uint8, @@ -181,6 +232,44 @@ def __init__( ) self._synchronize_ranks() + def prepare_direct_staging_views( + self, + shapes: Iterable[tuple[int, ...]], + *, + dtype: torch.dtype, + ) -> None: + """Materialize stable, graph-capturable views of the local IPC payload.""" + + if dtype not in _REDUCTION_DTYPE_BYTES: + raise TypeError(f"unsupported direct-staging dtype {dtype}") + element_size = _REDUCTION_DTYPE_BYTES[dtype] + for raw_shape in shapes: + shape = tuple(int(dim) for dim in raw_shape) + numel = 1 + for dim in shape: + if dim < 0: + raise ValueError(f"direct-staging dimensions must be non-negative, got {shape}") + numel *= dim + size_bytes = numel * element_size + if size_bytes > min(self.max_size_bytes, _DIRECT_STAGING_MAX_BYTES): + continue + byte_view = self._staging.narrow( + 0, + _COLLECTIVE_FRAME_METADATA_BYTES, + size_bytes, + ) + self._direct_staging_views[(shape, dtype)] = byte_view.view(dtype).view(shape) + + def direct_staging_view( + self, + shape: tuple[int, ...], + *, + dtype: torch.dtype, + ) -> torch.Tensor | None: + """Return a pre-bound direct-output view, or ``None`` for an uncaptured shape.""" + + return self._direct_staging_views.get((tuple(int(dim) for dim in shape), dtype)) + def all_reduce( self, input: torch.Tensor, @@ -239,13 +328,27 @@ def all_gather_many( *, validate_signature: bool = True, ) -> tuple[torch.Tensor, ...]: - """Gather several tensors through the available single-tensor ABI.""" + """Gather several tensors with one staging handshake and no packing.""" if not inputs: raise ValueError("all_gather_many requires at least one input") - return tuple( - self.all_gather(input, validate_signature=validate_signature) for input in inputs - ) + outputs = [] + for input in inputs: + self._validate_gather_input(input) + output_shape = (input.size(0) * self.world_size, *input.shape[1:]) + output = torch.empty(output_shape, dtype=input.dtype, device=input.device) + self._validate_sharded_output(output, input, output_shape) + outputs.append(output) + with self._lock: + if validate_signature: + self._validate_matching_many_signature("all_gather_many", inputs) + self._validate_many_capacity(inputs) + self._extension.deterministic_collective_all_gather_many( + self._handle, + list(inputs), + outputs, + ) + return tuple(outputs) def reduce_scatter( self, @@ -283,33 +386,17 @@ def reduce_scatter( def reduce_scatter_many( self, - inputs: tuple[torch.Tensor, ...] | list[torch.Tensor], + inputs: tuple[torch.Tensor, ...], *, - outs: tuple[torch.Tensor, ...] | list[torch.Tensor] | None = None, validate_signature: bool = True, ) -> tuple[torch.Tensor, ...]: - """Compatibility fallback for CUDA IPC collectives. + """Reduce-scatter several tensors through the single-tensor ABI.""" - The native CUDA IPC backend has no packed transport primitive yet, so - it preserves its established behavior by issuing the individual - fixed-tree calls. The ROCm transport subclass overrides this method - with a packed implementation. - """ - - values = tuple(inputs) - if not values: + if not inputs: raise ValueError("reduce_scatter_many requires at least one input") - if outs is not None and len(outs) != len(values): - raise ValueError("reduce_scatter_many outs must match the number of inputs") - results = tuple( - self.reduce_scatter( - value, - out=None if outs is None else outs[index], - validate_signature=validate_signature, - ) - for index, value in enumerate(values) + return tuple( + self.reduce_scatter(input, validate_signature=validate_signature) for input in inputs ) - return results def close(self) -> None: """Release imported CUDA IPC mappings after the last collective call.""" @@ -447,801 +534,13 @@ def _synchronize_ranks(self) -> None: dist.barrier(group=self.group) -class TorchDistributedDeterministicCollective: - """Correctness-first collectives using AllGather as transport only. - - Rank inputs are gathered without arithmetic and reduced locally as the - balanced tree ``((rank0 + rank1) + (rank2 + rank3)) + ...``. Consequently, - all ranks execute the exact same floating-point expression. TP sizes 1, - 2, 4, and 8 are nested prefixes of that expression and match the existing - CUDA IPC collective's ordering. - - The generic class also supports a CPU/Gloo process group, which is useful - as an executable reference. Production ROCm callers should use - :class:`RCCLDeterministicCollective` or - :func:`create_deterministic_collective` so backend validation fails closed. - All ranks must call methods in the same order with matching input shapes - and dtypes, and construct the instance with the same ``max_size_bytes``. - """ - - backend_id = "torch_distributed_balanced_tree" - transport_only = True - reduction_order = "balanced_rank_tree" - supports_async_overlap = False - supports_compute_communication_fusion = False - - def __init__( - self, - group: dist.ProcessGroup | None = None, - device: torch.device | str | int | None = None, - *, - max_size_bytes: int = _DEFAULT_MAX_SIZE_BYTES, - ) -> None: - if not dist.is_available() or not dist.is_initialized(): - raise RuntimeError("torch.distributed must be initialized before collectives") - if max_size_bytes <= 0: - raise ValueError("max_size_bytes must be positive") - - self.group = group if group is not None else dist.group.WORLD - self.rank = int(dist.get_rank(group=self.group)) - self.world_size = int(dist.get_world_size(group=self.group)) - if self.world_size not in _SUPPORTED_WORLD_SIZES: - raise ValueError( - "deterministic collectives require world_size in " - f"{_SUPPORTED_WORLD_SIZES}, got {self.world_size}" - ) - - self.device = self._normalize_device(device) - self.max_size_bytes = int(max_size_bytes) - self._backend = str(dist.get_backend(self.group)).lower() - self._lock = threading.Lock() - self._closed = False - # Keep a lifecycle marker for callers that historically inspected the - # CUDA IPC collective's ``_handle`` while managing the cache. Concrete - # transports own any native resource through their own state. - self._handle = id(self) - # One dtype-agnostic byte workspace is grown on demand and reused by - # reduction collectives. AllGather writes directly into its output. - self._workspace: torch.Tensor | None = None - # A Python-object collective is useful for catching a mismatched new - # signature, but running one on every hot-path call dominates small - # message latency. Validate each local signature once and then rely on - # the standard collective contract that ranks call operations in the - # same order. - self._validated_signatures: set[tuple[Any, ...]] = set() - self._validate_matching_capacity() - - @staticmethod - def _normalize_device( - device: torch.device | str | int | None, - ) -> torch.device: - if device is None: - if torch.cuda.is_available(): - normalized = torch.device("cuda", torch.cuda.current_device()) - else: - normalized = torch.device("cpu") - elif isinstance(device, int): - normalized = torch.device("cuda", device) - else: - normalized = torch.device(device) - - if normalized.type == "cuda": - if not torch.cuda.is_available(): - raise RuntimeError("a CUDA/ROCm device was requested but none is available") - current_device = torch.cuda.current_device() - if normalized.index is None: - normalized = torch.device("cuda", current_device) - if normalized.index != current_device: - raise ValueError( - "the collective device must be the current CUDA/ROCm device; call " - f"torch.cuda.set_device({normalized.index}) first" - ) - return normalized - - @property - def closed(self) -> bool: - """Whether this instance rejects further collective calls.""" - - return self._closed - - @property - def workspace_size_bytes(self) -> int: - """Currently retained reduction workspace size in bytes.""" - - workspace = self._workspace - return 0 if workspace is None else int(workspace.numel()) - - def all_reduce( - self, - input: torch.Tensor, - *, - out: torch.Tensor | None = None, - validate_signature: bool = True, - ) -> torch.Tensor: - """Return the fixed balanced-tree sum on every rank.""" - - self._check_open() - self._validate_reduction_input(input) - if out is None: - out = torch.empty_like(input) - self._validate_output(out, input, tuple(input.shape)) - if self.world_size == 1: - out.copy_(input) - return out - - with self._lock: - self._check_open() - if validate_signature: - self._validate_matching_signature("all_reduce", input) - if self._direct_all_reduce(input, out): - return out - rank_inputs = self._all_gather_transport(input) - if not self._fused_reduction( - rank_inputs, - out, - operation="all_reduce", - ): - reduced = self._balanced_tree_sum(rank_inputs) - out.copy_(reduced) - return out - - def all_gather( - self, - input: torch.Tensor, - *, - out: torch.Tensor | None = None, - validate_signature: bool = True, - ) -> torch.Tensor: - """Gather rank-ordered input bit patterns along dimension 0.""" - - self._check_open() - self._validate_gather_input(input) - output_shape = (input.size(0) * self.world_size, *input.shape[1:]) - if out is None: - out = torch.empty(output_shape, dtype=input.dtype, device=input.device) - self._validate_output(out, input, output_shape) - if self.world_size == 1: - out.copy_(input) - return out - - with self._lock: - self._check_open() - if validate_signature: - self._validate_matching_signature("all_gather", input) - if self._direct_all_gather(input, out): - return out - self._all_gather_transport(input, gathered_flat=out.view(-1)) - return out - - def all_gather_many( - self, - inputs: tuple[torch.Tensor, ...] | list[torch.Tensor], - *, - validate_signature: bool = True, - ) -> tuple[torch.Tensor, ...]: - """Gather several tensors through the platform transport.""" - - values = tuple(inputs) - if not values: - raise ValueError("all_gather_many requires at least one input") - return tuple( - self.all_gather(value, validate_signature=validate_signature) - for value in values - ) - - def reduce_scatter( - self, - input: torch.Tensor, - *, - out: torch.Tensor | None = None, - validate_signature: bool = True, - ) -> torch.Tensor: - """Fixed-tree sum followed by rank-ordered dimension-0 slicing.""" - - self._check_open() - self._validate_reduction_input(input) - if input.dim() == 0: - raise ValueError("reduce_scatter input must have at least one dimension") - if input.size(0) % self.world_size != 0: - raise ValueError( - "reduce_scatter input.size(0) must be divisible by " - f"world_size={self.world_size}; got {input.size(0)}" - ) - rows_per_rank = input.size(0) // self.world_size - output_shape = (rows_per_rank, *input.shape[1:]) - if out is None: - out = torch.empty(output_shape, dtype=input.dtype, device=input.device) - self._validate_output(out, input, output_shape) - if self.world_size == 1: - out.copy_(input) - return out - - with self._lock: - self._check_open() - if validate_signature: - self._validate_matching_signature("reduce_scatter", input) - if self._direct_reduce_scatter(input, out): - return out - rank_inputs = self._all_gather_transport(input) - begin = self.rank * rows_per_rank - # Only this rank's output shard participates in the reduction. The - # previous implementation reduced every global row and sliced the - # result afterwards, doing world_size times more arithmetic than - # ReduceScatter needs. The fixed rank tree is unchanged. - reduced = rank_inputs[:, begin : begin + rows_per_rank] - if not self._fused_reduction(reduced, out, operation="reduce_scatter"): - reduced = self._balanced_tree_sum(reduced) - out.copy_(reduced) - return out - - def reduce_scatter_many( - self, - inputs: tuple[torch.Tensor, ...] | list[torch.Tensor], - *, - outs: tuple[torch.Tensor, ...] | list[torch.Tensor] | None = None, - validate_signature: bool = True, - ) -> tuple[torch.Tensor, ...]: - """Reduce-scatter independent tensors in one fixed-tree collective. - - The tensors are packed along their final dimension, so each tensor's - element still follows the same balanced rank tree as an individual - ``reduce_scatter`` call. This is useful for independent gradient lanes: - packing them together removes one RCCL launch without changing the - floating-point expression for either lane. Inputs must have matching - shape/device/dtype except for the final dimension. - """ - - self._check_open() - values = tuple(inputs) - if not values: - raise ValueError("reduce_scatter_many requires at least one input") - if outs is not None and len(outs) != len(values): - raise ValueError("reduce_scatter_many outs must match the number of inputs") - if len(values) == 1: - return ( - self.reduce_scatter( - values[0], - out=None if outs is None else outs[0], - validate_signature=validate_signature, - ), - ) - - first = values[0] - self._validate_reduction_input(first) - if first.dim() < 2: - raise ValueError( - "reduce_scatter_many inputs must have at least two dimensions " - "when packing independent lanes" - ) - if first.size(0) % self.world_size != 0: - raise ValueError("reduce_scatter_many inputs must have a divisible leading dimension") - for value in values[1:]: - self._validate_reduction_input(value) - if value.dim() != first.dim() or value.shape[:-1] != first.shape[:-1]: - raise ValueError( - "reduce_scatter_many inputs must match in rank and all dimensions " - "except the final dimension" - ) - if value.device != first.device or value.dtype != first.dtype: - raise ValueError("reduce_scatter_many inputs must share device and dtype") - lane_sizes = tuple(int(value.size(-1)) for value in values) - rows_per_rank = first.size(0) // self.world_size - output_shape = (rows_per_rank, *first.shape[1:-1]) - if outs is not None: - for lane_size, out in zip(lane_sizes, outs, strict=True): - self._validate_output( - out, - first, - (*output_shape, lane_size), - ) - - packed_bytes = sum(value.numel() * value.element_size() for value in values) - if self._can_direct_reduce_scatter_many(): - if packed_bytes > self.max_size_bytes: - raise ValueError( - "reduce_scatter_many packed input requires " - f"{packed_bytes} bytes but max_size_bytes={self.max_size_bytes}" - ) - direct_outputs = tuple( - ( - outs[index] - if outs is not None - else torch.empty( - (*output_shape, lane_size), - dtype=first.dtype, - device=first.device, - ) - ) - for index, lane_size in enumerate(lane_sizes) - ) - with self._lock: - self._check_open() - if validate_signature: - self._validate_matching_signature( - f"reduce_scatter_many:{lane_sizes}", - first, - ) - if self._direct_reduce_scatter_many(values, direct_outputs): - return direct_outputs - - if packed_bytes > _PACKED_REDUCE_SCATTER_MAX_BYTES: - # A single packed AllGather moves the same bytes as two separate - # calls but loses RCCL's smaller-message algorithm. Use the - # established per-lane path above the measured crossover; this - # keeps the convenience API from regressing large FFN gradients. - return tuple( - self.reduce_scatter( - value, - out=None if outs is None else outs[index], - validate_signature=validate_signature, - ) - for index, value in enumerate(values) - ) - - if packed_bytes > self.max_size_bytes: - raise ValueError( - "reduce_scatter_many packed input requires " - f"{packed_bytes} bytes but max_size_bytes={self.max_size_bytes}" - ) - packed = torch.cat(values, dim=-1) - packed_out = torch.empty( - (packed.size(0) // self.world_size, *packed.shape[1:]), - dtype=packed.dtype, - device=packed.device, - ) - with self._lock: - self._check_open() - # Include lane boundaries in the signature. Equal packed shapes - # alone do not guarantee that every rank will split the result the - # same way, which could silently associate gradients with the - # wrong lane. - if validate_signature: - self._validate_matching_signature( - f"reduce_scatter_many:{lane_sizes}", - packed, - ) - if self._direct_reduce_scatter(packed, packed_out): - pieces = tuple(packed_out.split(tuple(value.size(-1) for value in values), dim=-1)) - if outs is None: - return pieces - result: list[torch.Tensor] = [] - for piece, out in zip(pieces, outs, strict=True): - out.copy_(piece) - result.append(out) - return tuple(result) - rank_inputs = self._all_gather_transport(packed) - begin = self.rank * rows_per_rank - reduced = rank_inputs[:, begin : begin + rows_per_rank] - if not self._fused_reduction( - reduced, - packed_out, - operation="reduce_scatter", - ): - reduced = self._balanced_tree_sum(reduced) - packed_out.copy_(reduced) - - pieces = tuple(packed_out.split(tuple(value.size(-1) for value in values), dim=-1)) - if outs is None: - return pieces - result: list[torch.Tensor] = [] - for piece, out in zip(pieces, outs, strict=True): - out.copy_(piece) - result.append(out) - return tuple(result) - - def close(self) -> None: - """Close the instance. - - Closing releases the lazily allocated reduction workspace and marks - the lifecycle boundary. Collective calls are blocking at this API. - """ - - with self._lock: - self._workspace = None - self._validated_signatures.clear() - self._closed = True - self._handle = 0 - - def __enter__(self) -> TorchDistributedDeterministicCollective: - self._check_open() - return self - - def __exit__( - self, - exc_type: type[BaseException] | None, - exc_value: BaseException | None, - traceback: TracebackType | None, - ) -> None: - self.close() - - def __del__(self) -> None: - try: - self.close() - except Exception: - pass - - def _check_open(self) -> None: - if getattr(self, "_closed", True): - raise RuntimeError("deterministic collective is closed") - - def _validate_tensor(self, input: torch.Tensor) -> None: - if not isinstance(input, torch.Tensor): - raise TypeError(f"input must be a torch.Tensor, got {type(input)!r}") - if input.device != self.device: - raise ValueError(f"input must be on {self.device}, got {input.device}") - if not input.is_contiguous(): - raise ValueError("input must be contiguous") - input_bytes = input.numel() * input.element_size() - if input_bytes > self.max_size_bytes: - raise ValueError( - f"input requires {input_bytes} bytes but max_size_bytes={self.max_size_bytes}" - ) - - def _validate_reduction_input(self, input: torch.Tensor) -> None: - self._validate_tensor(input) - if input.dtype not in _REDUCTION_DTYPES: - raise TypeError( - "deterministic reductions support float32, float16, and bfloat16; " - f"got {input.dtype}" - ) - - def _validate_gather_input(self, input: torch.Tensor) -> None: - self._validate_tensor(input) - if input.dim() == 0: - raise ValueError("all_gather input must have at least one dimension") - - @staticmethod - def _validate_output( - output: torch.Tensor, - input: torch.Tensor, - output_shape: tuple[int, ...], - ) -> None: - if not isinstance(output, torch.Tensor): - raise TypeError(f"out must be a torch.Tensor, got {type(output)!r}") - if output.device != input.device: - raise ValueError("out must be on the same device as input") - if output.dtype != input.dtype: - raise TypeError("out must have the same dtype as input") - if tuple(output.shape) != output_shape: - raise ValueError(f"out must have shape {output_shape}, got {tuple(output.shape)}") - if not output.is_contiguous(): - raise ValueError("out must be contiguous") - - def _validate_matching_signature(self, op_name: str, input: torch.Tensor) -> None: - if self.world_size == 1: - return - signature = (op_name, tuple(input.shape), str(input.dtype), input.numel()) - if signature in self._validated_signatures: - return - signatures: list[tuple[Any, ...] | None] = [None] * self.world_size - dist.all_gather_object(signatures, signature, group=self.group) - if any(peer_signature != signature for peer_signature in signatures): - raise ValueError( - f"all ranks must call {op_name} with matching shapes and dtypes; got {signatures}" - ) - self._validated_signatures.add(signature) - - def _validate_matching_capacity(self) -> None: - if self.world_size == 1: - return - capacities: list[int | None] = [None] * self.world_size - dist.all_gather_object(capacities, self.max_size_bytes, group=self.group) - if any(peer_capacity != self.max_size_bytes for peer_capacity in capacities): - raise ValueError(f"all ranks must use the same max_size_bytes; got {capacities}") - - def _all_gather_transport( - self, - input: torch.Tensor, - *, - gathered_flat: torch.Tensor | None = None, - ) -> torch.Tensor: - # Flattening makes the output contract independent of whether a given - # ProcessGroup implements the concatenation or stacking form of AG. - if self.world_size == 1: - if gathered_flat is None: - gathered_flat = input.clone().view(-1) - else: - gathered_flat.copy_(input.view(-1)) - return gathered_flat.reshape((1, *input.shape)) - input_flat = input.view(-1) - required_elements = self.world_size * input_flat.numel() - if gathered_flat is None: - gathered_flat = self._workspace_for(input, required_elements) - elif ( - gathered_flat.numel() != required_elements - or gathered_flat.dtype != input.dtype - or gathered_flat.device != input.device - or not gathered_flat.is_contiguous() - ): - raise ValueError("gathered transport output has an invalid layout") - if "nccl" in self._backend: - # PyTorch exposes RCCL through the NCCL ProcessGroup API. Keep this - # as a tensor-only transport; reduction happens below. - dist.all_gather_into_tensor(gathered_flat, input_flat, group=self.group) - else: - # Some reference backends (notably Gloo versions without - # all_gather_into_tensor) only implement the list API. - gathered_chunks = list( - gathered_flat.reshape(self.world_size, input_flat.numel()).unbind(0) - ) - dist.all_gather(gathered_chunks, input_flat, group=self.group) - return gathered_flat.reshape((self.world_size, *input.shape)) - - def _workspace_for(self, input: torch.Tensor, required_elements: int) -> torch.Tensor: - required_bytes = required_elements * input.element_size() - workspace = self._workspace - if workspace is None or workspace.numel() < required_bytes: - workspace = torch.empty(required_bytes, dtype=torch.uint8, device=self.device) - self._workspace = workspace - # Tensor.view(dtype) reinterprets the aligned byte allocation without - # an allocation or copy. Restrict the view to the current operation. - return workspace[:required_bytes].view(input.dtype) - - def _direct_all_reduce(self, input: torch.Tensor, output: torch.Tensor) -> bool: - return False - - def _direct_reduce_scatter(self, input: torch.Tensor, output: torch.Tensor) -> bool: - return False - - def _can_direct_reduce_scatter_many(self) -> bool: - return False - - def _direct_reduce_scatter_many( - self, - inputs: tuple[torch.Tensor, ...], - outputs: tuple[torch.Tensor, ...], - ) -> bool: - return False - - def _direct_all_gather(self, input: torch.Tensor, output: torch.Tensor) -> bool: - return False - - @staticmethod - def _balanced_tree_sum(rank_inputs: torch.Tensor) -> torch.Tensor: - world_size = rank_inputs.size(0) - if world_size not in _SUPPORTED_WORLD_SIZES: - raise ValueError( - "balanced reduction requires rank inputs for world_size in " - f"{_SUPPORTED_WORLD_SIZES}, got {world_size}" - ) - # ``rank_inputs`` is the private transport workspace for reductions, - # so fixed-tree nodes can be accumulated in place. This preserves the - # exact pairings while avoiding one temporary allocation per tree node. - stride = 1 - while stride < world_size: - for index in range(0, world_size, 2 * stride): - rank_inputs[index].add_(rank_inputs[index + stride]) - stride *= 2 - return rank_inputs[0] - - @staticmethod - def _fused_reduction( - rank_inputs: torch.Tensor, - output: torch.Tensor, - *, - operation: str, - ) -> bool: - """Use the optional ROCm fused fixed-tree kernel when available. - - The extension is deliberately optional: CPU/Gloo reference collectives - and installations built without the ROCm kernel retain the executable - Python implementation above. - """ - - if getattr(torch.version, "hip", None) is None or not rank_inputs.is_cuda: - return False - try: - from rl_engine import _C - except ImportError: - return False - if operation == "all_reduce": - fn = getattr(_C, "deterministic_collective_rocm_all_reduce", None) - if fn is not None: - fn(rank_inputs, output) - return True - elif operation == "reduce_scatter": - fn = getattr(_C, "deterministic_collective_rocm_reduce_scatter", None) - if fn is not None: - fn(rank_inputs, output) - return True - return False - - -class RCCLDeterministicCollective(TorchDistributedDeterministicCollective): - """Single-node ROCm fixed-tree collective using HIP IPC and RCCL.""" - - backend_id = "rocm_ipc_fixed_tree" - - def __init__( - self, - group: dist.ProcessGroup | None = None, - device: torch.device | str | int | None = None, - *, - max_size_bytes: int = _DEFAULT_MAX_SIZE_BYTES, - ) -> None: - if getattr(torch.version, "hip", None) is None: - raise RuntimeError("RCCL deterministic collectives require a ROCm PyTorch build") - if not torch.cuda.is_available(): - raise RuntimeError("RCCL deterministic collectives require an available ROCm device") - if device is not None and not isinstance(device, int): - requested_device = torch.device(device) - if requested_device.type != "cuda": - raise ValueError( - f"RCCL deterministic collectives require a ROCm device, got {device!r}" - ) - super().__init__(group=group, device=device, max_size_bytes=max_size_bytes) - if self.device.type != "cuda": - raise ValueError( - f"RCCL deterministic collectives require a ROCm device, got {device!r}" - ) - if "nccl" not in self._backend: - raise RuntimeError( - "RCCL deterministic collectives require PyTorch's NCCL process-group API" - ) - self._ipc_handle = 0 - self._ipc_staging: torch.Tensor | None = None - self._initialize_ipc_transport() - - @property - def workspace_size_bytes(self) -> int: - staging = self._ipc_staging - staging_bytes = 0 if staging is None else int(staging.numel()) - return staging_bytes + super().workspace_size_bytes - - def _initialize_ipc_transport(self) -> None: - if self.world_size == 1: - return - try: - from rl_engine import _C - except ImportError: - return - required_symbols = ( - "deterministic_collective_rocm_ipc_allocate", - "deterministic_collective_rocm_ipc_meta", - "deterministic_collective_rocm_ipc_create", - "deterministic_collective_rocm_ipc_synchronize", - "deterministic_collective_rocm_ipc_destroy", - "deterministic_collective_rocm_ipc_stage", - "deterministic_collective_rocm_ipc_all_reduce", - "deterministic_collective_rocm_ipc_all_reduce_input", - "deterministic_collective_rocm_ipc_reduce_scatter", - "deterministic_collective_rocm_ipc_reduce_scatter_input", - "deterministic_collective_rocm_ipc_reduce_scatter_many", - "deterministic_collective_rocm_ipc_all_gather", - "deterministic_collective_rocm_ipc_all_gather_input", - ) - if any(not hasattr(_C, symbol) for symbol in required_symbols): - return - - staging = _C.deterministic_collective_rocm_ipc_allocate(self.max_size_bytes) - handle, offset = _C.deterministic_collective_rocm_ipc_meta(staging) - local_metadata = (socket.gethostname(), handle, int(offset)) - gathered_metadata: list[tuple[str, list[int], int] | None] = [None] * self.world_size - dist.all_gather_object(gathered_metadata, local_metadata, group=self.group) - if any(metadata is None for metadata in gathered_metadata): - raise RuntimeError("failed to exchange ROCm IPC metadata") - complete_metadata = [metadata for metadata in gathered_metadata if metadata is not None] - if len({metadata[0] for metadata in complete_metadata}) != 1: - return - self._ipc_handle = int( - _C.deterministic_collective_rocm_ipc_create( - staging, - [metadata[1] for metadata in complete_metadata], - [metadata[2] for metadata in complete_metadata], - self.rank, - ) - ) - self._ipc_staging = staging - - def _direct_all_reduce(self, input: torch.Tensor, output: torch.Tensor) -> bool: - handle = self._ipc_handle - if not handle: - return False - from rl_engine import _C - - input_bytes = input.numel() * input.element_size() - if ( - _ROCM_IPC_DIRECT_ALL_REDUCE_MAX_BYTES - < input_bytes - < _ROCM_IPC_SHARDED_ALL_REDUCE_MIN_BYTES - and input.numel() % self.world_size == 0 - ): - return False - - if ( - input_bytes <= _ROCM_IPC_DIRECT_ALL_REDUCE_MAX_BYTES - or input.numel() % self.world_size != 0 - ): - _C.deterministic_collective_rocm_ipc_all_reduce_input( - handle, - input, - output, - ) - return True - - shard = self._workspace_for(input, input.numel() // self.world_size) - _C.deterministic_collective_rocm_ipc_reduce_scatter_input( - handle, - input, - shard, - ) - dist.all_gather_into_tensor(output.view(-1), shard, group=self.group) - return True - - def _direct_reduce_scatter(self, input: torch.Tensor, output: torch.Tensor) -> bool: - handle = self._ipc_handle - if not handle: - return False - from rl_engine import _C - - _C.deterministic_collective_rocm_ipc_reduce_scatter_input( - handle, - input, - output, - ) - return True - - def _can_direct_reduce_scatter_many(self) -> bool: - return bool(self._ipc_handle) - - def _direct_reduce_scatter_many( - self, - inputs: tuple[torch.Tensor, ...], - outputs: tuple[torch.Tensor, ...], - ) -> bool: - handle = self._ipc_handle - if not handle: - return False - from rl_engine import _C - - _C.deterministic_collective_rocm_ipc_reduce_scatter_many( - handle, - inputs, - outputs, - ) - return True - - def _direct_all_gather(self, input: torch.Tensor, output: torch.Tensor) -> bool: - handle = self._ipc_handle - input_bytes = input.numel() * input.element_size() - if not handle or input_bytes > _ROCM_IPC_ALL_GATHER_MAX_BYTES: - return False - from rl_engine import _C - - _C.deterministic_collective_rocm_ipc_all_gather_input( - handle, - input, - output, - ) - return True - - def close(self) -> None: - handle = getattr(self, "_ipc_handle", 0) - if handle: - from rl_engine import _C - - _C.deterministic_collective_rocm_ipc_synchronize(handle) - torch.cuda.synchronize(self.device) - self._ipc_handle = 0 - _C.deterministic_collective_rocm_ipc_destroy(handle) - self._ipc_staging = None - super().close() - - def create_deterministic_collective( group: dist.ProcessGroup | None = None, device: torch.device | str | int | None = None, *, max_size_bytes: int = _DEFAULT_MAX_SIZE_BYTES, ) -> Any: - """Create the platform-appropriate deterministic collective. - - CUDA uses the native ``DeterministicCollective`` implementation. ROCm uses - HIP IPC or RCCL for rank-ordered transport while preserving the fixed local - reduction tree. The returned object has independent ownership. Shared caches - may replace an entry without closing it immediately because active autograd - contexts can retain the previous instance until their work completes. - """ + """Create the deterministic collective for the active accelerator.""" if getattr(torch.version, "hip", None) is not None: return RCCLDeterministicCollective( @@ -1273,8 +572,8 @@ def collective_for_group( if minimum_capacity_bytes <= 0: raise ValueError("minimum_capacity_bytes must be positive") - rank = int(dist.get_rank(group=group)) - world_size = int(dist.get_world_size(group=group)) + rank = dist.get_rank(group=group) + world_size = dist.get_world_size(group=group) if device is None: device_index = torch.cuda.current_device() else: @@ -1303,6 +602,16 @@ def collective_for_group( return collective +# Keep the public API provided by the ROCm test branch while leaving the CUDA +# implementation above identical to PR377. The ROCm classes own no device +# resources until constructed, so importing their definitions has no hot-path +# or accelerator side effect. +from rl_engine.distributed.rocm_collectives import ( # noqa: E402 + RCCLDeterministicCollective, + TorchDistributedDeterministicCollective, +) + + __all__ = [ "DETERMINISTIC_ALL_REDUCE_OP", "DeterministicCollective", diff --git a/rl_engine/distributed/rocm_collectives.py b/rl_engine/distributed/rocm_collectives.py new file mode 100644 index 00000000..ace5dac2 --- /dev/null +++ b/rl_engine/distributed/rocm_collectives.py @@ -0,0 +1,800 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors + +from __future__ import annotations + +import socket +import threading +from types import TracebackType +from typing import Any + +import torch +import torch.distributed as dist + +_SUPPORTED_WORLD_SIZES = (1, 2, 4, 8) +_DEFAULT_MAX_SIZE_BYTES = 64 * 1024 * 1024 +_PACKED_REDUCE_SCATTER_MAX_BYTES = 8 * 1024 * 1024 +_ROCM_IPC_DIRECT_ALL_REDUCE_MAX_BYTES = 768 * 1024 +_ROCM_IPC_SHARDED_ALL_REDUCE_MIN_BYTES = 2176 * 1024 +_ROCM_IPC_ALL_GATHER_MAX_BYTES = 256 * 1024 +_REDUCTION_DTYPES = (torch.float32, torch.float16, torch.bfloat16) + +class TorchDistributedDeterministicCollective: + """Correctness-first collectives using AllGather as transport only. + + Rank inputs are gathered without arithmetic and reduced locally as the + balanced tree ``((rank0 + rank1) + (rank2 + rank3)) + ...``. Consequently, + all ranks execute the exact same floating-point expression. TP sizes 1, + 2, 4, and 8 are nested prefixes of that expression and match the existing + CUDA IPC collective's ordering. + + The generic class also supports a CPU/Gloo process group, which is useful + as an executable reference. Production ROCm callers should use + :class:`RCCLDeterministicCollective` or + :func:`create_deterministic_collective` so backend validation fails closed. + All ranks must call methods in the same order with matching input shapes + and dtypes, and construct the instance with the same ``max_size_bytes``. + """ + + backend_id = "torch_distributed_balanced_tree" + transport_only = True + reduction_order = "balanced_rank_tree" + supports_async_overlap = False + supports_compute_communication_fusion = False + + def __init__( + self, + group: dist.ProcessGroup | None = None, + device: torch.device | str | int | None = None, + *, + max_size_bytes: int = _DEFAULT_MAX_SIZE_BYTES, + ) -> None: + if not dist.is_available() or not dist.is_initialized(): + raise RuntimeError("torch.distributed must be initialized before collectives") + if max_size_bytes <= 0: + raise ValueError("max_size_bytes must be positive") + + self.group = group if group is not None else dist.group.WORLD + self.rank = int(dist.get_rank(group=self.group)) + self.world_size = int(dist.get_world_size(group=self.group)) + if self.world_size not in _SUPPORTED_WORLD_SIZES: + raise ValueError( + "deterministic collectives require world_size in " + f"{_SUPPORTED_WORLD_SIZES}, got {self.world_size}" + ) + + self.device = self._normalize_device(device) + self.max_size_bytes = int(max_size_bytes) + self._backend = str(dist.get_backend(self.group)).lower() + self._lock = threading.Lock() + self._closed = False + # Keep a lifecycle marker for callers that historically inspected the + # CUDA IPC collective's ``_handle`` while managing the cache. Concrete + # transports own any native resource through their own state. + self._handle = id(self) + # One dtype-agnostic byte workspace is grown on demand and reused by + # reduction collectives. AllGather writes directly into its output. + self._workspace: torch.Tensor | None = None + # A Python-object collective is useful for catching a mismatched new + # signature, but running one on every hot-path call dominates small + # message latency. Validate each local signature once and then rely on + # the standard collective contract that ranks call operations in the + # same order. + self._validated_signatures: set[tuple[Any, ...]] = set() + self._validate_matching_capacity() + + @staticmethod + def _normalize_device( + device: torch.device | str | int | None, + ) -> torch.device: + if device is None: + if torch.cuda.is_available(): + normalized = torch.device("cuda", torch.cuda.current_device()) + else: + normalized = torch.device("cpu") + elif isinstance(device, int): + normalized = torch.device("cuda", device) + else: + normalized = torch.device(device) + + if normalized.type == "cuda": + if not torch.cuda.is_available(): + raise RuntimeError("a CUDA/ROCm device was requested but none is available") + current_device = torch.cuda.current_device() + if normalized.index is None: + normalized = torch.device("cuda", current_device) + if normalized.index != current_device: + raise ValueError( + "the collective device must be the current CUDA/ROCm device; call " + f"torch.cuda.set_device({normalized.index}) first" + ) + return normalized + + @property + def closed(self) -> bool: + """Whether this instance rejects further collective calls.""" + + return self._closed + + @property + def workspace_size_bytes(self) -> int: + """Currently retained reduction workspace size in bytes.""" + + workspace = self._workspace + return 0 if workspace is None else int(workspace.numel()) + + def all_reduce( + self, + input: torch.Tensor, + *, + out: torch.Tensor | None = None, + validate_signature: bool = True, + ) -> torch.Tensor: + """Return the fixed balanced-tree sum on every rank.""" + + self._check_open() + self._validate_reduction_input(input) + if out is None: + out = torch.empty_like(input) + self._validate_output(out, input, tuple(input.shape)) + if self.world_size == 1: + out.copy_(input) + return out + + with self._lock: + self._check_open() + if validate_signature: + self._validate_matching_signature("all_reduce", input) + if self._direct_all_reduce(input, out): + return out + rank_inputs = self._all_gather_transport(input) + if not self._fused_reduction( + rank_inputs, + out, + operation="all_reduce", + ): + reduced = self._balanced_tree_sum(rank_inputs) + out.copy_(reduced) + return out + + def all_gather( + self, + input: torch.Tensor, + *, + out: torch.Tensor | None = None, + validate_signature: bool = True, + ) -> torch.Tensor: + """Gather rank-ordered input bit patterns along dimension 0.""" + + self._check_open() + self._validate_gather_input(input) + output_shape = (input.size(0) * self.world_size, *input.shape[1:]) + if out is None: + out = torch.empty(output_shape, dtype=input.dtype, device=input.device) + self._validate_output(out, input, output_shape) + if self.world_size == 1: + out.copy_(input) + return out + + with self._lock: + self._check_open() + if validate_signature: + self._validate_matching_signature("all_gather", input) + if self._direct_all_gather(input, out): + return out + self._all_gather_transport(input, gathered_flat=out.view(-1)) + return out + + def all_gather_many( + self, + inputs: tuple[torch.Tensor, ...] | list[torch.Tensor], + *, + validate_signature: bool = True, + ) -> tuple[torch.Tensor, ...]: + """Gather several tensors through the platform transport.""" + + values = tuple(inputs) + if not values: + raise ValueError("all_gather_many requires at least one input") + return tuple( + self.all_gather(value, validate_signature=validate_signature) + for value in values + ) + + def reduce_scatter( + self, + input: torch.Tensor, + *, + out: torch.Tensor | None = None, + validate_signature: bool = True, + ) -> torch.Tensor: + """Fixed-tree sum followed by rank-ordered dimension-0 slicing.""" + + self._check_open() + self._validate_reduction_input(input) + if input.dim() == 0: + raise ValueError("reduce_scatter input must have at least one dimension") + if input.size(0) % self.world_size != 0: + raise ValueError( + "reduce_scatter input.size(0) must be divisible by " + f"world_size={self.world_size}; got {input.size(0)}" + ) + rows_per_rank = input.size(0) // self.world_size + output_shape = (rows_per_rank, *input.shape[1:]) + if out is None: + out = torch.empty(output_shape, dtype=input.dtype, device=input.device) + self._validate_output(out, input, output_shape) + if self.world_size == 1: + out.copy_(input) + return out + + with self._lock: + self._check_open() + if validate_signature: + self._validate_matching_signature("reduce_scatter", input) + if self._direct_reduce_scatter(input, out): + return out + rank_inputs = self._all_gather_transport(input) + begin = self.rank * rows_per_rank + # Only this rank's output shard participates in the reduction. The + # previous implementation reduced every global row and sliced the + # result afterwards, doing world_size times more arithmetic than + # ReduceScatter needs. The fixed rank tree is unchanged. + reduced = rank_inputs[:, begin : begin + rows_per_rank] + if not self._fused_reduction(reduced, out, operation="reduce_scatter"): + reduced = self._balanced_tree_sum(reduced) + out.copy_(reduced) + return out + + def reduce_scatter_many( + self, + inputs: tuple[torch.Tensor, ...] | list[torch.Tensor], + *, + outs: tuple[torch.Tensor, ...] | list[torch.Tensor] | None = None, + validate_signature: bool = True, + ) -> tuple[torch.Tensor, ...]: + """Reduce-scatter independent tensors in one fixed-tree collective. + + The tensors are packed along their final dimension, so each tensor's + element still follows the same balanced rank tree as an individual + ``reduce_scatter`` call. This is useful for independent gradient lanes: + packing them together removes one RCCL launch without changing the + floating-point expression for either lane. Inputs must have matching + shape/device/dtype except for the final dimension. + """ + + self._check_open() + values = tuple(inputs) + if not values: + raise ValueError("reduce_scatter_many requires at least one input") + if outs is not None and len(outs) != len(values): + raise ValueError("reduce_scatter_many outs must match the number of inputs") + if len(values) == 1: + return ( + self.reduce_scatter( + values[0], + out=None if outs is None else outs[0], + validate_signature=validate_signature, + ), + ) + + first = values[0] + self._validate_reduction_input(first) + if first.dim() < 2: + raise ValueError( + "reduce_scatter_many inputs must have at least two dimensions " + "when packing independent lanes" + ) + if first.size(0) % self.world_size != 0: + raise ValueError("reduce_scatter_many inputs must have a divisible leading dimension") + for value in values[1:]: + self._validate_reduction_input(value) + if value.dim() != first.dim() or value.shape[:-1] != first.shape[:-1]: + raise ValueError( + "reduce_scatter_many inputs must match in rank and all dimensions " + "except the final dimension" + ) + if value.device != first.device or value.dtype != first.dtype: + raise ValueError("reduce_scatter_many inputs must share device and dtype") + lane_sizes = tuple(int(value.size(-1)) for value in values) + rows_per_rank = first.size(0) // self.world_size + output_shape = (rows_per_rank, *first.shape[1:-1]) + if outs is not None: + for lane_size, out in zip(lane_sizes, outs, strict=True): + self._validate_output( + out, + first, + (*output_shape, lane_size), + ) + + packed_bytes = sum(value.numel() * value.element_size() for value in values) + if self._can_direct_reduce_scatter_many(): + if packed_bytes > self.max_size_bytes: + raise ValueError( + "reduce_scatter_many packed input requires " + f"{packed_bytes} bytes but max_size_bytes={self.max_size_bytes}" + ) + direct_outputs = tuple( + ( + outs[index] + if outs is not None + else torch.empty( + (*output_shape, lane_size), + dtype=first.dtype, + device=first.device, + ) + ) + for index, lane_size in enumerate(lane_sizes) + ) + with self._lock: + self._check_open() + if validate_signature: + self._validate_matching_signature( + f"reduce_scatter_many:{lane_sizes}", + first, + ) + if self._direct_reduce_scatter_many(values, direct_outputs): + return direct_outputs + + if packed_bytes > _PACKED_REDUCE_SCATTER_MAX_BYTES: + # A single packed AllGather moves the same bytes as two separate + # calls but loses RCCL's smaller-message algorithm. Use the + # established per-lane path above the measured crossover; this + # keeps the convenience API from regressing large FFN gradients. + return tuple( + self.reduce_scatter( + value, + out=None if outs is None else outs[index], + validate_signature=validate_signature, + ) + for index, value in enumerate(values) + ) + + if packed_bytes > self.max_size_bytes: + raise ValueError( + "reduce_scatter_many packed input requires " + f"{packed_bytes} bytes but max_size_bytes={self.max_size_bytes}" + ) + packed = torch.cat(values, dim=-1) + packed_out = torch.empty( + (packed.size(0) // self.world_size, *packed.shape[1:]), + dtype=packed.dtype, + device=packed.device, + ) + with self._lock: + self._check_open() + # Include lane boundaries in the signature. Equal packed shapes + # alone do not guarantee that every rank will split the result the + # same way, which could silently associate gradients with the + # wrong lane. + if validate_signature: + self._validate_matching_signature( + f"reduce_scatter_many:{lane_sizes}", + packed, + ) + if self._direct_reduce_scatter(packed, packed_out): + pieces = tuple(packed_out.split(tuple(value.size(-1) for value in values), dim=-1)) + if outs is None: + return pieces + result: list[torch.Tensor] = [] + for piece, out in zip(pieces, outs, strict=True): + out.copy_(piece) + result.append(out) + return tuple(result) + rank_inputs = self._all_gather_transport(packed) + begin = self.rank * rows_per_rank + reduced = rank_inputs[:, begin : begin + rows_per_rank] + if not self._fused_reduction( + reduced, + packed_out, + operation="reduce_scatter", + ): + reduced = self._balanced_tree_sum(reduced) + packed_out.copy_(reduced) + + pieces = tuple(packed_out.split(tuple(value.size(-1) for value in values), dim=-1)) + if outs is None: + return pieces + result: list[torch.Tensor] = [] + for piece, out in zip(pieces, outs, strict=True): + out.copy_(piece) + result.append(out) + return tuple(result) + + def close(self) -> None: + """Close the instance. + + Closing releases the lazily allocated reduction workspace and marks + the lifecycle boundary. Collective calls are blocking at this API. + """ + + with self._lock: + self._workspace = None + self._validated_signatures.clear() + self._closed = True + self._handle = 0 + + def __enter__(self) -> TorchDistributedDeterministicCollective: + self._check_open() + return self + + def __exit__( + self, + exc_type: type[BaseException] | None, + exc_value: BaseException | None, + traceback: TracebackType | None, + ) -> None: + self.close() + + def __del__(self) -> None: + try: + self.close() + except Exception: + pass + + def _check_open(self) -> None: + if getattr(self, "_closed", True): + raise RuntimeError("deterministic collective is closed") + + def _validate_tensor(self, input: torch.Tensor) -> None: + if not isinstance(input, torch.Tensor): + raise TypeError(f"input must be a torch.Tensor, got {type(input)!r}") + if input.device != self.device: + raise ValueError(f"input must be on {self.device}, got {input.device}") + if not input.is_contiguous(): + raise ValueError("input must be contiguous") + input_bytes = input.numel() * input.element_size() + if input_bytes > self.max_size_bytes: + raise ValueError( + f"input requires {input_bytes} bytes but max_size_bytes={self.max_size_bytes}" + ) + + def _validate_reduction_input(self, input: torch.Tensor) -> None: + self._validate_tensor(input) + if input.dtype not in _REDUCTION_DTYPES: + raise TypeError( + "deterministic reductions support float32, float16, and bfloat16; " + f"got {input.dtype}" + ) + + def _validate_gather_input(self, input: torch.Tensor) -> None: + self._validate_tensor(input) + if input.dim() == 0: + raise ValueError("all_gather input must have at least one dimension") + + @staticmethod + def _validate_output( + output: torch.Tensor, + input: torch.Tensor, + output_shape: tuple[int, ...], + ) -> None: + if not isinstance(output, torch.Tensor): + raise TypeError(f"out must be a torch.Tensor, got {type(output)!r}") + if output.device != input.device: + raise ValueError("out must be on the same device as input") + if output.dtype != input.dtype: + raise TypeError("out must have the same dtype as input") + if tuple(output.shape) != output_shape: + raise ValueError(f"out must have shape {output_shape}, got {tuple(output.shape)}") + if not output.is_contiguous(): + raise ValueError("out must be contiguous") + + def _validate_matching_signature(self, op_name: str, input: torch.Tensor) -> None: + if self.world_size == 1: + return + signature = (op_name, tuple(input.shape), str(input.dtype), input.numel()) + if signature in self._validated_signatures: + return + signatures: list[tuple[Any, ...] | None] = [None] * self.world_size + dist.all_gather_object(signatures, signature, group=self.group) + if any(peer_signature != signature for peer_signature in signatures): + raise ValueError( + f"all ranks must call {op_name} with matching shapes and dtypes; got {signatures}" + ) + self._validated_signatures.add(signature) + + def _validate_matching_capacity(self) -> None: + if self.world_size == 1: + return + capacities: list[int | None] = [None] * self.world_size + dist.all_gather_object(capacities, self.max_size_bytes, group=self.group) + if any(peer_capacity != self.max_size_bytes for peer_capacity in capacities): + raise ValueError(f"all ranks must use the same max_size_bytes; got {capacities}") + + def _all_gather_transport( + self, + input: torch.Tensor, + *, + gathered_flat: torch.Tensor | None = None, + ) -> torch.Tensor: + # Flattening makes the output contract independent of whether a given + # ProcessGroup implements the concatenation or stacking form of AG. + if self.world_size == 1: + if gathered_flat is None: + gathered_flat = input.clone().view(-1) + else: + gathered_flat.copy_(input.view(-1)) + return gathered_flat.reshape((1, *input.shape)) + input_flat = input.view(-1) + required_elements = self.world_size * input_flat.numel() + if gathered_flat is None: + gathered_flat = self._workspace_for(input, required_elements) + elif ( + gathered_flat.numel() != required_elements + or gathered_flat.dtype != input.dtype + or gathered_flat.device != input.device + or not gathered_flat.is_contiguous() + ): + raise ValueError("gathered transport output has an invalid layout") + if "nccl" in self._backend: + # PyTorch exposes RCCL through the NCCL ProcessGroup API. Keep this + # as a tensor-only transport; reduction happens below. + dist.all_gather_into_tensor(gathered_flat, input_flat, group=self.group) + else: + # Some reference backends (notably Gloo versions without + # all_gather_into_tensor) only implement the list API. + gathered_chunks = list( + gathered_flat.reshape(self.world_size, input_flat.numel()).unbind(0) + ) + dist.all_gather(gathered_chunks, input_flat, group=self.group) + return gathered_flat.reshape((self.world_size, *input.shape)) + + def _workspace_for(self, input: torch.Tensor, required_elements: int) -> torch.Tensor: + required_bytes = required_elements * input.element_size() + workspace = self._workspace + if workspace is None or workspace.numel() < required_bytes: + workspace = torch.empty(required_bytes, dtype=torch.uint8, device=self.device) + self._workspace = workspace + # Tensor.view(dtype) reinterprets the aligned byte allocation without + # an allocation or copy. Restrict the view to the current operation. + return workspace[:required_bytes].view(input.dtype) + + def _direct_all_reduce(self, input: torch.Tensor, output: torch.Tensor) -> bool: + return False + + def _direct_reduce_scatter(self, input: torch.Tensor, output: torch.Tensor) -> bool: + return False + + def _can_direct_reduce_scatter_many(self) -> bool: + return False + + def _direct_reduce_scatter_many( + self, + inputs: tuple[torch.Tensor, ...], + outputs: tuple[torch.Tensor, ...], + ) -> bool: + return False + + def _direct_all_gather(self, input: torch.Tensor, output: torch.Tensor) -> bool: + return False + + @staticmethod + def _balanced_tree_sum(rank_inputs: torch.Tensor) -> torch.Tensor: + world_size = rank_inputs.size(0) + if world_size not in _SUPPORTED_WORLD_SIZES: + raise ValueError( + "balanced reduction requires rank inputs for world_size in " + f"{_SUPPORTED_WORLD_SIZES}, got {world_size}" + ) + # ``rank_inputs`` is the private transport workspace for reductions, + # so fixed-tree nodes can be accumulated in place. This preserves the + # exact pairings while avoiding one temporary allocation per tree node. + stride = 1 + while stride < world_size: + for index in range(0, world_size, 2 * stride): + rank_inputs[index].add_(rank_inputs[index + stride]) + stride *= 2 + return rank_inputs[0] + + @staticmethod + def _fused_reduction( + rank_inputs: torch.Tensor, + output: torch.Tensor, + *, + operation: str, + ) -> bool: + """Use the optional ROCm fused fixed-tree kernel when available. + + The extension is deliberately optional: CPU/Gloo reference collectives + and installations built without the ROCm kernel retain the executable + Python implementation above. + """ + + if getattr(torch.version, "hip", None) is None or not rank_inputs.is_cuda: + return False + try: + from rl_engine import _C + except ImportError: + return False + if operation == "all_reduce": + fn = getattr(_C, "deterministic_collective_rocm_all_reduce", None) + if fn is not None: + fn(rank_inputs, output) + return True + elif operation == "reduce_scatter": + fn = getattr(_C, "deterministic_collective_rocm_reduce_scatter", None) + if fn is not None: + fn(rank_inputs, output) + return True + return False + + +class RCCLDeterministicCollective(TorchDistributedDeterministicCollective): + """Single-node ROCm fixed-tree collective using HIP IPC and RCCL.""" + + backend_id = "rocm_ipc_fixed_tree" + + def __init__( + self, + group: dist.ProcessGroup | None = None, + device: torch.device | str | int | None = None, + *, + max_size_bytes: int = _DEFAULT_MAX_SIZE_BYTES, + ) -> None: + if getattr(torch.version, "hip", None) is None: + raise RuntimeError("RCCL deterministic collectives require a ROCm PyTorch build") + if not torch.cuda.is_available(): + raise RuntimeError("RCCL deterministic collectives require an available ROCm device") + if device is not None and not isinstance(device, int): + requested_device = torch.device(device) + if requested_device.type != "cuda": + raise ValueError( + f"RCCL deterministic collectives require a ROCm device, got {device!r}" + ) + super().__init__(group=group, device=device, max_size_bytes=max_size_bytes) + if self.device.type != "cuda": + raise ValueError( + f"RCCL deterministic collectives require a ROCm device, got {device!r}" + ) + if "nccl" not in self._backend: + raise RuntimeError( + "RCCL deterministic collectives require PyTorch's NCCL process-group API" + ) + self._ipc_handle = 0 + self._ipc_staging: torch.Tensor | None = None + self._initialize_ipc_transport() + + @property + def workspace_size_bytes(self) -> int: + staging = self._ipc_staging + staging_bytes = 0 if staging is None else int(staging.numel()) + return staging_bytes + super().workspace_size_bytes + + def _initialize_ipc_transport(self) -> None: + if self.world_size == 1: + return + try: + from rl_engine import _C + except ImportError: + return + required_symbols = ( + "deterministic_collective_rocm_ipc_allocate", + "deterministic_collective_rocm_ipc_meta", + "deterministic_collective_rocm_ipc_create", + "deterministic_collective_rocm_ipc_synchronize", + "deterministic_collective_rocm_ipc_destroy", + "deterministic_collective_rocm_ipc_stage", + "deterministic_collective_rocm_ipc_all_reduce", + "deterministic_collective_rocm_ipc_all_reduce_input", + "deterministic_collective_rocm_ipc_reduce_scatter", + "deterministic_collective_rocm_ipc_reduce_scatter_input", + "deterministic_collective_rocm_ipc_reduce_scatter_many", + "deterministic_collective_rocm_ipc_all_gather", + "deterministic_collective_rocm_ipc_all_gather_input", + ) + if any(not hasattr(_C, symbol) for symbol in required_symbols): + return + + staging = _C.deterministic_collective_rocm_ipc_allocate(self.max_size_bytes) + handle, offset = _C.deterministic_collective_rocm_ipc_meta(staging) + local_metadata = (socket.gethostname(), handle, int(offset)) + gathered_metadata: list[tuple[str, list[int], int] | None] = [None] * self.world_size + dist.all_gather_object(gathered_metadata, local_metadata, group=self.group) + if any(metadata is None for metadata in gathered_metadata): + raise RuntimeError("failed to exchange ROCm IPC metadata") + complete_metadata = [metadata for metadata in gathered_metadata if metadata is not None] + if len({metadata[0] for metadata in complete_metadata}) != 1: + return + self._ipc_handle = int( + _C.deterministic_collective_rocm_ipc_create( + staging, + [metadata[1] for metadata in complete_metadata], + [metadata[2] for metadata in complete_metadata], + self.rank, + ) + ) + self._ipc_staging = staging + + def _direct_all_reduce(self, input: torch.Tensor, output: torch.Tensor) -> bool: + handle = self._ipc_handle + if not handle: + return False + from rl_engine import _C + + input_bytes = input.numel() * input.element_size() + if ( + _ROCM_IPC_DIRECT_ALL_REDUCE_MAX_BYTES + < input_bytes + < _ROCM_IPC_SHARDED_ALL_REDUCE_MIN_BYTES + and input.numel() % self.world_size == 0 + ): + return False + + if ( + input_bytes <= _ROCM_IPC_DIRECT_ALL_REDUCE_MAX_BYTES + or input.numel() % self.world_size != 0 + ): + _C.deterministic_collective_rocm_ipc_all_reduce_input( + handle, + input, + output, + ) + return True + + shard = self._workspace_for(input, input.numel() // self.world_size) + _C.deterministic_collective_rocm_ipc_reduce_scatter_input( + handle, + input, + shard, + ) + dist.all_gather_into_tensor(output.view(-1), shard, group=self.group) + return True + + def _direct_reduce_scatter(self, input: torch.Tensor, output: torch.Tensor) -> bool: + handle = self._ipc_handle + if not handle: + return False + from rl_engine import _C + + _C.deterministic_collective_rocm_ipc_reduce_scatter_input( + handle, + input, + output, + ) + return True + + def _can_direct_reduce_scatter_many(self) -> bool: + return bool(self._ipc_handle) + + def _direct_reduce_scatter_many( + self, + inputs: tuple[torch.Tensor, ...], + outputs: tuple[torch.Tensor, ...], + ) -> bool: + handle = self._ipc_handle + if not handle: + return False + from rl_engine import _C + + _C.deterministic_collective_rocm_ipc_reduce_scatter_many( + handle, + inputs, + outputs, + ) + return True + + def _direct_all_gather(self, input: torch.Tensor, output: torch.Tensor) -> bool: + handle = self._ipc_handle + input_bytes = input.numel() * input.element_size() + if not handle or input_bytes > _ROCM_IPC_ALL_GATHER_MAX_BYTES: + return False + from rl_engine import _C + + _C.deterministic_collective_rocm_ipc_all_gather_input( + handle, + input, + output, + ) + return True + + def close(self) -> None: + handle = getattr(self, "_ipc_handle", 0) + if handle: + from rl_engine import _C + + _C.deterministic_collective_rocm_ipc_synchronize(handle) + torch.cuda.synchronize(self.device) + self._ipc_handle = 0 + _C.deterministic_collective_rocm_ipc_destroy(handle) + self._ipc_staging = None + super().close() diff --git a/rl_engine/integrations/framework_operators.py b/rl_engine/integrations/framework_operators.py index 3f19816b..ac981447 100644 --- a/rl_engine/integrations/framework_operators.py +++ b/rl_engine/integrations/framework_operators.py @@ -12,8 +12,9 @@ import os from dataclasses import replace +from functools import lru_cache from threading import Lock -from typing import Any, Mapping, cast +from typing import Any, Callable, Mapping, cast import torch @@ -32,10 +33,11 @@ from rl_engine.kernels.attention_contract import ReductionSpec as AttentionReductionSpec from rl_engine.kernels.attention_contract import ShardingSpec as AttentionShardingSpec from rl_engine.kernels.attention_contract import SplitKVSpec +from rl_engine.kernels.attention_projection import ROCM_DETERMINISTIC_PROJECTION_BACKEND_ID from rl_engine.kernels.logprob_contract import LogprobContract, LogprobDType, LogprobRole, MaskSpec from rl_engine.kernels.logprob_contract import ReductionSpec as LogprobReductionSpec from rl_engine.kernels.logprob_contract import ShardingSpec as LogprobShardingSpec -from rl_engine.kernels.ops.cuda.matmul.det_gemm import det_gemm_backend_id +from rl_engine.kernels.ops.matmul.det_gemm import DetGemmOp, det_gemm_backend_id from rl_engine.kernels.ops.pytorch.attention.ablation import AttentionAblationConfig from rl_engine.kernels.ops.pytorch.loss.vocab_parallel_logp import BACKEND_ID as LOGP_BACKEND_ID from rl_engine.kernels.ops.pytorch.loss.vocab_parallel_logp import DEFAULT_NUM_VOCAB_TILES @@ -44,6 +46,10 @@ ATTENTION_BACKEND_ID = "rlkernel.attention.deterministic.v1" FFN_BACKEND_ID = "rlkernel.ffn.qwen3.deterministic.v1" +_MEGATRON_TP_QKV_DGRAD_COLLECTIVE_ATTR = "__rl_kernel_tp_qkv_dgrad_collective_backend__" +_MEGATRON_TP_OUTPUT_PROJECTION_COLLECTIVE_ATTR = ( + "__rl_kernel_tp_output_projection_collective_backend__" +) def _device_name(tensor: torch.Tensor) -> str: @@ -118,8 +124,8 @@ def _attention_dtype(tensor: torch.Tensor) -> AttentionDType: def _require_nvidia_cuda(tensor: torch.Tensor, module: str) -> None: - if tensor.device.type != "cuda" or torch.version.hip is not None: - raise RuntimeError(f"strict {module} R/R requires NVIDIA CUDA tensors") + if tensor.device.type != "cuda": + raise RuntimeError(f"strict {module} R/R requires CUDA/ROCm GPU tensors") def _require_attention_accelerator(tensor: torch.Tensor) -> str: @@ -146,6 +152,32 @@ def _strict_attention_platform_contract(platform: str) -> tuple[str, str, str]: raise RuntimeError(f"unsupported strict Attention platform {platform!r}") +def _strict_attention_projection_backend_id(platform: str) -> str: + if platform == "rocm": + return ROCM_DETERMINISTIC_PROJECTION_BACKEND_ID + if platform == "cuda": + return det_gemm_backend_id() + raise RuntimeError(f"unsupported Attention projection platform {platform!r}") + + +def _strict_attention_projection_provenance(platform: str) -> dict[str, Any]: + return { + "backend_id": _strict_attention_projection_backend_id(platform), + "deterministic": True, + "accumulation_dtype": "fp32", + "reduction_order": "k_ascending", + "split_k": False, + "roles": ["qkv", "o_proj"], + "triton_used": platform == "rocm", + } + + +def _strict_attention_projection_op() -> Any: + """Construct the deterministic projection selected for this PyTorch build.""" + + return DetGemmOp() + + class SemanticOperatorHandle: """Resolve one exact semantic backend once for one framework process.""" @@ -281,6 +313,7 @@ def _megatron_parallel_state() -> Any: return parallel_state +@lru_cache(maxsize=512) def _megatron_zigzag_layout( local_tokens: int, *, @@ -436,6 +469,40 @@ def __init__(self, handle: SemanticOperatorHandle | None = None) -> None: self._packed_layout_owner: Any | None = None self._packed_layout_key: tuple[Any, ...] | None = None self._packed_layout_value: tuple[tuple[int, ...], tuple[int, ...]] | None = None + self._position_ids_cache: dict[tuple[Any, ...], torch.Tensor] = {} + + @staticmethod + def _tp_collective_backend(module: Any, attribute: str, tp_world: int) -> str: + value = getattr(module, attribute, None) + if isinstance(value, str) and value.strip(): + return value.strip() + return "none" if tp_world == 1 else "unbound" + + def _position_ids( + self, + positions: tuple[int, ...], + *, + batch_size: int, + cp_rank: int, + cp_world_size: int, + device: torch.device, + ) -> torch.Tensor: + key = ( + len(positions), + int(batch_size), + int(cp_rank), + int(cp_world_size), + device.type, + device.index, + ) + cached = self._position_ids_cache.get(key) + if cached is not None: + return cached + if len(self._position_ids_cache) >= 128: + self._position_ids_cache.pop(next(iter(self._position_ids_cache))) + value = torch.tensor(positions, dtype=torch.int64, device=device).repeat(batch_size, 1) + self._position_ids_cache[key] = value + return value def _packed_layout( self, @@ -538,11 +605,13 @@ def execute_sequence( cp_rank=cp_rank, cp_world_size=cp_world, ) - position_ids = torch.tensor( + position_ids = self._position_ids( positions, - dtype=torch.int64, + batch_size=q_ready.size(0), + cp_rank=cp_rank, + cp_world_size=cp_world, device=q_ready.device, - ).repeat(q_ready.size(0), 1) + ) return operator( q_ready, k_ready, @@ -612,20 +681,21 @@ def execute_sequence( sequence_provenance: list[dict[str, Any] | None] = [None] * len(global_lengths) launch_group_count = 0 for (local_length, global_length), sequences in grouped_sequences.items(): - q_ready = ( - torch.stack([query[start:end] for _index, start, end in sequences], dim=0) - .permute(0, 2, 1, 3) - .contiguous() + # Stack the cheap [H, T, D] views directly into FA4's + # [B, H, T, D] layout. Stacking before permuting would first + # materialize [B, T, H, D] and then copy the entire tensor a + # second time for every layer and microbatch. + q_ready = torch.stack( + [query[start:end].permute(1, 0, 2) for _index, start, end in sequences], + dim=0, ) - k_ready = ( - torch.stack([key[start:end] for _index, start, end in sequences], dim=0) - .permute(0, 2, 1, 3) - .contiguous() + k_ready = torch.stack( + [key[start:end].permute(1, 0, 2) for _index, start, end in sequences], + dim=0, ) - v_ready = ( - torch.stack([value[start:end] for _index, start, end in sequences], dim=0) - .permute(0, 2, 1, 3) - .contiguous() + v_ready = torch.stack( + [value[start:end].permute(1, 0, 2) for _index, start, end in sequences], + dim=0, ) result = execute_sequence( q_ready, @@ -665,7 +735,20 @@ def execute_sequence( "cp_world_size": cp_world, "tp_world_size": tp_world, "runtime_platform": runtime_platform, - "triton_used": False, + "triton_used": runtime_platform == "rocm", + "deterministic_projection": _strict_attention_projection_provenance( + runtime_platform + ), + "tp_qkv_dgrad_collective": self._tp_collective_backend( + module, + _MEGATRON_TP_QKV_DGRAD_COLLECTIVE_ATTR, + tp_world, + ), + "tp_output_projection_collective": self._tp_collective_backend( + module, + _MEGATRON_TP_OUTPUT_PROJECTION_COLLECTIVE_ATTR, + tp_world, + ), **execution_provenance, } return output @@ -750,12 +833,19 @@ def __call__( "framework_layout": "megatron_sequence_parallel", "cp_world_size": cp_world, "tp_world_size": tp_world, - "runtime_platform": "cuda", - "actual_backend": "rlkernel.cuda.det_gemm_swiglu", + "runtime_platform": _device_name(hidden_states), + "actual_backend": "rlkernel.rocm.det_gemm_swiglu" if torch.version.hip is not None else "rlkernel.cuda.det_gemm_swiglu", "gemm_backend": det_gemm_backend_id(), "fallback": False, "gate_up_projection": "separate_strict_launches", - "triton_used": False, + "deterministic_all_reduce_backend": ( + "none" + if tp_world == 1 + else "rocm_ipc_fixed_tree" + if torch.version.hip is not None + else "deterministic_all_reduce.ipc_localized_fixed_tree.v1" + ), + "triton_used": torch.version.hip is not None, } return output, None @@ -798,6 +888,10 @@ def normalize(plane: torch.Tensor) -> torch.Tensor: "vLLM K/V cache layout does not expose the declared number of KV heads" ) + # Match vLLM's native AITER layout before the legacy flattened layout. + # With two KV heads both layouts otherwise look like [B, 2, N, 2 * head]. + if kv_cache.ndim == 4 and kv_cache.size(-1) == 2 * head_size: + return kv_cache.transpose(1, 2).split(head_size, dim=-1) if ( kv_cache.ndim == 4 and platform == "rocm" @@ -811,8 +905,6 @@ def normalize(plane: torch.Tensor) -> torch.Tensor: key_cache.view(blocks, block_size, num_kv_heads, head_size), value_cache.view(blocks, block_size, num_kv_heads, head_size), ) - if kv_cache.ndim == 4 and kv_cache.size(-1) == 2 * head_size: - return kv_cache.transpose(1, 2).split(head_size, dim=-1) # The K/V axis is leading in CUDA FlashAttention caches and follows the # block axis in the ROCm AITER layouts. Prefer the platform convention in # the ambiguous two-block case where both dimensions happen to equal two. @@ -847,10 +939,16 @@ class VllmAttentionOperator: backend_id = ATTENTION_BACKEND_ID - def __init__(self, handle: SemanticOperatorHandle | None = None) -> None: + def __init__( + self, + handle: SemanticOperatorHandle | None = None, + *, + projection_collective_backend: Callable[[], str | None] | None = None, + ) -> None: self._handle = handle or SemanticOperatorHandle( target="rollout", semantic_op="attention", backend_id=self.backend_id ) + self._projection_collective_backend = projection_collective_backend self._last_provenance: dict[str, Any] = {} self._tp_coordinates: tuple[int, int, Any] | None = None self._metadata_cache_key: tuple[Any, ...] | None = None @@ -917,24 +1015,6 @@ def _materialization_groups( "metadata_source": "vllm_gpu", } - query_starts = query_starts_source.to(device=query.device, dtype=torch.int32) - seq_lens = seq_lens_source.to(device=query.device, dtype=torch.int32) - query_indices = torch.arange(num_actual, dtype=torch.int32, device=query.device) - query_ends = query_starts[1:] - request_indices = torch.searchsorted(query_ends, query_indices, right=True).to( - dtype=torch.long - ) - active_queries = query_indices < query_starts[-1] - request_indices = request_indices.clamp_max(seq_lens.numel() - 1) - request_query_ends = query_ends.index_select(0, request_indices) - request_seq_lens = seq_lens.index_select(0, request_indices) - seqused_k = request_seq_lens - (request_query_ends - query_indices) + 1 - seqused_k = torch.where( - active_queries, - seqused_k, - torch.ones_like(seqused_k), - ) - max_seq_len = int(getattr(attn_metadata, "max_seq_len", block_table.size(1) * block_size)) page_count = min(block_table.size(1), (max_seq_len + block_size - 1) // block_size) cache_key = ( @@ -948,6 +1028,8 @@ def _materialization_groups( page_count, ) owner_id = id(cache_owner) if cache_owner is not None else None + # Resolve the cross-layer cache before scheduling any GPU metadata work. + # Repeating an owner still forces the first layer of the next forward to miss. if ( owner_id is not None and self._metadata_cache_key == cache_key @@ -958,6 +1040,24 @@ def _materialization_groups( groups, cached_summary = self._metadata_cache_value return groups, {**cached_summary, "metadata_reused_across_layers": True} + query_starts = query_starts_source.to(device=query.device, dtype=torch.int32) + seq_lens = seq_lens_source.to(device=query.device, dtype=torch.int32) + query_indices = torch.arange(num_actual, dtype=torch.int32, device=query.device) + query_ends = query_starts[1:] + request_indices = torch.searchsorted(query_ends, query_indices, right=True).to( + dtype=torch.long + ) + active_queries = query_indices < query_starts[-1] + request_indices = request_indices.clamp_max(seq_lens.numel() - 1) + request_query_ends = query_ends.index_select(0, request_indices) + request_seq_lens = seq_lens.index_select(0, request_indices) + seqused_k = request_seq_lens - (request_query_ends - query_indices) + 1 + seqused_k = torch.where( + active_queries, + seqused_k, + torch.ones_like(seqused_k), + ) + pages = ( block_table.index_select(0, request_indices)[:, :page_count] .to(dtype=torch.int32) @@ -1094,6 +1194,13 @@ def __call__( else: output_heads.index_copy_(0, query_indices, result.out.squeeze(2)) last_operator_provenance = _compact_attention_provenance(result.provenance) + projection_collective_backend = "none" + if runtime_platform == "rocm" and tp_world > 1: + projection_collective_backend = "unbound" + if self._projection_collective_backend is not None: + projection_collective_backend = ( + self._projection_collective_backend() or "unbound" + ) self._last_provenance = { "framework_layout": "vllm_paged_kv", "materialization": ( @@ -1104,7 +1211,11 @@ def __call__( "tp_world_size": tp_world, "tp_group_bound": tp_group is not None, "runtime_platform": runtime_platform, - "triton_used": False, + "triton_used": runtime_platform == "rocm", + "deterministic_projection": _strict_attention_projection_provenance( + runtime_platform + ), + "deterministic_all_reduce_backend": projection_collective_backend, "direct_output_buffer": direct_output_buffer, **metadata_summary, "operator": last_operator_provenance, @@ -1171,16 +1282,17 @@ def bind_packed_inference(self, module: Any) -> tuple[int, int]: def _set_runtime_provenance( self, tp_world: int, deterministic_all_reduce_backend: str = "unbound" ) -> None: + runtime_platform = "rocm" if torch.version.hip is not None else "cuda" self._last_provenance = { "framework_layout": "vllm_tensor_parallel", "tp_world_size": tp_world, - "runtime_platform": "cuda", - "actual_backend": "rlkernel.cuda.det_gemm_swiglu", + "runtime_platform": runtime_platform, + "actual_backend": f"rlkernel.{runtime_platform}.det_gemm_swiglu", "gemm_backend": det_gemm_backend_id(), "fallback": False, "gate_up_projection": "packed_single_launch", "deterministic_all_reduce_backend": deterministic_all_reduce_backend, - "triton_used": False, + "triton_used": runtime_platform == "rocm", } def __call__(self, module: Any, hidden_states: torch.Tensor) -> torch.Tensor: @@ -1230,7 +1342,7 @@ def __call__(self, module: Any, hidden_states: torch.Tensor) -> torch.Tensor: class MegatronLogpOperator: - """Require CUDA while reusing the structural Vime Logp provider.""" + """Route structural Vime logp requests through the strict GPU backend.""" def __init__( self, @@ -1269,8 +1381,8 @@ def __call__(self, request: Any) -> Any: "operator": self.backend_id, "actual_backend": self.backend_id, "fallback": False, - "runtime_platform": "cuda", - "triton_used": False, + "runtime_platform": _device_name(hidden), + "triton_used": torch.version.hip is not None, "provider": dict(getattr(result, "provenance", {})), "linear_logp": dict(self._linear_logp.provenance), "logits_materialized": False, @@ -1289,8 +1401,8 @@ def __call__(self, request: Any) -> Any: "operator": self.backend_id, "actual_backend": self.backend_id, "fallback": False, - "runtime_platform": "cuda", - "triton_used": False, + "runtime_platform": _device_name(logits), + "triton_used": torch.version.hip is not None, "provider": dict(getattr(result, "provenance", {})), } return result @@ -1533,13 +1645,20 @@ def __call__( real_vocab_size=context.real_vocab_size, temperature=float(os.getenv("RL_KERNEL_VLLM_TEMPERATURE", "1.0")), target="rollout", + diagnostics_hidden=context.hidden, + diagnostics_lm_head_weight=context.lm_head_weight, ) strict_provenance = self._linear_logp.provenance + expected_entrypoint = ( + "rocm_vocab_parallel_logp_from_local_logits_tp" + if torch.version.hip is not None + else "sm90_deterministic_logp_from_local_logits_tp" + ) if ( strict_provenance.get("deterministic_linear_logp") is not True or strict_provenance.get("actual_backend") != self._linear_logp.backend_id or strict_provenance.get("strict_entrypoint") - != "sm90_deterministic_logp_from_local_logits_tp" + != expected_entrypoint ): raise RuntimeError( "strict vLLM rollout linear_logp did not execute the deterministic " @@ -1547,8 +1666,8 @@ def __call__( ) provenance = { **dict(strict_provenance), - "runtime_platform": "cuda", - "triton_used": False, + "runtime_platform": _device_name(local_logits), + "triton_used": torch.version.hip is not None, "execution": { "role": "vllm_rollout_linear_logprob", "strict_backend": True, @@ -1660,8 +1779,8 @@ def __call__( provenance = { **dict(dispatch.provenance), - "runtime_platform": "cuda", - "triton_used": False, + "runtime_platform": _device_name(source_logits), + "triton_used": torch.version.hip is not None, "source_logits_shape": list(source_logits.shape), "source_logits_dtype": _dtype_name(source_logits), "contract_real_vocab_size": contract.sharding.real_vocab_size, diff --git a/rl_engine/integrations/linear_logp.py b/rl_engine/integrations/linear_logp.py index 55d3892d..667f989f 100644 --- a/rl_engine/integrations/linear_logp.py +++ b/rl_engine/integrations/linear_logp.py @@ -1,6 +1,8 @@ from __future__ import annotations import os +import pathlib +import threading from collections.abc import Mapping from dataclasses import dataclass from typing import Any @@ -8,6 +10,138 @@ import torch from rl_engine.integrations.ablation import operator_ablation_case +from rl_engine.kernels.logprob_contract import ( + LogprobContract, + LogprobDType, + LogprobRole, + MaskSpec, + ReductionSpec, + ShardingSpec, +) +from rl_engine.kernels.ops.pytorch.loss.vocab_parallel_logp import DEFAULT_NUM_VOCAB_TILES + + +_ALIGNMENT_DIAGNOSTIC_LOCK = threading.Lock() +_ALIGNMENT_DIAGNOSTIC_CALLS = 0 + + +def _alignment_diagnostics_enabled() -> bool: + return os.getenv("RL_KERNEL_ALIGNMENT_DIAGNOSTICS", "").strip().lower() in { + "1", + "true", + "yes", + "on", + } + + +def _alignment_diagnostic_rank() -> int: + try: + if torch.distributed.is_available() and torch.distributed.is_initialized(): + return int(torch.distributed.get_rank()) + except (RuntimeError, ValueError): + pass + try: + return int(os.getenv("RANK", "0")) + except ValueError: + return 0 + + +def _alignment_tensor_summary(value: torch.Tensor) -> dict[str, Any]: + detached = value.detach() + flat = detached.reshape(-1) + first = flat[:16].to(device="cpu") + last = flat[-16:].to(device="cpu") if flat.numel() else first + summary: dict[str, Any] = { + "shape": list(detached.shape), + "dtype": str(detached.dtype).replace("torch.", ""), + "device": str(detached.device), + "first16": first.tolist(), + "last16": last.tolist(), + } + if detached.numel(): + values = detached.float() + summary.update( + { + "sum_fp32": float(values.sum().item()), + "abs_sum_fp32": float(values.abs().sum().item()), + "min_fp32": float(values.min().item()), + "max_fp32": float(values.max().item()), + } + ) + return summary + + +def _record_alignment_diagnostics( + *, + target: str, + local_logits: torch.Tensor, + target_ids: torch.Tensor, + result: torch.Tensor, + lse: torch.Tensor, + vocab_start_index: int, + hidden: torch.Tensor | None = None, + lm_head_weight: torch.Tensor | None = None, +) -> None: + """Persist bounded row diagnostics only when explicitly enabled. + + This branch performs device-to-host copies and reductions for debugging; + production runs leave it cold and do no additional tensor work. + """ + + if not _alignment_diagnostics_enabled(): + return + root = os.getenv("RL_KERNEL_ALIGNMENT_DIAGNOSTICS_DIR", "").strip() + if not root: + return + global _ALIGNMENT_DIAGNOSTIC_CALLS + with _ALIGNMENT_DIAGNOSTIC_LOCK: + call_index = _ALIGNMENT_DIAGNOSTIC_CALLS + _ALIGNMENT_DIAGNOSTIC_CALLS += 1 + detached_logits = local_logits.detach() + ids = target_ids.detach().to(device="cpu", dtype=torch.int64) + local_ids = target_ids.detach().to(dtype=torch.long) - int(vocab_start_index) + local_mask = (local_ids >= 0) & (local_ids < detached_logits.size(1)) + safe_ids = local_ids.clamp(0, max(int(detached_logits.size(1)) - 1, 0)) + selected_local = detached_logits.gather(1, safe_ids.reshape(-1, 1)).reshape(-1) + selected_local = torch.where( + local_mask, + selected_local, + torch.full_like(selected_local, float("nan")), + ) + row_sum = detached_logits.float().sum(dim=1) + row_abs_sum = detached_logits.float().abs().sum(dim=1) + payload = { + "schema_version": "rlkernel.linear_logp_alignment_diagnostic.v1", + "target": target, + "rank": _alignment_diagnostic_rank(), + "pid": os.getpid(), + "call_index": call_index, + "vocab_start_index": int(vocab_start_index), + "target_ids": ids, + "selected_local_logits": selected_local.cpu(), + "result": result.detach().cpu(), + "lse": lse.detach().cpu(), + "row_sum_fp32": row_sum.cpu(), + "row_abs_sum_fp32": row_abs_sum.cpu(), + "logits_head": detached_logits[:, :8].cpu(), + "logits_tail": detached_logits[:, -8:].cpu(), + "logits_summary": _alignment_tensor_summary(detached_logits), + } + if hidden is not None: + payload["hidden"] = hidden.detach().cpu() + payload["hidden_summary"] = _alignment_tensor_summary(hidden) + if lm_head_weight is not None: + detached_weight = lm_head_weight.detach() + payload["lm_head_weight_head"] = detached_weight[:8].cpu() + payload["lm_head_weight_tail"] = detached_weight[-8:].cpu() + payload["lm_head_weight_summary"] = _alignment_tensor_summary(detached_weight) + output_dir = pathlib.Path(root) + output_dir.mkdir(parents=True, exist_ok=True) + path = output_dir / ( + f"{target}-pid{os.getpid()}-rank{_alignment_diagnostic_rank():05d}-" + f"call{call_index:08d}.pt" + ) + torch.save(payload, path) @dataclass(frozen=True) @@ -101,6 +235,8 @@ class LinearLogpWrapper: def __init__(self) -> None: self._op: Any | None = None self._tp_op: Any | None = None + self._rocm_op: Any | None = None + self._rocm_linear: Any | None = None self._last_provenance: dict[str, Any] = {} @property @@ -143,6 +279,22 @@ def _resolve(self, *, tensor_parallel: bool) -> Any: self._op = op return op + def _resolve_rocm(self) -> Any: + if self._rocm_op is None: + from rl_engine.kernels.ops.rocm.loss.vocab_parallel_logp import ( + RocmVocabParallelLogprobOp, + ) + + self._rocm_op = RocmVocabParallelLogprobOp() + return self._rocm_op + + def _resolve_rocm_linear(self) -> Any: + if self._rocm_linear is None: + from rl_engine.kernels.ops.matmul.det_gemm import DetGemmOp + + self._rocm_linear = DetGemmOp() + return self._rocm_linear + @staticmethod def _tp_coordinates(tp_group: Any) -> tuple[int, int]: if tp_group is None: @@ -219,12 +371,12 @@ def _validate_contract( raise ValueError( f"linear_logp wrapper expects hidden [tokens, hidden], got {tuple(hidden.shape)}" ) - if not hidden.is_cuda or torch.version.hip is not None: - raise RuntimeError("strict linear_logp requires NVIDIA CUDA tensors") + if not hidden.is_cuda: + raise RuntimeError("strict linear_logp requires CUDA/ROCm GPU tensors") if lm_head_weight.ndim != 2: raise ValueError("linear_logp LM-head weight must be [vocab_local, hidden]") if hidden.dtype != torch.bfloat16 or lm_head_weight.dtype != torch.bfloat16: - raise TypeError("strict SM90 linear_logp requires bfloat16 hidden and LM-head") + raise TypeError("strict linear_logp requires bfloat16 hidden and LM-head") if lm_head_weight.device != hidden.device: raise ValueError("linear_logp hidden and LM-head must share a device") if hidden.size(1) != lm_head_weight.size(1): @@ -268,6 +420,108 @@ def _validate_contract( cls._validate_targets(target_ids, rows=hidden.size(0), real_vocab_size=real) return tensor_parallel, requested_global, real + @staticmethod + def _rocm_contract( + local_logits: torch.Tensor, + *, + rank: int, + world: int, + global_vocab_size: int, + real_vocab_size: int, + target: str, + ) -> LogprobContract: + local_vocab = int(local_logits.size(1)) + dtype = { + torch.bfloat16: LogprobDType.BF16, + torch.float16: LogprobDType.FP16, + torch.float32: LogprobDType.FP32, + }.get(local_logits.dtype) + if dtype is None: + raise TypeError(f"unsupported ROCm linear_logp dtype {local_logits.dtype}") + role = LogprobRole.TRAIN if target == "training" else LogprobRole.INFER + return LogprobContract( + role=role, + dtype=dtype, + mask=MaskSpec( + num_tokens=int(local_logits.size(0)), + active_mask=(True,) * int(local_logits.size(0)), + ), + sharding=ShardingSpec( + tp_rank=rank, + tp_world_size=world, + vocab_shard_bounds=tuple( + (index * local_vocab, (index + 1) * local_vocab) + for index in range(world) + ), + real_vocab_size=real_vocab_size, + padded_vocab_size=global_vocab_size, + ), + reduction=ReductionSpec(), + ) + + def _rocm_from_local_logits( + self, + local_logits: torch.Tensor, + target_ids: torch.Tensor, + *, + tp_group: Any, + rank: int, + world: int, + global_vocab_size: int, + real_vocab_size: int, + target: str, + temperature: float | torch.Tensor | None, + diagnostics_hidden: torch.Tensor | None = None, + diagnostics_lm_head_weight: torch.Tensor | None = None, + ) -> tuple[torch.Tensor, torch.Tensor, str]: + if temperature is None or ( + not isinstance(temperature, torch.Tensor) and float(temperature) == 1.0 + ): + effective_logits = local_logits.contiguous() + # Megatron exposes the materialized TP LM-head through a squeezed + # view. The loss path may reuse that storage while autograd still + # needs it for the logp backward. Keep one gradient-preserving + # snapshot for the strict training kernel; rollout is inference + # only and remains zero-copy. + if target == "training" and effective_logits.requires_grad: + effective_logits = effective_logits.clone() + else: + temperature_tensor = self._temperature_tensor( + temperature, + rows=local_logits.size(0), + device=local_logits.device, + ) + assert temperature_tensor is not None + effective_logits = local_logits.float() / temperature_tensor.unsqueeze(1) + contract = self._rocm_contract( + effective_logits, + rank=rank, + world=world, + global_vocab_size=global_vocab_size, + real_vocab_size=real_vocab_size, + target=target, + ) + op = self._resolve_rocm() + result, lse = op.apply( + effective_logits, + target_ids, + contract=contract, + tp_group=tp_group, + num_vocab_tiles=DEFAULT_NUM_VOCAB_TILES, + deterministic=True, + ) + _record_alignment_diagnostics( + target=target, + local_logits=effective_logits, + target_ids=target_ids, + result=result, + lse=lse, + vocab_start_index=rank * int(effective_logits.size(1)), + hidden=diagnostics_hidden, + lm_head_weight=diagnostics_lm_head_weight, + ) + return result, lse, str(op.backend_id) + def from_local_logits( self, local_logits: torch.Tensor, @@ -279,18 +533,23 @@ def from_local_logits( real_vocab_size: int, target: str = "rollout", temperature: float | torch.Tensor | None = None, - ) -> torch.Tensor: + return_lse: bool = False, + diagnostics_hidden: torch.Tensor | None = None, + diagnostics_lm_head_weight: torch.Tensor | None = None, + ) -> torch.Tensor | tuple[torch.Tensor, torch.Tensor]: """Score deterministic local LM-head logits without a duplicate GEMM.""" - if local_logits.ndim != 2 or local_logits.dtype != torch.bfloat16: - raise TypeError("strict reused LM-head logits must be 2-D bfloat16") - if not local_logits.is_cuda or torch.version.hip is not None: - raise RuntimeError("strict reused LM-head logits require NVIDIA CUDA") + if local_logits.ndim != 2 or local_logits.dtype not in ( + torch.bfloat16, + torch.float16, + torch.float32, + ): + raise TypeError("strict reused LM-head logits must be 2-D bf16/fp16/fp32") + if not local_logits.is_cuda: + raise RuntimeError("strict reused LM-head logits require CUDA/ROCm") if target_ids.device != local_logits.device: raise ValueError("linear_logp target_ids must share the logits device") rank, world = self._tp_coordinates(tp_group) - if tp_group is None or world <= 1: - raise ValueError("reused rollout LM-head logits require a multi-rank TP group") local_vocab = int(local_logits.size(1)) requested_global = int(global_vocab_size) expected_global = local_vocab * world @@ -305,32 +564,59 @@ def from_local_logits( rows=local_logits.size(0), real_vocab_size=real, ) - temperature_tensor = self._temperature_tensor( - temperature, - rows=local_logits.size(0), - device=local_logits.device, - ) - from rl_engine.kernels.ops.cuda.loss.linear_logp import ( - sm90_deterministic_logp_from_local_logits_tp, - ) + if torch.version.hip is not None: + result, lse, kernel_backend = self._rocm_from_local_logits( + local_logits, + target_ids, + tp_group=tp_group, + rank=rank, + world=world, + global_vocab_size=requested_global, + real_vocab_size=real, + target=target, + temperature=temperature, + diagnostics_hidden=diagnostics_hidden, + diagnostics_lm_head_weight=diagnostics_lm_head_weight, + ) + runtime_platform = "rocm" + strict_entrypoint = ( + "rocm_deterministic_linear_logp_tp" + if target == "training" + else "rocm_vocab_parallel_logp_from_local_logits_tp" + ) + contract_version = "rocm-det-gemm-vocab-parallel-logp-ws2-v1" + else: + temperature_tensor = self._temperature_tensor( + temperature, + rows=local_logits.size(0), + device=local_logits.device, + ) + from rl_engine.kernels.ops.cuda.loss.linear_logp import ( + sm90_deterministic_logp_from_local_logits_tp, + ) - result, _lse = sm90_deterministic_logp_from_local_logits_tp( - local_logits.contiguous(), - target_ids, - tp_group=tp_group, - vocab_start_index=int(vocab_start_index), - global_vocab_size=requested_global, - real_vocab_size=real, - temperature=temperature_tensor, - ) + result, lse = sm90_deterministic_logp_from_local_logits_tp( + local_logits.contiguous(), + target_ids, + tp_group=tp_group, + vocab_start_index=int(vocab_start_index), + global_vocab_size=requested_global, + real_vocab_size=real, + temperature=temperature_tensor, + ) + kernel_backend = self.backend_id + runtime_platform = "cuda" + strict_entrypoint = "sm90_deterministic_logp_from_local_logits_tp" + contract_version = "cuda-det-gemm-linear-logp-sm90-contract-v2" self._last_provenance = { **self._mismatch_provenance(), "target": target, - "runtime_platform": "cuda", - "triton_used": False, + "runtime_platform": runtime_platform, + "triton_used": runtime_platform == "rocm", "actual_backend": self.backend_id, + "logprob_kernel_backend": kernel_backend, "deterministic_linear_logp": True, - "strict_entrypoint": "sm90_deterministic_logp_from_local_logits_tp", + "strict_entrypoint": strict_entrypoint, "local_logits_shape": list(local_logits.shape), "target_shape": list(target_ids.shape), "tp_group_present": True, @@ -338,11 +624,11 @@ def from_local_logits( "global_vocab_size": requested_global, "real_vocab_size": real, "temperature": None if temperature is None else "provided", - "contract_version": "cuda-det-gemm-linear-logp-sm90-contract-v2", + "contract_version": contract_version, "logits_materialized": True, "lm_head_result_reused": True, } - return result + return (result, lse) if return_lse else result @staticmethod def _mismatch_provenance() -> dict[str, Any]: @@ -400,31 +686,60 @@ def __call__( ) effective_weight = lm_head_weight effective_bias = bias - temperature_tensor = self._temperature_tensor( - temperature, rows=hidden.size(0), device=hidden.device - ) - op = self._resolve(tensor_parallel=tensor_parallel) - if tensor_parallel: - result, _lse = op( - hidden, - effective_weight, + if torch.version.hip is not None: + local_logits = self._resolve_rocm_linear().linear(hidden, effective_weight) + if effective_bias is not None: + local_logits = local_logits + effective_bias + result, _lse, kernel_backend = self._rocm_from_local_logits( + local_logits, target_ids, - effective_bias, tp_group=tp_group, - vocab_start_index=int(vocab_start_index), + rank=self._tp_coordinates(tp_group)[0], + world=self._tp_coordinates(tp_group)[1], global_vocab_size=requested_global_vocab, real_vocab_size=real, - temperature=temperature_tensor, + target=target, + temperature=temperature, + diagnostics_hidden=hidden, + diagnostics_lm_head_weight=effective_weight, ) + runtime_platform = "rocm" + strict_entrypoint = "rocm_deterministic_linear_logp_tp" + contract_version = "rocm-det-gemm-vocab-parallel-logp-ws2-v1" else: - result, _lse = op( - hidden, - effective_weight, - target_ids, - effective_bias, - real_vocab_size=real, - temperature=temperature_tensor, + temperature_tensor = self._temperature_tensor( + temperature, rows=hidden.size(0), device=hidden.device ) + op = self._resolve(tensor_parallel=tensor_parallel) + if tensor_parallel: + result, _lse = op( + hidden, + effective_weight, + target_ids, + effective_bias, + tp_group=tp_group, + vocab_start_index=int(vocab_start_index), + global_vocab_size=requested_global_vocab, + real_vocab_size=real, + temperature=temperature_tensor, + ) + else: + result, _lse = op( + hidden, + effective_weight, + target_ids, + effective_bias, + real_vocab_size=real, + temperature=temperature_tensor, + ) + kernel_backend = self.backend_id + runtime_platform = "cuda" + strict_entrypoint = ( + "sm90_deterministic_linear_logp_tp" + if tensor_parallel + else "sm90_deterministic_linear_logp" + ) + contract_version = "cuda-det-gemm-linear-logp-sm90-contract-v2" if not isinstance(result, torch.Tensor): raise RuntimeError("linear_logp backend returned a non-tensor result") if result.numel() != hidden.size(0): @@ -437,15 +752,12 @@ def __call__( self._last_provenance = { **self._mismatch_provenance(), "target": target, - "runtime_platform": "cuda", - "triton_used": False, + "runtime_platform": runtime_platform, + "triton_used": runtime_platform == "rocm", "actual_backend": actual_backend, + "logprob_kernel_backend": kernel_backend, "deterministic_linear_logp": True, - "strict_entrypoint": ( - "sm90_deterministic_linear_logp_tp" - if tensor_parallel - else "sm90_deterministic_linear_logp" - ), + "strict_entrypoint": strict_entrypoint, "hidden_shape": list(hidden.shape), "hidden_dtype": str(hidden.dtype).replace("torch.", ""), "lm_head_weight_shape": list(lm_head_weight.shape), @@ -457,7 +769,7 @@ def __call__( "real_vocab_size": real, "requested_real_vocab_size": None if real_vocab_size is None else int(real_vocab_size), "temperature": None if temperature is None else "provided", - "contract_version": "cuda-det-gemm-linear-logp-sm90-contract-v2", + "contract_version": contract_version, "logits_materialized": True, } return result diff --git a/rl_engine/integrations/megatron_runtime.py b/rl_engine/integrations/megatron_runtime.py index 72ac8919..596b0645 100644 --- a/rl_engine/integrations/megatron_runtime.py +++ b/rl_engine/integrations/megatron_runtime.py @@ -6,7 +6,9 @@ from __future__ import annotations import importlib +import io import os +import pathlib from collections.abc import Callable, Iterable from types import MethodType from typing import Any @@ -15,10 +17,13 @@ from rl_engine.integrations.ablation import Implementation, IntegrationPlan from rl_engine.integrations.framework_operators import ( + _MEGATRON_TP_OUTPUT_PROJECTION_COLLECTIVE_ATTR, + _MEGATRON_TP_QKV_DGRAD_COLLECTIVE_ATTR, MegatronAttentionOperator, MegatronFFNOperator, MegatronLogpOperator, _fused_rms_norm_input, + _strict_attention_projection_op, ) from rl_engine.integrations.linear_logp import LinearLogpWrapper from rl_engine.integrations.megatron import MegatronIntegration @@ -28,7 +33,530 @@ _PATCH_MARKER = "__rl_kernel_original_forward__" _STRICT_ATTENTION_PATCH_MARKER = "__rl_kernel_original_strict_attention_init__" _STRICT_ATTENTION_PROJECTION_MARKER = "__rl_kernel_strict_attention_projection__" +_STRICT_ATTENTION_CORE_MARKER = "__rl_kernel_strict_attention_core__" _STRICT_TE_RMS_NORM_PATCH_MARKER = "__rl_kernel_original_strict_rms_norm_forward__" +_STRICT_LOGP_OUTPUT_PATCH_MARKER = "__rl_kernel_original_strict_logp_output_forward__" +_STRICT_LOGP_REUSABLE_MARKER = "__rl_kernel_reusable_local_logits__" +_STRICT_ROCM_ROPE_PATCH_MARKER = "__rl_kernel_original_rocm_rope_apply__" +_STRICT_LAYER_DIAGNOSTIC_PATCH_MARKER = "__rl_kernel_original_layer_diagnostic_forward__" + + +def _alignment_diagnostics_enabled() -> bool: + value = os.getenv( + "RL_KERNEL_LAYER_ALIGNMENT_DIAGNOSTICS", + os.getenv("RL_KERNEL_ALIGNMENT_DIAGNOSTICS", ""), + ) + return value.strip().lower() in { + "1", + "true", + "yes", + "on", + } + + +def _diagnostic_rank() -> int: + if torch.distributed.is_available() and torch.distributed.is_initialized(): + return int(torch.distributed.get_rank()) + return int(os.getenv("RANK", "0")) + + +def _diagnostic_rows(row_count: int) -> torch.Tensor: + requested = os.getenv("RL_KERNEL_ALIGNMENT_ROWS", "").strip() + rows: list[int] = [] + for part in requested.split(",") if requested else (): + start_text, separator, end_text = part.strip().partition("-") + try: + start = int(start_text) + end = int(end_text) if separator else start + except ValueError: + continue + rows.extend(range(max(start, 0), min(end + 1, row_count))) + if not rows: + rows = list(range(max(0, row_count - min(row_count, 64)), row_count)) + return torch.tensor(sorted(set(rows)), dtype=torch.long) + + +def _save_megatron_layer_diagnostic( + instance: Any, + input_value: torch.Tensor, + output_value: torch.Tensor, + intermediates: dict[str, torch.Tensor] | None = None, +) -> None: + if not _alignment_diagnostics_enabled(): + return + root = os.getenv("RL_KERNEL_ALIGNMENT_DIAGNOSTICS_DIR", "").strip() + if not root: + return + layer = int(getattr(instance, "layer_number", 0)) - 1 + requested_layers = os.getenv("RL_KERNEL_ALIGNMENT_LAYERS", "").strip() + if requested_layers: + try: + enabled_layers = {int(value.strip()) for value in requested_layers.split(",")} + except ValueError as exc: + raise RuntimeError("RL_KERNEL_ALIGNMENT_LAYERS must be comma-separated integers") from exc + if layer not in enabled_layers: + return + if input_value.ndim == 3: + if input_value.size(1) != 1 or output_value.size(1) != 1: + raise RuntimeError("Megatron alignment diagnostics require batch size one") + input_rows = input_value[:, 0] + output_rows = output_value[:, 0] + elif input_value.ndim == 2: + input_rows = input_value + output_rows = output_value + else: + raise RuntimeError("Megatron layer diagnostics require [S,B,H] or [T,H]") + call_index = int(getattr(instance, "__rl_kernel_layer_diagnostic_call__", 0)) + setattr(instance, "__rl_kernel_layer_diagnostic_call__", call_index + 1) + indices_cpu = _diagnostic_rows(int(input_rows.size(0))) + indices = indices_cpu.to(device=input_rows.device) + rank = _diagnostic_rank() + payload = { + "schema_version": "rlkernel.layer_alignment_diagnostic.v1", + "framework": "megatron", + "rank": rank, + "layer": layer, + "call_index": call_index, + "row_indices": indices_cpu, + "input": input_rows.detach().index_select(0, indices).cpu(), + "output": output_rows.detach().index_select(0, indices).cpu(), + } + for name, value in (intermediates or {}).items(): + rows = value.detach() + if rows.ndim >= 3 and rows.size(1) == 1: + rows = rows[:, 0] + if rows.ndim < 2 or rows.size(0) != input_rows.size(0): + raise RuntimeError( + f"Megatron layer diagnostic {name} does not expose the token axis first" + ) + payload[name] = rows.index_select(0, indices).cpu() + output_dir = pathlib.Path(root) / "layers" + output_dir.mkdir(parents=True, exist_ok=True) + torch.save( + payload, + output_dir + / ( + f"megatron-pid{os.getpid()}-rank{rank:05d}-layer{layer:02d}-" + f"call{call_index:08d}.pt" + ), + ) + + +def _patch_layer_alignment_diagnostics() -> None: + if not _alignment_diagnostics_enabled(): + return + from megatron.core.transformer.transformer_layer import TransformerLayer + + if hasattr(TransformerLayer, _STRICT_LAYER_DIAGNOSTIC_PATCH_MARKER): + return + original = TransformerLayer.forward + + def wrapped(instance: Any, *args: Any, **kwargs: Any) -> Any: + input_value = args[0] if args else kwargs.get("hidden_states") + intermediates: dict[str, torch.Tensor] = {} + handles: list[Any] = [] + + def tensor_result(value: Any) -> torch.Tensor | None: + candidate = value[0] if isinstance(value, tuple) else value + return candidate if isinstance(candidate, torch.Tensor) else None + + if int(getattr(instance, "layer_number", 0)) == 1: + attention = instance.self_attention + qkv_projection = attention.linear_qkv + core_attention = attention.core_attention + mlp = instance.mlp + + def qkv_hook(_module: Any, hook_args: tuple[Any, ...], hook_result: Any) -> None: + qkv = tensor_result(hook_result) + if qkv is not None and hook_args and isinstance(hook_args[0], torch.Tensor): + intermediates["attention_norm"] = _fused_rms_norm_input( + qkv_projection, hook_args[0], "linear_qkv" + ) + intermediates["qkv"] = qkv + + def core_hook( + _module: Any, + hook_args: tuple[Any, ...], + hook_kwargs: dict[str, Any], + hook_result: Any, + ) -> None: + names = ("query", "key", "value") + for index, name in enumerate(names): + value = hook_args[index] if len(hook_args) > index else hook_kwargs.get(name) + if isinstance(value, torch.Tensor): + intermediates[name] = value + output = tensor_result(hook_result) + if output is not None: + intermediates["attention_core"] = output + + def attention_hook(_module: Any, _args: tuple[Any, ...], hook_result: Any) -> None: + output = tensor_result(hook_result) + if output is not None: + intermediates["attention_output"] = output + + def mlp_hook(_module: Any, hook_args: tuple[Any, ...], hook_result: Any) -> None: + output = tensor_result(hook_result) + if output is not None and hook_args and isinstance(hook_args[0], torch.Tensor): + intermediates["mlp_norm"] = _fused_rms_norm_input( + mlp.linear_fc1, hook_args[0], "linear_fc1" + ) + intermediates["mlp_output"] = output + + handles.extend( + ( + qkv_projection.register_forward_hook(qkv_hook), + core_attention.register_forward_hook(core_hook, with_kwargs=True), + attention.register_forward_hook(attention_hook), + mlp.register_forward_hook(mlp_hook), + ) + ) + try: + result = original(instance, *args, **kwargs) + finally: + for handle in handles: + handle.remove() + output_value = result[0] if isinstance(result, tuple) else result + if isinstance(input_value, torch.Tensor) and isinstance(output_value, torch.Tensor): + attention_output = intermediates.get("attention_output") + if attention_output is not None: + intermediates["attention_residual"] = input_value + attention_output + _save_megatron_layer_diagnostic( + instance, + input_value, + output_value, + intermediates, + ) + return result + + setattr(TransformerLayer, _STRICT_LAYER_DIAGNOSTIC_PATCH_MARKER, original) + TransformerLayer.forward = wrapped + + +def _strict_rocm_rope_positions( + tensor: torch.Tensor, + freqs: torch.Tensor, + cu_seqlens: torch.Tensor | None, + cp_group: Any, +) -> torch.Tensor: + """Recover Megatron's logical RoPE positions without a per-head loop. + + Megatron's packed THD input is stored as two CP-owned sequence chunks. The + deterministic ROCm kernel indexes a head-major view with ``row % S``; + returning one position per token lets every head share one table and one + launch while preserving the framework's zigzag ownership. + """ + + cp_size = 1 if cp_group is None else int(cp_group.size()) + cp_rank = 0 if cp_group is None else int(cp_group.rank()) + if tensor.ndim == 3: + if cu_seqlens is None: + raise RuntimeError("strict ROCm THD RoPE requires cu_seqlens") + values = tuple( + int(value) + for value in cu_seqlens.detach().to(device="cpu", dtype=torch.int64).tolist() + ) + if len(values) < 2 or values[0] != 0: + raise RuntimeError("strict ROCm THD RoPE received invalid cu_seqlens") + exact_global_freqs = int(freqs.size(0)) == values[-1] + positions: list[int] = [] + for index, (start, end) in enumerate(zip(values[:-1], values[1:], strict=True)): + length = end - start + if length <= 0 or length % cp_size: + raise RuntimeError("strict ROCm THD RoPE sequence length is not CP divisible") + local = length // cp_size + if local % 2: + raise RuntimeError("strict ROCm THD RoPE local length must be even") + half = local // 2 + base = start if exact_global_freqs else 0 + positions.extend(range(base + cp_rank * half, base + (cp_rank + 1) * half)) + second = 2 * cp_size - cp_rank - 1 + positions.extend(range(base + second * half, base + (second + 1) * half)) + if len(positions) != tensor.size(0): + raise RuntimeError( + "strict ROCm THD RoPE position count does not match local token rows: " + f"{len(positions)} != {tensor.size(0)}" + ) + return torch.tensor(positions, dtype=torch.int64, device=tensor.device) + + if tensor.ndim != 4: + raise RuntimeError(f"strict ROCm RoPE expects SBHD or THD tensors, got {tensor.shape}") + sequence = int(tensor.size(0)) + if sequence <= 0 or (cp_size > 1 and sequence % 2): + raise RuntimeError("strict ROCm SBHD RoPE received an invalid sequence length") + if cp_size == 1: + local_positions = list(range(sequence)) + else: + half = sequence // 2 + second = 2 * cp_size - cp_rank - 1 + local_positions = list(range(cp_rank * half, (cp_rank + 1) * half)) + local_positions.extend(range(second * half, (second + 1) * half)) + positions = torch.tensor(local_positions, dtype=torch.int64, device=tensor.device) + return positions.unsqueeze(0).expand(tensor.size(1), -1).contiguous() + + +def _apply_strict_rocm_rope( + tensor: torch.Tensor, + positions: torch.Tensor, + operator: Any, +) -> torch.Tensor: + """Apply the shared ROCm RoPE kernel in its head-major indexing layout.""" + + if tensor.size(-1) % 2 or tensor.size(-1) <= 0: + raise RuntimeError("strict ROCm RoPE requires an even head dimension") + if tensor.ndim == 3: + # Megatron THD is token-major; the kernel's table index is shared by + # all heads, so transpose once and launch the operator over [H,T,D]. + head_major = tensor.transpose(0, 1).contiguous() + return operator(head_major, positions).transpose(0, 1).contiguous() + if tensor.ndim == 4: + # SelfAttention uses [S,B,H,D]. The operator accepts [B,H,S,D] with + # one position table per batch row. + batch_head_major = tensor.permute(1, 2, 0, 3).contiguous() + return operator(batch_head_major, positions).permute(2, 0, 1, 3).contiguous() + raise RuntimeError(f"strict ROCm RoPE expects 3-D or 4-D tensors, got {tensor.shape}") + + +def _patch_strict_rocm_rope() -> None: + """Use one deterministic RoPE implementation for Megatron and vLLM.""" + + if torch.version.hip is None: + return + attention_module = importlib.import_module("megatron.core.transformer.attention") + if hasattr(attention_module, _STRICT_ROCM_ROPE_PATCH_MARKER): + return + from rl_engine.kernels.ops.rocm.rotary_embedding.rope import RocmDeterministicRoPEOp + + operator = RocmDeterministicRoPEOp() + original = attention_module.apply_rotary_pos_emb + + def strict_apply_rotary_pos_emb( + tensor: torch.Tensor, + freqs: torch.Tensor, + config: Any, + cu_seqlens: torch.Tensor | None = None, + mscale: float = 1.0, + cp_group: Any = None, + ) -> torch.Tensor: + if torch.version.hip is None: + return original(tensor, freqs, config, cu_seqlens, mscale, cp_group) + if bool(getattr(config, "rotary_interleaved", False)): + raise RuntimeError("strict ROCm RoPE requires non-interleaved Qwen3 layout") + if bool(getattr(config, "multi_latent_attention", False)) or float(mscale) != 1.0: + raise RuntimeError("strict ROCm RoPE does not support MLA or mscale != 1") + if freqs.ndim < 1 or int(freqs.shape[-1]) != int(tensor.shape[-1]): + raise RuntimeError( + "strict ROCm RoPE requires full-dimension frequency tensors: " + f"freqs={tuple(freqs.shape)}, tensor={tuple(tensor.shape)}" + ) + positions = _strict_rocm_rope_positions(tensor, freqs, cu_seqlens, cp_group) + return _apply_strict_rocm_rope(tensor, positions, operator) + + strict_apply_rotary_pos_emb.__name__ = getattr(original, "__name__", "apply_rotary_pos_emb") + setattr(attention_module, _STRICT_ROCM_ROPE_PATCH_MARKER, original) + attention_module.apply_rotary_pos_emb = strict_apply_rotary_pos_emb + + +def _install_torch_dist_object_compatibility() -> None: + """Normalize the PyTorch DCP object shape expected by this Megatron revision.""" + + strategy = importlib.import_module( + "megatron.core.dist_checkpointing.strategies.torch" + ) + original = strategy._replace_sharded_keys_with_state_dict_keys + if getattr(original, "__rl_kernel_dcp_object_compatibility__", False): + return + + def wrapped(state_dict: dict[str, Any], flat_mapping: Any, rename_mapping: Any): + normalized = {} + for key, value in state_dict.items(): + if isinstance(value, io.BytesIO): + value.seek(0) + value = torch.load(value, weights_only=False) + normalized[key] = value + return original(normalized, flat_mapping, rename_mapping) + + wrapped.__rl_kernel_dcp_object_compatibility__ = True + strategy._replace_sharded_keys_with_state_dict_keys = wrapped + + +class _DeterministicTensorParallelReduce(torch.autograd.Function): + """Megatron reduce-from-TP semantics backed by the shared fixed tree.""" + + @staticmethod + def forward(ctx: Any, input_value: torch.Tensor, collective: Any | None) -> torch.Tensor: + del ctx + if collective is None: + return input_value + return collective.all_reduce(input_value.contiguous()) + + @staticmethod + def backward(ctx: Any, grad_output: torch.Tensor) -> tuple[torch.Tensor, None]: + del ctx + # ``reduce_from_tensor_model_parallel_region`` is an all-reduce in the + # forward pass and identity in the backward pass. + return grad_output, None + + +class _DeterministicCopyToTensorParallelRegion(torch.autograd.Function): + """Megatron copy-to-TP semantics with a fixed-tree dgrad reduction.""" + + @staticmethod + def forward(ctx: Any, input_value: torch.Tensor, collective: Any | None) -> torch.Tensor: + ctx.collective = collective + return input_value + + @staticmethod + def backward(ctx: Any, grad_output: torch.Tensor) -> tuple[torch.Tensor, None]: + collective = ctx.collective + if collective is None: + return grad_output, None + return collective.all_reduce(grad_output.contiguous()), None + + +def _deterministic_reduce_from_tensor_model_parallel_region( + input_value: torch.Tensor, + collective: Any | None, +) -> torch.Tensor: + resolved = collective + if resolved is not None and not callable(getattr(resolved, "all_reduce", None)): + from rl_engine.distributed.collectives import collective_for_group + + resolved = collective_for_group( + resolved, + min_size_bytes=input_value.numel() * input_value.element_size(), + ) + return _DeterministicTensorParallelReduce.apply(input_value, resolved) + + +def _deterministic_copy_to_tensor_model_parallel_region( + input_value: torch.Tensor, + collective: Any | None, +) -> torch.Tensor: + return _DeterministicCopyToTensorParallelRegion.apply(input_value, collective) + + +def _module_tp_group(module: Any) -> Any | None: + for name in ("tp_group", "_tp_group"): + group = getattr(module, name, None) + if group is not None: + return group + try: + from megatron.core import parallel_state + + return parallel_state.get_tensor_model_parallel_group() + except (ImportError, AssertionError, RuntimeError): + return None + + +def _tp_world_size(group: Any | None) -> int: + if ( + group is None + or not torch.distributed.is_available() + or not torch.distributed.is_initialized() + ): + return 1 + return int(torch.distributed.get_world_size(group=group)) + + +def _fixed_tree_collective( + module: Any, + input_value: torch.Tensor | None = None, +) -> Any | None: + group = _module_tp_group(module) + if _tp_world_size(group) == 1: + return None + from rl_engine.distributed.collectives import collective_for_group + + min_size_bytes = 0 if input_value is None else input_value.numel() * input_value.element_size() + collective = collective_for_group(group, min_size_bytes=min_size_bytes) + if collective is None: + raise RuntimeError("strict Attention TP collective is not initialized") + backend_id = getattr(collective, "backend_id", None) + if not isinstance(backend_id, str) or not backend_id.strip(): + raise RuntimeError("strict Attention TP collective has no backend identity") + return collective + + +def _collective_backend_id(collective: Any | None) -> str: + if collective is None: + return "none" + backend_id = getattr(collective, "backend_id", None) + if not isinstance(backend_id, str) or not backend_id.strip(): + raise RuntimeError("strict Attention TP collective has no backend identity") + return backend_id.strip() + + +class _DeterministicTPOutputProjection(torch.autograd.Function): + """Materialize the strict TP LM head once and preserve its dgrad contract.""" + + @staticmethod + def forward( + ctx: Any, + input_value: torch.Tensor, + weight: torch.Tensor, + bias: torch.Tensor | None, + tp_group: Any, + ) -> torch.Tensor: + from rl_engine.kernels.ops.matmul.det_gemm import det_gemm_linear + + if input_value.ndim == 3: + input_2d = input_value.transpose(0, 1).contiguous().reshape(-1, input_value.shape[-1]) + ctx.batch_major = True + else: + input_2d = input_value.reshape(-1, input_value.shape[-1]).contiguous() + ctx.batch_major = False + weight_2d = weight.contiguous() + output_2d = det_gemm_linear(input_2d, weight_2d) + if bias is not None: + output_2d = (output_2d.float() + bias.float().reshape(1, -1)).to(torch.bfloat16) + ctx.save_for_backward(input_2d, weight_2d) + ctx.input_shape = input_value.shape + ctx.input_dtype = input_value.dtype + ctx.weight_dtype = weight.dtype + ctx.bias_dtype = None if bias is None else bias.dtype + ctx.has_bias = bias is not None + ctx.tp_group = tp_group + if ctx.batch_major: + return ( + output_2d.reshape(input_value.shape[1], input_value.shape[0], weight.size(0)) + .transpose(0, 1) + .contiguous() + ) + return output_2d.reshape(*input_value.shape[:-1], weight.size(0)) + + @staticmethod + def backward(ctx: Any, grad_output: torch.Tensor): + from rl_engine.kernels.ops.cuda.loss.linear_logp import _deterministic_tp_all_reduce_ + from rl_engine.kernels.ops.matmul.det_gemm import ( + det_gemm_linear_input_gradient, + det_gemm_linear_weight_gradient, + ) + + input_2d, weight = ctx.saved_tensors + if ctx.batch_major: + dlogits = grad_output.transpose(0, 1).contiguous().reshape(-1, grad_output.shape[-1]) + else: + dlogits = grad_output.reshape(-1, grad_output.shape[-1]).contiguous() + if dlogits.dtype != torch.bfloat16: + raise TypeError("strict TP LM-head backward requires BF16 dlogits") + grad_input = grad_weight = grad_bias = None + if ctx.needs_input_grad[0]: + grad_input = det_gemm_linear_input_gradient(dlogits, weight) + _deterministic_tp_all_reduce_(grad_input, ctx.tp_group) + if ctx.batch_major: + grad_input = ( + grad_input.reshape(ctx.input_shape[1], ctx.input_shape[0], ctx.input_shape[2]) + .transpose(0, 1) + .contiguous() + ) + else: + grad_input = grad_input.reshape(ctx.input_shape) + grad_input = grad_input.to(ctx.input_dtype) + if ctx.needs_input_grad[1]: + grad_weight = det_gemm_linear_weight_gradient(input_2d, dlogits).to(ctx.weight_dtype) + if ctx.has_bias and ctx.needs_input_grad[2]: + grad_bias = dlogits.float().sum(dim=0).to(ctx.bias_dtype) + return grad_input, grad_weight, grad_bias, None def _optional_class(path: str) -> type[Any] | None: @@ -101,30 +629,12 @@ def _patch_strict_attention_projections( if hasattr(self_attention_cls, _STRICT_ATTENTION_PATCH_MARKER): return if det_gemm is None: - from rl_engine.kernels.ops.cuda.matmul.det_gemm import DetGemmOp - - det_gemm = DetGemmOp() + det_gemm = _strict_attention_projection_op() attention_init = self_attention_cls.__init__ column_forward_impl = column_linear_cls._forward_impl row_forward_impl = row_linear_cls._forward_impl - def tp_mappings() -> tuple[ - Callable[[torch.Tensor], torch.Tensor], - Callable[[torch.Tensor], torch.Tensor], - ]: - if copy_to_tp is not None and reduce_from_tp is not None: - return copy_to_tp, reduce_from_tp - from megatron.core.tensor_parallel.mappings import ( - copy_to_tensor_model_parallel_region, - reduce_from_tensor_model_parallel_region, - ) - - return ( - copy_to_tensor_model_parallel_region, - reduce_from_tensor_model_parallel_region, - ) - def deterministic_projection( input_value: torch.Tensor, weight: torch.Tensor, @@ -140,29 +650,167 @@ def deterministic_projection( output = output_2d.reshape(*input_value.shape[:-1], weight.shape[0]) return output if bias is None else output + bias + def record_collective_backend(core_attention: Any, attribute: str, backend: str) -> None: + if core_attention is not None: + setattr(core_attention, attribute, backend) + + def callback_backend(callback: Callable[[torch.Tensor], torch.Tensor]) -> str: + backend_id = getattr(callback, "backend_id", None) + return ( + backend_id.strip() + if isinstance(backend_id, str) and backend_id.strip() + else "test_override" + ) + + def strict_tp_copy( + module: Any, + core_attention: Any, + input_value: torch.Tensor, + ) -> torch.Tensor: + if copy_to_tp is not None: + record_collective_backend( + core_attention, + _MEGATRON_TP_QKV_DGRAD_COLLECTIVE_ATTR, + callback_backend(copy_to_tp), + ) + return copy_to_tp(input_value) + collective = _fixed_tree_collective(module, input_value) + record_collective_backend( + core_attention, + _MEGATRON_TP_QKV_DGRAD_COLLECTIVE_ATTR, + _collective_backend_id(collective), + ) + return _deterministic_copy_to_tensor_model_parallel_region(input_value, collective) + + def strict_tp_reduce( + module: Any, + core_attention: Any, + input_value: torch.Tensor, + ) -> torch.Tensor: + if reduce_from_tp is not None: + record_collective_backend( + core_attention, + _MEGATRON_TP_OUTPUT_PROJECTION_COLLECTIVE_ATTR, + callback_backend(reduce_from_tp), + ) + return reduce_from_tp(input_value) + collective = _fixed_tree_collective(module, input_value) + record_collective_backend( + core_attention, + _MEGATRON_TP_OUTPUT_PROJECTION_COLLECTIVE_ATTR, + _collective_backend_id(collective), + ) + return _deterministic_reduce_from_tensor_model_parallel_region(input_value, collective) + + def bind_collective_identity(module: Any, core_attention: Any, attribute: str) -> None: + collective = _fixed_tree_collective(module) + record_collective_backend( + core_attention, + attribute, + _collective_backend_id(collective), + ) + + def sequence_parallel_enabled(module: Any) -> bool: + config = getattr(module, "config", None) + return bool(getattr(module, "sequence_parallel", False)) or bool( + getattr(config, "sequence_parallel", False) + ) + + def local_qkv_forward( + module: Any, + input_value: torch.Tensor, + weight: torch.Tensor | None = None, + runtime_gather_output: bool | None = None, + ) -> tuple[torch.Tensor, torch.Tensor | None]: + gather_output = bool(getattr(module, "gather_output", False)) + if runtime_gather_output is not None: + gather_output = bool(runtime_gather_output) + if gather_output: + raise RuntimeError("strict Attention QKV does not support gathered column output") + selected_weight = getattr(module, "weight", None) if weight is None else weight + if not isinstance(selected_weight, torch.Tensor): + raise RuntimeError("strict Attention QKV requires an allocated weight") + core_attention = getattr(module, _STRICT_ATTENTION_CORE_MARKER, None) + copied = strict_tp_copy(module, core_attention, input_value) + skip_bias_add = bool(getattr(module, "skip_bias_add", False)) + bias = None if skip_bias_add else getattr(module, "bias", None) + output = deterministic_projection(copied, selected_weight, bias) + return output, getattr(module, "bias", None) if skip_bias_add else None + + def local_projection_forward( + module: Any, + input_value: torch.Tensor, + ) -> tuple[torch.Tensor, torch.Tensor | None]: + if not bool(getattr(module, "input_is_parallel", False)): + raise RuntimeError("strict Attention output projection requires TP-sharded input") + weight = getattr(module, "weight", None) + if not isinstance(weight, torch.Tensor): + raise RuntimeError("strict Attention output projection requires an allocated weight") + core_attention = getattr(module, _STRICT_ATTENTION_CORE_MARKER, None) + output = deterministic_projection(input_value, weight, None) + output = strict_tp_reduce(module, core_attention, output) + skip_bias_add = bool(getattr(module, "skip_bias_add", False)) + bias = getattr(module, "bias", None) + if not skip_bias_add and bias is not None: + output = output + bias + return output, bias if skip_bias_add else None + def attention_init_wrapped(instance: Any, *args: Any, **kwargs: Any) -> None: attention_init(instance, *args, **kwargs) qkv = instance.linear_qkv projection = instance.linear_proj + core_attention = getattr(instance, "core_attention", None) + if sequence_parallel_enabled(qkv) or sequence_parallel_enabled(projection): + raise RuntimeError( + "strict Attention projection collectives do not support sequence parallelism" + ) setattr(qkv, _STRICT_ATTENTION_PROJECTION_MARKER, "qkv") setattr(projection, _STRICT_ATTENTION_PROJECTION_MARKER, "o_proj") + setattr(qkv, _STRICT_ATTENTION_CORE_MARKER, core_attention) + setattr(projection, _STRICT_ATTENTION_CORE_MARKER, core_attention) + if copy_to_tp is None: + bind_collective_identity( + qkv, + core_attention, + _MEGATRON_TP_QKV_DGRAD_COLLECTIVE_ATTR, + ) + else: + record_collective_backend( + core_attention, + _MEGATRON_TP_QKV_DGRAD_COLLECTIVE_ATTR, + callback_backend(copy_to_tp), + ) + if reduce_from_tp is None: + bind_collective_identity( + projection, + core_attention, + _MEGATRON_TP_OUTPUT_PROJECTION_COLLECTIVE_ATTR, + ) + else: + record_collective_backend( + core_attention, + _MEGATRON_TP_OUTPUT_PROJECTION_COLLECTIVE_ATTR, + callback_backend(reduce_from_tp), + ) if hasattr(qkv, "layer_norm_weight"): - tp_copy, tp_reduce = tp_mappings() - def te_qkv_forward(module: Any, input_value: torch.Tensor) -> Any: normalized = _fused_rms_norm_input(module, input_value, "linear_qkv") - normalized = tp_copy(normalized) + normalized = strict_tp_copy(module, core_attention, normalized) return deterministic_projection(normalized, module.weight, None), None def te_projection_forward(module: Any, input_value: torch.Tensor) -> Any: output = deterministic_projection(input_value, module.weight, None) - return tp_reduce(output), None + return strict_tp_reduce(module, core_attention, output), None qkv.forward = MethodType(te_qkv_forward, qkv) projection.forward = MethodType(te_projection_forward, projection) else: - # The local ColumnParallelLinear wrapper already owns TP dgrad. + # Per-instance forwards bypass Megatron's arithmetic collectives. qkv.allreduce_dgrad = False + if callable(getattr(qkv, "forward", None)): + qkv.forward = MethodType(local_qkv_forward, qkv) + if callable(getattr(projection, "forward", None)): + projection.forward = MethodType(local_projection_forward, projection) def column_forward_impl_wrapped( instance: Any, @@ -171,7 +819,9 @@ def column_forward_impl_wrapped( *args: Any, **kwargs: Any, ) -> torch.Tensor: - if hasattr(instance, _STRICT_ATTENTION_PROJECTION_MARKER): + if getattr(instance, _STRICT_ATTENTION_PROJECTION_MARKER, None) == "qkv": + core_attention = getattr(instance, _STRICT_ATTENTION_CORE_MARKER, None) + input = strict_tp_copy(instance, core_attention, input) return deterministic_projection(input, weight, kwargs.get("bias")) return column_forward_impl(instance, input, weight, *args, **kwargs) @@ -222,6 +872,69 @@ def wrapped(instance: Any, input_value: torch.Tensor, *args: Any, **kwargs: Any) rms_norm_cls.forward = wrapped +def _patch_strict_logp_output_layer( + linear_cross_entropy_cls: type[Any] | None = None, +) -> None: + """Replace Megatron's duplicate LM-head path with one reusable strict result.""" + + if linear_cross_entropy_cls is None: + from megatron.core.transformer.linear_cross_entropy import LinearCrossEntropyModule + + linear_cross_entropy_cls = LinearCrossEntropyModule + if hasattr(linear_cross_entropy_cls, _STRICT_LOGP_OUTPUT_PATCH_MARKER): + return + original = linear_cross_entropy_cls.forward + + def wrapped( + instance: Any, + input_: torch.Tensor, + weight: torch.Tensor | None = None, + runtime_gather_output: bool | None = None, + output_cross_entropy_loss: bool = False, + labels: torch.Tensor | None = None, + reduction: str = "none", + ignore_index: int = -100, + ) -> Any: + if output_cross_entropy_loss: + return original( + instance, + input_, + weight, + runtime_gather_output, + output_cross_entropy_loss, + labels, + reduction, + ignore_index, + ) + output_weight = instance.weight if weight is None else weight + if output_weight is None: + raise RuntimeError("strict TP LM head requires an explicit weight") + gather_output = ( + instance.gather_output if runtime_gather_output is None else bool(runtime_gather_output) + ) + if gather_output: + raise RuntimeError("strict reusable TP LM head does not support gathered logits") + if bool(getattr(instance, "sequence_parallel", False)): + raise RuntimeError("strict reusable TP LM head does not support sequence parallelism") + if bool(getattr(instance, "explicit_expert_comm", False)) or bool( + getattr(instance, "disable_grad_reduce", False) + ): + raise RuntimeError("strict reusable TP LM head requires ordinary TP dgrad reduction") + bias = instance.bias if not instance.skip_bias_add else None + output = _DeterministicTPOutputProjection.apply( + input_, output_weight, bias, instance.tp_group + ) + instance._rl_kernel_local_logits = output + output_bias = instance.bias if instance.skip_bias_add else None + return output, output_bias + + wrapped.__name__ = getattr(original, "__name__", "forward") + wrapped.__doc__ = getattr(original, "__doc__", None) + setattr(linear_cross_entropy_cls, _STRICT_LOGP_OUTPUT_PATCH_MARKER, original) + setattr(linear_cross_entropy_cls, _STRICT_LOGP_REUSABLE_MARKER, True) + linear_cross_entropy_cls.forward = wrapped + + def install_megatron_integration( plan: IntegrationPlan, *, @@ -262,7 +975,11 @@ def install_megatron_integration( raise RuntimeError("R/R Megatron FFN selected but no supported class was found") set_active_integration("megatron", integration) + _patch_layer_alignment_diagnostics() + if plan.implementation_for("logp", "training") is Implementation.RL_KERNEL: + _patch_strict_logp_output_layer() if plan.implementation_for("attention", "training") is Implementation.RL_KERNEL: + _patch_strict_rocm_rope() _patch_strict_attention_projections() if plan.implementation_for("ffn", "training") is Implementation.RL_KERNEL: _patch_strict_te_rms_norm() @@ -280,9 +997,12 @@ def install_megatron_integration( "ffn", ",".join(f"{cls.__module__}.{cls.__name__}.forward" for cls in resolved_ffn), ) - integration.record_installed_hook( - "logp", "rl_engine.integrations.vime.linear_logp_provider.provider" - ) + if plan.implementation_for("logp", "training") is Implementation.RL_KERNEL: + integration.record_installed_hook( + "logp", + "rl_engine.integrations.vime.linear_logp_provider.provider," + "megatron.core.transformer.linear_cross_entropy.LinearCrossEntropyModule.forward", + ) return integration @@ -291,6 +1011,7 @@ def initialize_from_environment(_args: Any = None) -> MegatronIntegration: from rl_engine.integrations.ablation import integration_plan_from_environment + _install_torch_dist_object_compatibility() plan = integration_plan_from_environment() integration = install_megatron_integration(plan) if plan.implementation_for("attention", "training") is Implementation.RL_KERNEL: diff --git a/rl_engine/integrations/runtime.py b/rl_engine/integrations/runtime.py index ae65a2f5..3807ad52 100644 --- a/rl_engine/integrations/runtime.py +++ b/rl_engine/integrations/runtime.py @@ -124,7 +124,16 @@ def execute( if _contains_profile_only(raw_provenance): self.record_profile_call(normalized) else: - self.record_execution(normalized, selected) + self.record_execution( + normalized, + selected, + execution_provenance=_execution_provenance( + selected=selected, + args=args, + kwargs=kwargs, + result=result, + ), + ) return result def record_profile_call(self, module: str) -> None: @@ -145,6 +154,7 @@ def record_execution( selected: Callable[..., Any], *, execution_mode: str = "eager", + execution_provenance: Mapping[str, Any] | None = None, ) -> None: """Record one eager or custom-op execution outside Dynamo tracing.""" @@ -166,8 +176,13 @@ def record_execution( if implementation is Implementation.PRODUCTION else f"rlkernel.{normalized}.unidentified" ) - raw_provenance = getattr(selected, "provenance", {}) - provenance = dict(raw_provenance) if isinstance(raw_provenance, Mapping) else {} + if execution_provenance is None: + raw_provenance = getattr(selected, "provenance", {}) + provenance = ( + dict(raw_provenance) if isinstance(raw_provenance, Mapping) else {} + ) + else: + provenance = dict(execution_provenance) actual_backend = _actual_backend(provenance) or backend_id fallback = _contains_fallback(provenance) should_report = False @@ -210,7 +225,9 @@ def record_execution( flush=True, ) if implementation is Implementation.RL_KERNEL and fallback: - self.record_fallback(normalized, f"operator provenance selected {actual_backend}") + self.record_fallback( + normalized, f"operator provenance selected {actual_backend}" + ) raise RuntimeError( f"{self.framework} {normalized} strict RL-Kernel route reported fallback" ) @@ -225,7 +242,8 @@ def readback(self) -> dict[str, Any]: "fallbacks": list(self._fallbacks), "profile_calls": dict(self._profile_counts), "operators": { - module: readback.to_dict() for module, readback in self._readbacks.items() + module: readback.to_dict() + for module, readback in self._readbacks.items() }, } @@ -371,4 +389,76 @@ def _runtime_platform(value: Any) -> str | None: return None +def _execution_provenance( + *, + selected: Callable[..., Any], + args: tuple[Any, ...], + kwargs: Mapping[str, Any], + result: Any, +) -> dict[str, Any]: + """Capture backend evidence from the callable, result, and real tensors.""" + + raw = getattr(selected, "provenance", {}) + provenance = dict(raw) if isinstance(raw, Mapping) else {} + result_provenance = getattr(result, "provenance", {}) + if isinstance(result_provenance, Mapping) and result_provenance: + if not provenance: + provenance.update(result_provenance) + else: + provenance.setdefault("result", dict(result_provenance)) + if _runtime_platform(provenance) is None: + inferred = _infer_runtime_platform((args, kwargs, result)) + if inferred is not None: + provenance["runtime_platform"] = inferred + return provenance + + +def _infer_runtime_platform(value: Any) -> str | None: + """Infer the device from bounded traversal of execution inputs and outputs.""" + + platforms: set[str] = set() + seen: set[int] = set() + structural_fields = ( + "context", + "hidden", + "hidden_states", + "logits", + "target_ids", + "logp", + "entropy", + ) + + def visit(item: Any, depth: int) -> None: + if depth > 4 or len(seen) > 256: + return + if isinstance(item, torch.Tensor): + platforms.add(item.device.type) + return + item_id = id(item) + if item_id in seen: + return + seen.add(item_id) + if isinstance(item, Mapping): + for nested in item.values(): + visit(nested, depth + 1) + return + if isinstance(item, (list, tuple)): + for nested in item: + visit(nested, depth + 1) + return + for field_name in structural_fields: + try: + nested = getattr(item, field_name) + except (AttributeError, RuntimeError): + continue + visit(nested, depth + 1) + + visit(value, 0) + if "cuda" in platforms: + return "cuda" + if len(platforms) == 1: + return next(iter(platforms)) + return None + + __all__ = ["FrameworkOperatorIntegration", "OperatorReadback"] diff --git a/rl_engine/integrations/vime/linear_logp_provider.py b/rl_engine/integrations/vime/linear_logp_provider.py index d70bf6c7..476ba5db 100644 --- a/rl_engine/integrations/vime/linear_logp_provider.py +++ b/rl_engine/integrations/vime/linear_logp_provider.py @@ -216,6 +216,43 @@ def _contract_for_request(request: Any) -> tuple[LogprobContract, int]: return contract, _tile_count(metadata, padded_vocab_size) +def _is_identity_temperature(value: Any) -> bool: + return value is None or (not isinstance(value, torch.Tensor) and float(value) == 1.0) + + +@torch.no_grad() +def _metric_entropy_from_strict_lse( + local_logits: torch.Tensor, + lse: torch.Tensor, + *, + vocab_start_index: int, + real_vocab_size: int, + tp_group: Any, +) -> torch.Tensor: + """Compute metric-only entropy without re-running selected-logp/LSE.""" + + local_real = max( + 0, + min(local_logits.size(1), int(real_vocab_size) - int(vocab_start_index)), + ) + centered = local_logits[:, :local_real].float() + centered.sub_(lse.reshape(-1, 1).float()) + probabilities = centered.exp() + centered.neg_() + local_entropy = torch.einsum("ij,ij->i", probabilities, centered).contiguous() + _rank, world = _tp_coordinates(tp_group) + if world == 1: + return local_entropy + import torch.distributed as dist + + gathered = [torch.empty_like(local_entropy) for _ in range(world)] + dist.all_gather(gathered, local_entropy, group=tp_group) + entropy = gathered[0].clone() + for partial in gathered[1:]: + entropy.add_(partial) + return entropy + + def _provider_impl(request: Any, *, linear_logp: Any = None) -> LinearLogpResult: """Compute Vime log-probabilities on the explicit TP/CP contract.""" @@ -240,20 +277,94 @@ def _provider_impl(request: Any, *, linear_logp: Any = None) -> LinearLogpResult raise RuntimeError("linear_logp context must expose projection.weight") if partition is None: raise RuntimeError("linear_logp context must expose vocab_partition") - logp = linear_logp( - hidden, - projection.weight, - request.target_ids, - getattr(projection, "bias", None), - tp_group=getattr(request, "tensor_parallel_group", None), - vocab_start_index=int(partition.local_start), - global_vocab_size=int(partition.padded_size), - real_vocab_size=int(partition.real_size), - temperature=getattr(request, "temperature", None), + reuse_local_logits = bool(getattr(context, "reuse_local_logits", False)) + # Megatron's strict output-layer patch materializes the same local + # padded-vocabulary logits that vLLM serves. The Vime context type in + # this revision predates the explicit ``local_logits`` field, so use + # the request tensor when it is the matching ROCm TP shard. This + # avoids a second GEMM and keeps the backward graph attached to the + # output projection. + request_logits = getattr(request, "logits", None) + materialized_local_logits = ( + torch.version.hip is not None + and isinstance(request_logits, torch.Tensor) + and request_logits.ndim == 2 + and request_logits.dtype in (torch.bfloat16, torch.float16, torch.float32) + and request_logits.shape == ( + hidden.size(0), + projection.weight.size(0), + ) + ) + if materialized_local_logits: + reuse_local_logits = True + with_entropy = bool(getattr(request, "with_entropy", False)) + with_entropy_grad = bool(getattr(request, "with_entropy_grad", False)) + fast_metric_entropy = ( + reuse_local_logits + and with_entropy + and not with_entropy_grad + and _is_identity_temperature(getattr(request, "temperature", None)) ) + strict_lse = None + if reuse_local_logits: + local_logits = getattr(context, "local_logits", None) + if materialized_local_logits: + local_logits = request_logits + if not isinstance(local_logits, torch.Tensor): + raise RuntimeError("strict reusable LM-head context is missing local logits") + from_local_logits = getattr(linear_logp, "from_local_logits", None) + if not callable(from_local_logits): + raise RuntimeError( + "strict reusable LM-head logits require a from_local_logits provider" + ) + scored = from_local_logits( + local_logits, + request.target_ids, + tp_group=getattr(request, "tensor_parallel_group", None), + vocab_start_index=int(partition.local_start), + global_vocab_size=int(partition.padded_size), + real_vocab_size=int(partition.real_size), + target="training", + temperature=getattr(request, "temperature", None), + return_lse=fast_metric_entropy, + diagnostics_hidden=hidden, + diagnostics_lm_head_weight=projection.weight, + ) + if fast_metric_entropy: + logp, strict_lse = scored + else: + logp = scored + else: + logp = linear_logp( + hidden, + projection.weight, + request.target_ids, + getattr(projection, "bias", None), + tp_group=getattr(request, "tensor_parallel_group", None), + vocab_start_index=int(partition.local_start), + global_vocab_size=int(partition.padded_size), + real_vocab_size=int(partition.real_size), + temperature=getattr(request, "temperature", None), + ) entropy = None entropy_provenance: dict[str, Any] = {} - if getattr(request, "with_entropy", False): + if fast_metric_entropy: + if strict_lse is None: + raise RuntimeError("metric-only entropy requires the strict local-logits LSE") + entropy = _metric_entropy_from_strict_lse( + local_logits, + strict_lse, + vocab_start_index=int(partition.local_start), + real_vocab_size=int(partition.real_size), + tp_group=getattr(request, "tensor_parallel_group", None), + ) + entropy_provenance = { + "backend_id": "rlkernel.strict-lse-metric-entropy.v1", + "logits_materialized": True, + "strict_lse_reused": True, + "with_entropy_grad": False, + } + elif with_entropy: entropy_contract, entropy_tiles = _contract_for_request(request) entropy_dispatch = kernel_registry.get_logprob_op( entropy_contract, requested_backend=BACKEND_ID @@ -271,7 +382,7 @@ def _provider_impl(request: Any, *, linear_logp: Any = None) -> LinearLogpResult contract=entropy_contract, tp_group=getattr(request, "tensor_parallel_group", None), num_vocab_tiles=entropy_tiles, - with_entropy_grad=bool(getattr(request, "with_entropy_grad", False)), + with_entropy_grad=with_entropy_grad, ) entropy_provenance = { "backend_id": entropy_dispatch.capability.backend_id, @@ -285,7 +396,10 @@ def _provider_impl(request: Any, *, linear_logp: Any = None) -> LinearLogpResult "strict_backend": True, "top_p_replay": False, "cp_is_merge_axis": False, - "logits_materialized": bool(getattr(request, "with_entropy", False)), + "logits_materialized": bool( + reuse_local_logits or getattr(request, "with_entropy", False) + ), + "lm_head_result_reused": reuse_local_logits, "entropy": entropy_provenance, } provenance["request"] = { @@ -382,7 +496,11 @@ def provider(request: Any) -> LinearLogpResult: case = operator_ablation_case("logp", os.getenv("RL_KERNEL_LOGP_CASE", "P/P")) if case.training is Implementation.PRODUCTION: - return _provider_impl(request) + raise RuntimeError( + "the RL-Kernel linear_logp provider must not be configured for a " + "production Megatron logp route; omit --linear-logp-provider so " + "Vime executes calculate_log_probs_and_entropy directly" + ) from rl_engine.integrations.state import get_active_integration diff --git a/rl_engine/integrations/vllm_runtime.py b/rl_engine/integrations/vllm_runtime.py index 37e2c521..63e43424 100644 --- a/rl_engine/integrations/vllm_runtime.py +++ b/rl_engine/integrations/vllm_runtime.py @@ -7,7 +7,10 @@ import importlib import os +import pathlib +import re import sys +from types import MethodType from typing import Any import torch @@ -16,6 +19,8 @@ DETERMINISTIC_ALL_REDUCE_OP, collective_for_group, deterministic_all_reduce_inplace, + deterministic_all_reduce_staged, + deterministic_staging_reserve, ) from rl_engine.integrations.ablation import ( Implementation, @@ -27,6 +32,7 @@ VllmAttentionOperator, VllmFFNOperator, VllmLogpOperator, + _strict_attention_projection_op, ) from rl_engine.integrations.linear_logp import ( clear_rollout_linear_logp_context, @@ -42,13 +48,216 @@ _STRICT_PROJECTION_MARKER = "__rl_kernel_strict_attention_projection__" _STRICT_FFN_INIT_MARKER = "__rl_kernel_original_strict_ffn_init__" _STRICT_RMS_NORM_INIT_MARKER = "__rl_kernel_original_strict_rms_norm_init__" +_STRICT_ATTENTION_RMS_NORM_MARKER = "__rl_kernel_strict_attention_rms_norm__" _STRICT_ROTARY_INIT_MARKER = "__rl_kernel_original_strict_rotary_init__" +_STRICT_ROCM_ROPE_PATCH_MARKER = "__rl_kernel_original_strict_rocm_rope_forward__" _STRICT_LM_HEAD_LINEAR_PATCH_MARKER = "__rl_kernel_original_lm_head_linear_apply__" _STRICT_O_PROJ_COLLECTIVE_MARKER = "__rl_kernel_o_proj_collective__" _STRICT_ROW_PARALLEL_PATCH_MARKER = "__rl_kernel_original_row_parallel_forward__" +_STRICT_DIRECT_STAGING_MARKER = "__rl_kernel_direct_staging_active__" +_STRICT_LAYER_DIAGNOSTIC_PATCH_MARKER = "__rl_kernel_original_layer_diagnostic_forward__" _RLK_ATTENTION_BACKEND: type[Any] | None = None _RLK_ATTENTION_IMPL: type[Any] | None = None _RLK_ATTENTION_BUILDER: type[Any] | None = None +_RLK_O_PROJ_COLLECTIVE_BACKEND: str | None = None +_VLLM_LAYER_DIAGNOSTIC_BUFFER: dict[str, Any] | None = None +_VLLM_LAYER_DIAGNOSTIC_CALLS = 0 +_VLLM_LAYER_DIAGNOSTIC_ACTIVE_LAYER: int | None = None + + +def _alignment_diagnostics_enabled() -> bool: + value = os.getenv( + "RL_KERNEL_LAYER_ALIGNMENT_DIAGNOSTICS", + os.getenv("RL_KERNEL_ALIGNMENT_DIAGNOSTICS", ""), + ) + return value.strip().lower() in { + "1", + "true", + "yes", + "on", + } + + +def _diagnostic_rank() -> int: + if torch.distributed.is_available() and torch.distributed.is_initialized(): + return int(torch.distributed.get_rank()) + return int(os.getenv("RANK", "0")) + + +def _patch_qwen3_layer_alignment_diagnostics() -> None: + """Record one bounded semantic decoder trace per vLLM model forward.""" + + if not _alignment_diagnostics_enabled(): + return + from vllm.model_executor.models.qwen3 import Qwen3Attention, Qwen3DecoderLayer + from rl_engine.kernels.ops.rocm.attention.strict_runtime import StrictRocmAttentionRuntime + + if hasattr(Qwen3DecoderLayer, _STRICT_LAYER_DIAGNOSTIC_PATCH_MARKER): + return + original_init = Qwen3DecoderLayer.__init__ + original_forward = Qwen3DecoderLayer.forward + original_attention_forward = Qwen3Attention.forward + original_gather_paged_row = StrictRocmAttentionRuntime._gather_paged_row + + def init_wrapped(instance: Any, *args: Any, **kwargs: Any) -> None: + original_init(instance, *args, **kwargs) + config = args[0] if args else kwargs.get("config") + prefix = kwargs.get("prefix", args[3] if len(args) > 3 else "") + match = re.search(r"(?:^|\.)layers\.(\d+)(?:\.|$)", str(prefix)) + if match is None: + raise RuntimeError(f"cannot recover Qwen3 decoder layer from prefix {prefix!r}") + instance._rl_kernel_layer_diagnostic_index = int(match.group(1)) + instance._rl_kernel_layer_diagnostic_count = int(config.num_hidden_layers) + + def forward_wrapped( + instance: Any, + positions: torch.Tensor, + hidden_states: torch.Tensor, + residual: torch.Tensor | None, + ) -> tuple[torch.Tensor, torch.Tensor]: + global _VLLM_LAYER_DIAGNOSTIC_BUFFER + global _VLLM_LAYER_DIAGNOSTIC_CALLS + global _VLLM_LAYER_DIAGNOSTIC_ACTIVE_LAYER + + layer = int(instance._rl_kernel_layer_diagnostic_index) + layer_count = int(instance._rl_kernel_layer_diagnostic_count) + max_rows = int(os.getenv("RL_KERNEL_ALIGNMENT_MAX_ROWS", "64")) + if layer == 0: + semantic_input = hidden_states if residual is None else hidden_states + residual + if semantic_input.ndim != 2: + raise RuntimeError("vLLM layer diagnostics require [T,H] hidden states") + _VLLM_LAYER_DIAGNOSTIC_BUFFER = ( + { + "positions": positions.detach(), + "input": semantic_input.detach(), + "outputs": [], + } + if int(semantic_input.size(0)) <= max_rows + else None + ) + + _VLLM_LAYER_DIAGNOSTIC_ACTIVE_LAYER = layer + try: + result = original_forward(instance, positions, hidden_states, residual) + finally: + _VLLM_LAYER_DIAGNOSTIC_ACTIVE_LAYER = None + output, output_residual = result + buffer = _VLLM_LAYER_DIAGNOSTIC_BUFFER + if buffer is None: + return result + semantic_output = output + output_residual + buffer["outputs"].append(semantic_output.detach()) + if layer == 0: + layer_zero = buffer.setdefault("layer0", {}) + layer_zero["attention_residual"] = output_residual.detach() + layer_zero["mlp_norm"] = torch.nn.functional.rms_norm( + output_residual, + (output_residual.shape[-1],), + instance.post_attention_layernorm.weight, + instance.post_attention_layernorm.variance_epsilon, + ).detach() + layer_zero["mlp_output"] = output.detach() + if layer != layer_count - 1: + return result + if len(buffer["outputs"]) != layer_count: + raise RuntimeError( + "vLLM layer diagnostics did not observe every decoder layer: " + f"{len(buffer['outputs'])} != {layer_count}" + ) + call_index = _VLLM_LAYER_DIAGNOSTIC_CALLS + _VLLM_LAYER_DIAGNOSTIC_CALLS += 1 + rank = _diagnostic_rank() + payload = { + "schema_version": "rlkernel.layer_alignment_diagnostic.v1", + "framework": "vllm", + "pid": os.getpid(), + "rank": rank, + "call_index": call_index, + "positions": buffer["positions"].cpu(), + "input": buffer["input"].cpu(), + "outputs": torch.stack(buffer["outputs"]).cpu(), + } + if "layer0" in buffer: + payload["layer0"] = { + name: value.cpu() for name, value in buffer["layer0"].items() + } + _VLLM_LAYER_DIAGNOSTIC_BUFFER = None + root = os.getenv("RL_KERNEL_ALIGNMENT_DIAGNOSTICS_DIR", "").strip() + if root: + output_dir = pathlib.Path(root) / "layers" + output_dir.mkdir(parents=True, exist_ok=True) + torch.save( + payload, + output_dir + / ( + f"vllm-pid{os.getpid()}-rank{rank:05d}-" + f"call{call_index:08d}.pt" + ), + ) + return result + + def attention_forward_wrapped( + instance: Any, + positions: torch.Tensor, + hidden_states: torch.Tensor, + ) -> torch.Tensor: + buffer = _VLLM_LAYER_DIAGNOSTIC_BUFFER + if buffer is None or _VLLM_LAYER_DIAGNOSTIC_ACTIVE_LAYER != 0: + return original_attention_forward(instance, positions, hidden_states) + qkv, _ = instance.qkv_proj(hidden_states) + q, k, v = qkv.split([instance.q_size, instance.kv_size, instance.kv_size], dim=-1) + q_by_head = q.view(*q.shape[:-1], q.shape[-1] // instance.head_dim, instance.head_dim) + q_by_head = instance.q_norm(q_by_head) + q = q_by_head.view(q.shape) + k_by_head = k.view(*k.shape[:-1], k.shape[-1] // instance.head_dim, instance.head_dim) + k_by_head = instance.k_norm(k_by_head) + k = k_by_head.view(k.shape) + q, k = instance.rotary_emb(positions, q, k) + attention_core = instance.attn(q, k, v) + output, _ = instance.o_proj(attention_core) + buffer.setdefault("layer0", {}).update({ + "attention_norm": hidden_states.detach(), + "qkv": qkv.detach(), + "query": q.view(*q.shape[:-1], -1, instance.head_dim).detach(), + "key": k.view(*k.shape[:-1], -1, instance.head_dim).detach(), + "value": v.view(*v.shape[:-1], -1, instance.head_dim).detach(), + "attention_core": attention_core.detach(), + "attention_output": output.detach(), + }) + return output + + def gather_paged_row_wrapped( + k_cache: torch.Tensor, + v_cache: torch.Tensor, + page_row: torch.Tensor, + cached_length: int, + ) -> tuple[torch.Tensor, torch.Tensor]: + key, value = original_gather_paged_row( + k_cache, + v_cache, + page_row, + cached_length, + ) + buffer = _VLLM_LAYER_DIAGNOSTIC_BUFFER + if buffer is not None and _VLLM_LAYER_DIAGNOSTIC_ACTIVE_LAYER == 0: + buffer.setdefault("layer0", {}).update( + { + "logical_key": key.detach(), + "logical_value": value.detach(), + } + ) + return key, value + + setattr(Qwen3DecoderLayer, _STRICT_LAYER_DIAGNOSTIC_PATCH_MARKER, original_forward) + Qwen3DecoderLayer.__init__ = init_wrapped + Qwen3DecoderLayer.forward = forward_wrapped + setattr(Qwen3Attention, _STRICT_LAYER_DIAGNOSTIC_PATCH_MARKER, original_attention_forward) + Qwen3Attention.forward = attention_forward_wrapped + StrictRocmAttentionRuntime._gather_paged_row = staticmethod(gather_paged_row_wrapped) + + +def _o_proj_collective_backend() -> str | None: + return _RLK_O_PROJ_COLLECTIVE_BACKEND def _is_worker_sampler_profile_batch(value: Any) -> bool: @@ -261,7 +470,7 @@ def _patch_strict_lm_head_linear( return previous_apply = linear_method_cls.apply if det_gemm is None: - from rl_engine.kernels.ops.cuda.matmul.det_gemm import DetGemmOp + from rl_engine.kernels.ops.matmul.det_gemm import DetGemmOp det_gemm = DetGemmOp() @@ -286,6 +495,59 @@ def wrapped( setattr(linear_method_cls, "apply", wrapped) +def _patch_strict_rocm_rotary_embedding(rotary_cls: type[Any]) -> None: + """Bind vLLM Qwen3 RotaryEmbedding to the shared ROCm RoPE kernel.""" + + if torch.version.hip is None or hasattr(rotary_cls, _STRICT_ROCM_ROPE_PATCH_MARKER): + return + from rl_engine.kernels.ops.rocm.rotary_embedding.rope import RocmDeterministicRoPEOp + + operator = RocmDeterministicRoPEOp() + original = rotary_cls.forward_cuda + + def strict_forward_cuda( + instance: Any, + positions: torch.Tensor, + query: torch.Tensor, + key: torch.Tensor | None = None, + ) -> tuple[torch.Tensor, torch.Tensor | None]: + head_size = int(getattr(instance, "head_size", 0)) + rotary_dim = int(getattr(instance, "rotary_dim", head_size)) + if head_size <= 0 or rotary_dim != head_size: + raise RuntimeError( + "strict ROCm Qwen3 RoPE requires full-dimension rotation: " + f"rotary_dim={rotary_dim}, head_size={head_size}" + ) + if positions.ndim not in (1, 2) or positions.numel() != query.shape[0]: + raise RuntimeError( + "strict ROCm vLLM RoPE positions must align with flattened query rows: " + f"positions={tuple(positions.shape)}, query={tuple(query.shape)}" + ) + flat_positions = positions.reshape(-1).to(device=query.device, dtype=torch.int64) + + def apply(value: torch.Tensor | None) -> torch.Tensor | None: + if value is None: + return None + if value.ndim != 2 or value.shape[1] % head_size: + raise RuntimeError( + "strict ROCm vLLM RoPE expects flattened [tokens, heads*head_dim] tensors" + ) + tokens = value.shape[0] + heads = value.shape[1] // head_size + # The HIP kernel indexes one position table across rows. A + # head-major view avoids duplicating the table for every head and + # keeps the dispatch to one deterministic launch per Q/K tensor. + head_major = value.view(tokens, heads, head_size).permute(1, 0, 2).contiguous() + rotated = operator(head_major, flat_positions) + return rotated.permute(1, 0, 2).reshape_as(value).contiguous() + + return apply(query), apply(key) + + strict_forward_cuda.__name__ = getattr(original, "__name__", "forward_cuda") + setattr(rotary_cls, _STRICT_ROCM_ROPE_PATCH_MARKER, original) + rotary_cls.forward_cuda = strict_forward_cuda + + def _configure_strict_ffn_compilation(vllm_config: Any | None = None) -> None: """Keep the graph-safe TP reduction inside vLLM CUDA graphs.""" @@ -328,13 +590,17 @@ def wrapped_init(instance: Any, *args: Any, **kwargs: Any) -> None: original_init(instance, *args, **kwargs) _handle, tp_world_size = operator.bind_packed_inference(instance) if not compiled_evidence_armed: + execution_mode = ( + "eager" if getattr(torch.version, "hip", None) is not None + else "compiled_cuda_graph" + ) register_packed_inference_observer( lambda: integration.record_execution( - "ffn", operator, execution_mode="compiled_cuda_graph" + "ffn", operator, execution_mode=execution_mode ) ) compiled_evidence_armed = True - if tp_world_size > 1: + if tp_world_size > 1 and getattr(torch.version, "hip", None) is None: _configure_strict_ffn_compilation() setattr(Qwen2MLP, _STRICT_FFN_INIT_MARKER, original_init) @@ -395,10 +661,10 @@ def _patch_qwen3_strict_model( assert rms_norm_cls is not None assert linear_method_cls is not None assert attention_cls is not None + if rotary_cls is not None: + _patch_strict_rocm_rotary_embedding(rotary_cls) if det_gemm is None: - from rl_engine.kernels.ops.cuda.matmul.det_gemm import DetGemmOp - - det_gemm = DetGemmOp() + det_gemm = _strict_attention_projection_op() attention_init = attention_cls.__init__ unquantized_apply = linear_method_cls.apply @@ -414,14 +680,79 @@ def deterministic_linear_apply( return unquantized_apply(method, layer, x, bias) x_2d = x.reshape(-1, x.shape[-1]) linear = getattr(det_gemm, "linear", None) + collective = getattr(layer, _STRICT_O_PROJ_COLLECTIVE_MARKER, None) + direct_output = None + if ( + torch.version.hip is None + and collective is not None + and bias is None + and linear is not None + ): + direct_output = collective.direct_staging_view( + (x_2d.size(0), layer.weight.shape[0]), + dtype=x.dtype, + ) + if direct_output is not None: + deterministic_staging_reserve( + direct_output, + collective_handle=int(collective._handle), + ) output_2d = ( - linear(x_2d, layer.weight) + linear(x_2d, layer.weight, out=direct_output) if linear is not None else det_gemm(x_2d, layer.weight.t().contiguous()) ) + setattr(layer, _STRICT_DIRECT_STAGING_MARKER, direct_output is not None) output = output_2d.reshape(*x.shape[:-1], layer.weight.shape[0]) return output if bias is None else output + bias + def strict_attention_rms_norm_forward( + instance: Any, + x: torch.Tensor, + residual: torch.Tensor | None = None, + ) -> torch.Tensor: + if residual is not None: + raise RuntimeError("strict Attention Q/K RMSNorm does not accept a residual") + if instance.variance_size_override is not None: + raise RuntimeError("strict Attention Q/K RMSNorm requires the full head dimension") + weight = instance.weight.data if instance.has_weight else None + if weight is None: + raise RuntimeError("strict Attention Q/K RMSNorm requires a weight") + return strict_rms_norm(x, weight, eps=instance.variance_epsilon) + + def require_rocm_eager_runtime() -> None: + if not production_classes or torch.version.hip is None: + return + from vllm.config import ( + CUDAGraphMode, + CompilationMode, + get_current_vllm_config_or_none, + ) + + config = get_current_vllm_config_or_none() + model_config = None if config is None else config.model_config + compilation_config = None if config is None else config.compilation_config + if ( + model_config is None + or model_config.enforce_eager is not True + or compilation_config is None + or compilation_config.mode != CompilationMode.NONE + or compilation_config.cudagraph_mode != CUDAGraphMode.NONE + ): + raise RuntimeError( + "strict ROCm vLLM Attention requires enforce_eager with compilation " + "and CUDA/HIP graph capture disabled" + ) + + def bind_attention_rms_norm(attention: Any, name: str) -> None: + norm = getattr(attention, name, None) + if not isinstance(norm, rms_norm_cls): + raise RuntimeError(f"strict Qwen3 Attention requires an RMSNorm {name} instance") + if getattr(norm, _STRICT_ATTENTION_RMS_NORM_MARKER, False): + raise RuntimeError(f"strict Qwen3 Attention {name} is already bound") + norm._forward_method = MethodType(strict_attention_rms_norm_forward, norm) + setattr(norm, _STRICT_ATTENTION_RMS_NORM_MARKER, True) + def strict_rms_norm_forward_cuda( instance: Any, x: torch.Tensor, @@ -446,7 +777,10 @@ def strict_rms_norm_forward_cuda( ) def bind_o_proj_collective(module: Any) -> None: + global _RLK_O_PROJ_COLLECTIVE_BACKEND + if int(getattr(module, "tp_size", 1)) <= 1: + _RLK_O_PROJ_COLLECTIVE_BACKEND = "none" return from vllm.distributed.parallel_state import get_tp_group @@ -456,6 +790,21 @@ def bind_o_proj_collective(module: Any) -> None: if collective is None: raise RuntimeError("strict rollout o_proj requires an initialized TP process group") setattr(module, _STRICT_O_PROJ_COLLECTIVE_MARKER, collective) + backend_id = getattr(collective, "backend_id", None) + if not isinstance(backend_id, str) or not backend_id.strip(): + raise RuntimeError("strict rollout o_proj collective has no backend identity") + _RLK_O_PROJ_COLLECTIVE_BACKEND = backend_id.strip() + if torch.version.hip is not None: + return + max_capture = int(os.getenv("RL_KERNEL_VLLM_CUDAGRAPH_MAX_CAPTURE_SIZE", "0")) + if max_capture <= 0: + raise RuntimeError( + "strict rollout direct staging requires a positive graph capture size" + ) + collective.prepare_direct_staging_views( + ((batch, int(module.weight.shape[0])) for batch in range(1, max_capture + 1)), + dtype=module.weight.dtype, + ) if row_parallel_cls is not None and not hasattr( row_parallel_cls, _STRICT_ROW_PARALLEL_PATCH_MARKER @@ -481,11 +830,19 @@ def strict_row_parallel_forward(instance: Any, input_: torch.Tensor) -> Any: output_parallel = instance.quant_method.apply(instance, input_parallel, bias_) if instance.reduce_results and instance.tp_size > 1: - deterministic_all_reduce_inplace( - output_parallel, - collective_handle=int(collective._handle), - ) - output = output_parallel + if torch.version.hip is not None: + output = collective.all_reduce(output_parallel) + elif bool(getattr(instance, _STRICT_DIRECT_STAGING_MARKER, False)): + output = deterministic_all_reduce_staged( + output_parallel, + collective_handle=int(collective._handle), + ) + else: + deterministic_all_reduce_inplace( + output_parallel, + collective_handle=int(collective._handle), + ) + output = output_parallel else: output = output_parallel @@ -522,9 +879,13 @@ def rotary_init_wrapped(instance: Any, *args: Any, **kwargs: Any) -> None: return def attention_init_wrapped(instance: Any, *args: Any, **kwargs: Any) -> None: + require_rocm_eager_runtime() attention_init(instance, *args, **kwargs) setattr(instance.qkv_proj, _STRICT_PROJECTION_MARKER, "qkv") setattr(instance.o_proj, _STRICT_PROJECTION_MARKER, "o_proj") + if torch.version.hip is not None and production_classes: + bind_attention_rms_norm(instance, "q_norm") + bind_attention_rms_norm(instance, "k_norm") bind_o_proj_collective(instance.o_proj) setattr(attention_cls, _STRICT_MODEL_PATCH_MARKER, attention_init) @@ -633,7 +994,9 @@ def _register_attention_backend(integration: VllmIntegration) -> None: operator: VllmAttentionOperator | None = None if integration.plan.implementation_for("attention", "rollout") is Implementation.RL_KERNEL: - operator = VllmAttentionOperator() + operator = VllmAttentionOperator( + projection_collective_backend=_o_proj_collective_backend + ) integration.install_operator("attention", operator) class RlKernelAttentionImpl(PlatformAttentionImpl): @@ -700,11 +1063,13 @@ def install_vllm_integration(plan: IntegrationPlan) -> VllmIntegration: # vLLM imports this legacy rotary path even when every operator is routed # to production. FA4-only installations need the bundled compatibility # namespace before any attention backend is imported. - _install_flash_attn_ops_compatibility() + if torch.version.hip is None: + _install_flash_attn_ops_compatibility() strict_linear_logp = plan.implementation_for("logp", "rollout") is Implementation.RL_KERNEL strict_attention = plan.implementation_for("attention", "rollout") is Implementation.RL_KERNEL if strict_attention: _patch_qwen3_strict_model() + _patch_qwen3_layer_alignment_diagnostics() if strict_linear_logp: _patch_qwen_lm_head_padding() _patch_strict_lm_head_linear() @@ -714,15 +1079,21 @@ def install_vllm_integration(plan: IntegrationPlan) -> VllmIntegration: # integration object chooses native versus RL-Kernel per module. _register_attention_backend(integration) _patch_qwen_ffn(integration) - _patch_sampler(integration, strict_linear_logp=strict_linear_logp) - # vLLM 0.16's GPU model runner invokes vllm.v1.sample.sampler.Sampler. - # Its separate worker sampler has an incompatible seven-argument API; it - # must not replace the single logp route used by the active V1 runner. + if torch.version.hip is not None: + # Current ROCm deployments use vLLM's V2 runner and its worker sampler. + _patch_worker_sampler(integration, strict_linear_logp=strict_linear_logp) + else: + _patch_sampler(integration, strict_linear_logp=strict_linear_logp) if strict_linear_logp: + sampler_hook = ( + "vllm.v1.worker.gpu.sample.sampler.Sampler.__call__" + if torch.version.hip is not None + else "vllm.v1.sample.sampler.Sampler.forward" + ) integration.record_installed_hook( "logp", "vllm.model_executor.models.qwen3.Qwen3ForCausalLM.compute_logits," - "vllm.v1.sample.sampler.Sampler.forward", + f"{sampler_hook}", ) return integration diff --git a/rl_engine/kernels/attention_preprocess.py b/rl_engine/kernels/attention_preprocess.py index 746990de..d82d9df3 100644 --- a/rl_engine/kernels/attention_preprocess.py +++ b/rl_engine/kernels/attention_preprocess.py @@ -370,17 +370,11 @@ def __init__( if self.device.index is None: self.device = torch.device("cuda", current_device) - from rl_engine.kernels.ops.cuda.rotary_embedding import rope as rope_module + from rl_engine.kernels.ops.rocm.rotary_embedding.rope import RocmDeterministicRoPEOp from rl_engine.kernels.ops.triton.rmsnorm_triton import RMSNormTritonOp - rocm_rope_type = getattr(rope_module, "RocmDeterministicRoPEOp", None) - if rocm_rope_type is None: - raise RuntimeError( - "ROCm Attention preprocessing requires RocmDeterministicRoPEOp " - "from the ROCm Attention integration" - ) self.rmsnorm = RMSNormTritonOp() - self.rope = rocm_rope_type() + self.rope = RocmDeterministicRoPEOp() self.deterministic_backend_ids = ROCM_ATTENTION_PREPROCESS_BACKENDS self.device_capability = (0, 0) if native_qk_norm is None and reuse_transformer_engine_qk_norm: diff --git a/rl_engine/kernels/ops/cuda/attention/cp_comm.py b/rl_engine/kernels/ops/cuda/attention/cp_comm.py index 6766bb67..f2895d57 100644 --- a/rl_engine/kernels/ops/cuda/attention/cp_comm.py +++ b/rl_engine/kernels/ops/cuda/attention/cp_comm.py @@ -338,6 +338,165 @@ def backward(ctx, grad_global: torch.Tensor) -> tuple[torch.Tensor, None, None]: return grad_local.movedim(0, ctx.sequence_dim).contiguous(), None, None +class _AllGatherKVSequence(torch.autograd.Function): + """Gather equal-shape K/V shards through one deterministic collective.""" + + @staticmethod + def forward( + ctx, + local_k: torch.Tensor, + local_v: torch.Tensor, + collective: Any, + sequence_dim: int, + ) -> tuple[torch.Tensor, torch.Tensor]: + ctx.collective = collective + ctx.sequence_dim = int(sequence_dim) + packed = torch.stack( + ( + local_k.movedim(ctx.sequence_dim, 0), + local_v.movedim(ctx.sequence_dim, 0), + ), + dim=1, + ) + gathered = collective.all_gather(packed) + return ( + gathered[:, 0].movedim(0, ctx.sequence_dim).contiguous(), + gathered[:, 1].movedim(0, ctx.sequence_dim).contiguous(), + ) + + @staticmethod + def backward( + ctx, + grad_global_k: torch.Tensor, + grad_global_v: torch.Tensor, + ) -> tuple[torch.Tensor, torch.Tensor, None, None]: + packed = torch.stack( + ( + grad_global_k.movedim(ctx.sequence_dim, 0), + grad_global_v.movedim(ctx.sequence_dim, 0), + ), + dim=1, + ) + local = ctx.collective.reduce_scatter(packed) + return ( + local[:, 0].movedim(0, ctx.sequence_dim).contiguous(), + local[:, 1].movedim(0, ctx.sequence_dim).contiguous(), + None, + None, + ) + + +class _AllGatherSequenceFirst(torch.autograd.Function): + """Gather a local Attention tensor and keep the sequence-leading result.""" + + @staticmethod + def forward(ctx, local: torch.Tensor, collective: Any, sequence_dim: int) -> torch.Tensor: + ctx.collective = collective + ctx.sequence_dim = int(sequence_dim) + packed = local.movedim(ctx.sequence_dim, 0).contiguous() + return collective.all_gather(packed) + + @staticmethod + def backward(ctx, grad_global: torch.Tensor) -> tuple[torch.Tensor, None, None]: + grad_local = ctx.collective.reduce_scatter(grad_global.contiguous()) + return grad_local.movedim(0, ctx.sequence_dim).contiguous(), None, None + + +class _AllGatherKVSequenceFirst(torch.autograd.Function): + """Gather K/V together and retain sequence-leading views for direct reorder.""" + + @staticmethod + def forward( + ctx, + local_k: torch.Tensor, + local_v: torch.Tensor, + collective: Any, + sequence_dim: int, + ) -> tuple[torch.Tensor, torch.Tensor]: + ctx.collective = collective + ctx.sequence_dim = int(sequence_dim) + packed = torch.stack( + ( + local_k.movedim(ctx.sequence_dim, 0), + local_v.movedim(ctx.sequence_dim, 0), + ), + dim=1, + ) + gathered = collective.all_gather(packed) + return gathered[:, 0], gathered[:, 1] + + @staticmethod + def backward( + ctx, + grad_global_k: torch.Tensor, + grad_global_v: torch.Tensor, + ) -> tuple[torch.Tensor, torch.Tensor, None, None]: + packed = torch.stack((grad_global_k, grad_global_v), dim=1) + local = ctx.collective.reduce_scatter(packed) + return ( + local[:, 0].movedim(0, ctx.sequence_dim).contiguous(), + local[:, 1].movedim(0, ctx.sequence_dim).contiguous(), + None, + None, + ) + + +class _AllGatherQKVSequenceFirst(torch.autograd.Function): + """Gather Q/K/V through one collective and expose sequence-leading views.""" + + @staticmethod + def forward( + ctx, + local_q: torch.Tensor, + local_k: torch.Tensor, + local_v: torch.Tensor, + collective: Any, + sequence_dim: int, + ) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor]: + ctx.collective = collective + ctx.sequence_dim = int(sequence_dim) + ctx.q_heads = int(local_q.size(1)) + ctx.kv_heads = int(local_k.size(1)) + packed = torch.cat( + ( + local_q.movedim(ctx.sequence_dim, 0), + local_k.movedim(ctx.sequence_dim, 0), + local_v.movedim(ctx.sequence_dim, 0), + ), + dim=2, + ) + gathered = collective.all_gather(packed) + k_start = ctx.q_heads + v_start = k_start + ctx.kv_heads + return ( + gathered[:, :, :k_start], + gathered[:, :, k_start:v_start], + gathered[:, :, v_start:], + ) + + @staticmethod + def backward( + ctx, + grad_global_q: torch.Tensor, + grad_global_k: torch.Tensor, + grad_global_v: torch.Tensor, + ) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor, None, None]: + packed = torch.cat( + (grad_global_q, grad_global_k, grad_global_v), + dim=2, + ) + local = ctx.collective.reduce_scatter(packed) + k_start = ctx.q_heads + v_start = k_start + ctx.kv_heads + return ( + local[:, :, :k_start].movedim(0, ctx.sequence_dim).contiguous(), + local[:, :, k_start:v_start].movedim(0, ctx.sequence_dim).contiguous(), + local[:, :, v_start:].movedim(0, ctx.sequence_dim).contiguous(), + None, + None, + ) + + class _RootReduceScatterSequence(torch.autograd.Function): """Scatter one authoritative full result and gather its gradient back to the root.""" @@ -377,6 +536,35 @@ def backward(ctx, grad_local: torch.Tensor) -> tuple[torch.Tensor, None, None, N return grad_full, None, None, None, None +class _RootReduceScatterSequenceFirst(torch.autograd.Function): + """RS an already sequence-leading full tensor without another layout copy.""" + + @staticmethod + def forward( + ctx, + full: torch.Tensor, + collective: Any, + output_sequence_dim: int, + rank: int, + root: int, + ) -> torch.Tensor: + ctx.collective = collective + ctx.output_sequence_dim = int(output_sequence_dim) + ctx.rank = int(rank) + ctx.root = int(root) + packed = full if ctx.rank == ctx.root else torch.zeros_like(full) + local = collective.reduce_scatter(packed) + return local.movedim(0, ctx.output_sequence_dim).contiguous() + + @staticmethod + def backward(ctx, grad_local: torch.Tensor) -> tuple[torch.Tensor, None, None, None, None]: + packed = grad_local.movedim(ctx.output_sequence_dim, 0).contiguous() + grad_full = ctx.collective.all_gather(packed) + if ctx.rank != ctx.root: + grad_full.zero_() + return grad_full, None, None, None, None + + class CUDAAGRSAttentionCPCommunication: """Deterministic CUDA AG/RS adapter backed by PR311/PR312.""" @@ -448,9 +636,75 @@ def all_gather_kv( _validate_local_kv_shard(local_k, local_v, plan) _require_equal_kv_owner_widths(plan, "self-owned CUDA AG") collective = self._get_collective(plan) - global_k = self._all_gather_sequence_tensor(local_k, collective, sequence_dim=2) - global_v = self._all_gather_sequence_tensor(local_v, collective, sequence_dim=2) - return global_k, global_v + return _AllGatherKVSequence.apply(local_k, local_v, collective, 2) + + def all_gather_query_sequence_first( + self, + local_q: torch.Tensor, + plan: AttentionCPCommunicationPlan, + ) -> torch.Tensor: + """Gather Q without copying the sequence-leading collective result back.""" + + self._validate_cuda_plan(plan) + _validate_query_shard(local_q, plan) + ranges = plan.query_token_ranges + if len({end - start for start, end in ranges}) != 1: + raise AttentionCPCommunicationUnavailable( + "self-owned CUDA AG requires equal query shard lengths" + ) + return _AllGatherSequenceFirst.apply(local_q, self._get_collective(plan), 2) + + def all_gather_kv_sequence_first( + self, + local_k: torch.Tensor, + local_v: torch.Tensor, + plan: AttentionCPCommunicationPlan, + ) -> tuple[torch.Tensor, torch.Tensor]: + """Gather K/V once and expose sequence-leading views to strict FA4.""" + + self._validate_cuda_plan(plan) + _validate_local_kv_shard(local_k, local_v, plan) + _require_equal_kv_owner_widths(plan, "self-owned CUDA AG") + return _AllGatherKVSequenceFirst.apply( + local_k, + local_v, + self._get_collective(plan), + 2, + ) + + def all_gather_qkv_sequence_first( + self, + local_q: torch.Tensor, + local_k: torch.Tensor, + local_v: torch.Tensor, + plan: AttentionCPCommunicationPlan, + ) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor]: + """Gather equal-length Q/K/V shards with one collective invocation.""" + + self._validate_cuda_plan(plan) + _validate_query_shard(local_q, plan) + _validate_local_kv_shard(local_k, local_v, plan) + _require_equal_kv_owner_widths(plan, "self-owned CUDA AG") + ranges = plan.query_token_ranges + if len({end - start for start, end in ranges}) != 1: + raise AttentionCPCommunicationUnavailable( + "self-owned CUDA AG requires equal query shard lengths" + ) + if ( + local_q.size(0) != local_k.size(0) + or local_q.size(2) != local_k.size(2) + or local_q.size(3) != local_k.size(3) + ): + raise AttentionCPCommunicationUnavailable( + "fused Q/K/V gather requires matching batch, sequence, and head dimensions" + ) + return _AllGatherQKVSequenceFirst.apply( + local_q, + local_k, + local_v, + self._get_collective(plan), + 2, + ) def all_gather_position_ids( self, @@ -462,6 +716,11 @@ def all_gather_position_ids( _validate_local_position_ids(local_query_positions, local_key_positions, plan) _require_equal_kv_owner_widths(plan, "self-owned CUDA AG") collective = self._get_collective(plan) + if local_query_positions is local_key_positions: + positions = self._all_gather_sequence_tensor( + local_query_positions, collective, sequence_dim=1 + ) + return positions, positions query_positions = self._all_gather_sequence_tensor( local_query_positions, collective, sequence_dim=1 ) @@ -570,6 +829,53 @@ def reduce_scatter_strict_result( result.validate() return result + def reduce_scatter_strict_result_sequence_first( + self, + out: torch.Tensor, + lse: torch.Tensor, + plan: AttentionCPCommunicationPlan, + ) -> AttentionCPOutputShard: + """RS strict outputs already laid out as [S, B, H, ...].""" + + self._validate_cuda_plan(plan) + ranges = plan.query_token_ranges + if len({end - start for start, end in ranges}) != 1: + raise AttentionCPCommunicationUnavailable( + "self-owned CUDA RS requires equal contiguous query ranges" + ) + full_tokens = ranges[-1][1] + if ( + out.ndim != 4 + or lse.ndim != 3 + or out.size(0) != full_tokens + or lse.size(0) != full_tokens + or out.shape[1:3] != lse.shape[1:3] + ): + raise ValueError("sequence-first strict outputs must use [S,B,H,D] and [S,B,H]") + if not out.is_cuda or not lse.is_cuda or not out.is_contiguous() or not lse.is_contiguous(): + raise ValueError("sequence-first strict outputs must be contiguous CUDA tensors") + collective = self._get_collective(plan) + rank = plan.parallel.cp_rank + root = plan.merge_root_cp_rank + result = AttentionCPOutputShard( + out=_RootReduceScatterSequenceFirst.apply( + out, + collective, + 2, + rank, + root, + ), + lse=_RootReduceScatterSequenceFirst.apply( + lse, + collective, + 2, + rank, + root, + ), + ) + result.validate() + return result + def _dist(self): import torch.distributed as dist @@ -695,18 +1001,6 @@ class RCCLAGRSAttentionCPCommunication(CUDAAGRSAttentionCPCommunication): supports_async_overlap = False supports_compute_communication_fusion = False - def _get_collective(self, plan: AttentionCPCommunicationPlan): - if self._collective is None: - self._collective = _RCCLRankOrderedTransport( - process_group=self._process_group, - root=plan.merge_root_cp_rank, - ) - if self._collective.world_size != plan.parallel.cp_world_size: - raise AttentionCPCommunicationUnavailable( - "self-owned RCCL world size does not match the CP plan" - ) - return self._collective - def _dist(self): import torch.distributed as dist diff --git a/rl_engine/kernels/ops/cuda/attention/flash_attn.py b/rl_engine/kernels/ops/cuda/attention/flash_attn.py index e57cdeb5..5b75d601 100644 --- a/rl_engine/kernels/ops/cuda/attention/flash_attn.py +++ b/rl_engine/kernels/ops/cuda/attention/flash_attn.py @@ -108,6 +108,52 @@ def __init__( self.api_source = api_source self._op = op self._paged_op = paged_op + self._validated_position_layouts: dict[tuple[Any, ...], None] = {} + + def _validate_positions_cached( + self, + q: torch.Tensor, + k: torch.Tensor, + *, + causal: bool, + query_position_ids: torch.Tensor | None, + key_position_ids: torch.Tensor | None, + ) -> None: + key = ( + tuple(q.shape), + tuple(k.shape), + bool(causal), + ( + None + if query_position_ids is None + else ( + query_position_ids.data_ptr(), + int(query_position_ids._version), + tuple(query_position_ids.shape), + ) + ), + ( + None + if key_position_ids is None + else ( + key_position_ids.data_ptr(), + int(key_position_ids._version), + tuple(key_position_ids.shape), + ) + ), + ) + if key in self._validated_position_layouts: + return + RLKernelDeterministicAttentionCore._validate_positions( + q, + k, + causal=causal, + query_position_ids=query_position_ids, + key_position_ids=key_position_ids, + ) + if len(self._validated_position_layouts) >= 128: + self._validated_position_layouts.pop(next(iter(self._validated_position_layouts))) + self._validated_position_layouts[key] = None @classmethod def precompile_training( @@ -230,7 +276,7 @@ def forward_with_lse( output_dtype: torch.dtype | None = None, ) -> DeterministicAttentionCoreResult: self._validate_inputs(q, k, v, key_padding_mask) - RLKernelDeterministicAttentionCore._validate_positions( + self._validate_positions_cached( q, k, causal=causal, @@ -276,7 +322,7 @@ def forward_bshd_with_lse( ) -> DeterministicAttentionCoreResult: """Run strict FA4 on tensors already laid out as [B, S, H, D].""" self._validate_bshd_inputs(q, k, v, key_padding_mask) - RLKernelDeterministicAttentionCore._validate_positions( + self._validate_positions_cached( q.transpose(1, 2), k.transpose(1, 2), causal=causal, diff --git a/rl_engine/kernels/ops/cuda/attention/strict_runtime.py b/rl_engine/kernels/ops/cuda/attention/strict_runtime.py index d7d95f9b..af5d9ef0 100644 --- a/rl_engine/kernels/ops/cuda/attention/strict_runtime.py +++ b/rl_engine/kernels/ops/cuda/attention/strict_runtime.py @@ -61,6 +61,54 @@ def __init__( if getattr(self._core, "strict_schedule", None) != self.strict_schedule: raise RuntimeError("strict CUDA Attention runtime requires the FA4 fixed schedule") self.communication_executed = False + self._position_layout_cache: dict[tuple[Any, ...], tuple[torch.Tensor, ...]] = {} + self._validated_global_position_layouts: dict[tuple[Any, ...], None] = {} + + def _position_layout( + self, + query_position_ids: torch.Tensor, + key_position_ids: torch.Tensor, + plan: AttentionCPCommunicationPlan, + ) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor]: + key = ( + query_position_ids.data_ptr(), + int(query_position_ids._version), + key_position_ids.data_ptr(), + int(key_position_ids._version), + tuple(query_position_ids.shape), + tuple(key_position_ids.shape), + plan.parallel.cp_rank, + plan.parallel.cp_world_size, + ) + cached = self._position_layout_cache.get(key) + if cached is not None: + return cached # type: ignore[return-value] + global_q_positions, global_k_positions = self._communication.all_gather_position_ids( + query_position_ids, + key_position_ids, + plan, + ) + q_sort = torch.argsort(global_q_positions, dim=1) + if global_q_positions is global_k_positions: + k_sort = q_sort + q_positions_sorted = torch.gather(global_q_positions, 1, q_sort) + k_positions_sorted = q_positions_sorted + else: + k_sort = torch.argsort(global_k_positions, dim=1) + q_positions_sorted = torch.gather(global_q_positions, 1, q_sort) + k_positions_sorted = torch.gather(global_k_positions, 1, k_sort) + inverse_q_sort = torch.argsort(q_sort, dim=1) + value = ( + q_positions_sorted, + k_positions_sorted, + q_sort, + k_sort, + inverse_q_sort, + ) + if len(self._position_layout_cache) >= 128: + self._position_layout_cache.pop(next(iter(self._position_layout_cache))) + self._position_layout_cache[key] = value + return value def forward_with_lse( self, @@ -87,15 +135,47 @@ def forward_with_lse( global_k_positions = key_position_ids communication_backend = "none" self.communication_executed = False + gathered_tensors_are_sequence_first = False else: plan = self._communication_plan(contract, q.size(2), k.size(2)) - global_q = self._communication.all_gather_query(q, plan) - global_k, global_v = self._communication.all_gather_kv(k, v, plan) - global_q_positions, global_k_positions = self._communication.all_gather_position_ids( - query_position_ids, - key_position_ids, - plan, + gather_query_sequence_first = getattr( + self._communication, + "all_gather_query_sequence_first", + None, + ) + gather_kv_sequence_first = getattr( + self._communication, + "all_gather_kv_sequence_first", + None, ) + gather_qkv_sequence_first = getattr( + self._communication, + "all_gather_qkv_sequence_first", + None, + ) + if callable(gather_qkv_sequence_first): + global_q, global_k, global_v = gather_qkv_sequence_first( + q, + k, + v, + plan, + ) + gathered_tensors_are_sequence_first = True + elif callable(gather_query_sequence_first) and callable(gather_kv_sequence_first): + global_q = gather_query_sequence_first(q, plan) + global_k, global_v = gather_kv_sequence_first(k, v, plan) + gathered_tensors_are_sequence_first = True + else: + global_q = self._communication.all_gather_query(q, plan) + global_k, global_v = self._communication.all_gather_kv(k, v, plan) + gathered_tensors_are_sequence_first = False + ( + q_positions_sorted, + k_positions_sorted, + q_sort, + k_sort, + inverse_q_sort, + ) = self._position_layout(query_position_ids, key_position_ids, plan) communication_backend = "cuda_ag_rs" self.communication_executed = True @@ -105,21 +185,47 @@ def forward_with_lse( q_sorted, k_sorted, v_sorted = global_q, global_k, global_v q_positions_sorted, k_positions_sorted = global_q_positions, global_k_positions q_sort = None + sorted_tensors_are_bshd = False else: - q_sorted, q_positions_sorted, q_sort = self._sort_by_position( - global_q, global_q_positions + if cp_world_size > 1: + # Gather directly into FA4's [B, S, H, D] layout. Gathering + # into [B, H, S, D] and transposing afterward materializes + # every global Q/K/V tensor twice on every Attention call. + gather_for_fa4 = ( + self._gather_sequence_first_bshd + if gathered_tensors_are_sequence_first + else self._gather_sequence_bshd + ) + q_sorted = gather_for_fa4(global_q, q_sort) + k_sorted = gather_for_fa4(global_k, k_sort) + sorted_tensors_are_bshd = True + else: + q_sorted, q_positions_sorted, q_sort = self._sort_by_position( + global_q, global_q_positions + ) + k_sorted, k_positions_sorted, k_sort = self._sort_by_position( + global_k, global_k_positions + ) + sorted_tensors_are_bshd = False + v_sorted = ( + gather_for_fa4(global_v, k_sort) + if sorted_tensors_are_bshd + else self._gather_sequence(global_v, k_sort) ) - k_sorted, k_positions_sorted, k_sort = self._sort_by_position( - global_k, global_k_positions + self._validate_global_positions_cached( + q_positions_sorted, + k_positions_sorted, + causal, ) - v_sorted = self._gather_sequence(global_v, k_sort) - self._validate_global_positions(q_positions_sorted, k_positions_sorted, causal) # FA4 consumes [B, S, H, D]. Materialize that layout once for every # logical sequence instead of once per causal prefix. - q_fa = q_sorted.transpose(1, 2).contiguous() - k_fa = k_sorted.transpose(1, 2).contiguous() - v_fa = v_sorted.transpose(1, 2).contiguous() + if sorted_tensors_are_bshd: + q_fa, k_fa, v_fa = q_sorted, k_sorted, v_sorted + else: + q_fa = q_sorted.transpose(1, 2).contiguous() + k_fa = k_sorted.transpose(1, 2).contiguous() + v_fa = v_sorted.transpose(1, 2).contiguous() # FA4's full causal schedule produces the same output, LSE, and dQ as # launching one single-query prefix at a time. Keeping all query rows @@ -135,8 +241,6 @@ def forward_with_lse( key_position_ids=k_positions_sorted, output_dtype=q.dtype, ) - out_sorted = result.out.transpose(1, 2).contiguous() - lse_sorted = result.lse backend = ( result.provenance.get("actual_backend") or result.provenance.get("attention_backend") @@ -146,17 +250,39 @@ def forward_with_lse( if cp_world_size > 1: if q_sort is None: raise RuntimeError("CP Attention requires a framework position reorder") - inverse_q_sort = torch.argsort(q_sort, dim=1) - out_rank_packed = self._gather_sequence(out_sorted, inverse_q_sort) - lse_rank_packed = self._gather_sequence(lse_sorted, inverse_q_sort) - shard = self._communication.reduce_scatter_strict_result( - out_rank_packed, - lse_rank_packed, - plan, + reduce_sequence_first = getattr( + self._communication, + "reduce_scatter_strict_result_sequence_first", + None, ) + if callable(reduce_sequence_first): + out_rank_packed = self._gather_bshd_sequence_first( + result.out, + inverse_q_sort, + ) + lse_rank_packed = self._gather_bhs_sequence_first( + result.lse, + inverse_q_sort, + ) + shard = reduce_sequence_first( + out_rank_packed, + lse_rank_packed, + plan, + ) + else: + out_sorted = result.out.transpose(1, 2).contiguous() + lse_sorted = result.lse + out_rank_packed = self._gather_sequence(out_sorted, inverse_q_sort) + lse_rank_packed = self._gather_sequence(lse_sorted, inverse_q_sort) + shard = self._communication.reduce_scatter_strict_result( + out_rank_packed, + lse_rank_packed, + plan, + ) out, lse = shard.out, shard.lse else: - out, lse = out_sorted, lse_sorted + out = result.out.transpose(1, 2).contiguous() + lse = result.lse return StrictCUDAAttentionResult( out=out, @@ -176,10 +302,10 @@ def forward_with_lse( "framework_position_reorder": True, "query_schedule": "full_sequence_causal_single_launch", "backward_schedule": "fa4_deterministic_full_sequence", - "core_row_count": q_sorted.size(0) * q_sorted.size(2), + "core_row_count": q_fa.size(0) * q_fa.size(1), "core_launch_count": 1, - "core_batch_size": q_sorted.size(0), - "core_query_length": q_sorted.size(2), + "core_batch_size": q_fa.size(0), + "core_query_length": q_fa.size(1), "core_actual_backends": [] if backend is None else [str(backend)], }, ) @@ -334,6 +460,58 @@ def _gather_sequence(tensor: torch.Tensor, order: torch.Tensor) -> torch.Tensor: raise RuntimeError("strict Attention sequence reorder expects a 3-D or 4-D tensor") return torch.gather(tensor, 2, index).contiguous() + @staticmethod + def _gather_sequence_bshd( + tensor: torch.Tensor, + order: torch.Tensor, + ) -> torch.Tensor: + if tensor.ndim != 4: + raise RuntimeError("strict Attention BSHD reorder expects a 4-D tensor") + transposed = tensor.transpose(1, 2) + index = order[:, :, None, None].expand( + tensor.size(0), order.size(1), tensor.size(1), tensor.size(3) + ) + return torch.gather(transposed, 1, index).contiguous() + + @staticmethod + def _gather_sequence_first_bshd( + tensor: torch.Tensor, + order: torch.Tensor, + ) -> torch.Tensor: + if tensor.ndim != 4: + raise RuntimeError("strict Attention sequence-first BSHD reorder expects a 4-D tensor") + batch_major = tensor.permute(1, 0, 2, 3) + index = order[:, :, None, None].expand( + tensor.size(1), order.size(1), tensor.size(2), tensor.size(3) + ) + return torch.gather(batch_major, 1, index).contiguous() + + @staticmethod + def _gather_bshd_sequence_first( + tensor: torch.Tensor, + order: torch.Tensor, + ) -> torch.Tensor: + if tensor.ndim != 4: + raise RuntimeError("strict Attention output reorder expects a BSHD tensor") + sequence_first = tensor.permute(1, 0, 2, 3) + index = order.transpose(0, 1)[:, :, None, None].expand( + order.size(1), tensor.size(0), tensor.size(2), tensor.size(3) + ) + return torch.gather(sequence_first, 0, index).contiguous() + + @staticmethod + def _gather_bhs_sequence_first( + tensor: torch.Tensor, + order: torch.Tensor, + ) -> torch.Tensor: + if tensor.ndim != 3: + raise RuntimeError("strict Attention LSE reorder expects a BHS tensor") + sequence_first = tensor.permute(2, 0, 1) + index = order.transpose(0, 1)[:, :, None].expand( + order.size(1), tensor.size(0), tensor.size(1) + ) + return torch.gather(sequence_first, 0, index).contiguous() + @staticmethod def _validate_global_positions( query_positions: torch.Tensor, @@ -355,5 +533,37 @@ def _validate_global_positions( "causal Attention queries must be the trailing global KV positions", ) + def _validate_global_positions_cached( + self, + query_positions: torch.Tensor, + key_positions: torch.Tensor, + causal: bool, + ) -> None: + """Validate each immutable sorted position layout only once. + + The position-layout cache owns these tensors and reuses their exact + storage across layers and recomputation. Re-launching three device + assertions on every attention call adds work without adding coverage. + Tensor versions keep the cache fail-safe if a layout is ever mutated. + """ + + key = ( + query_positions.data_ptr(), + int(query_positions._version), + tuple(query_positions.shape), + key_positions.data_ptr(), + int(key_positions._version), + tuple(key_positions.shape), + bool(causal), + ) + if key in self._validated_global_position_layouts: + return + self._validate_global_positions(query_positions, key_positions, causal) + if len(self._validated_global_position_layouts) >= 128: + self._validated_global_position_layouts.pop( + next(iter(self._validated_global_position_layouts)) + ) + self._validated_global_position_layouts[key] = None + __all__ = ["StrictCUDAAttentionResult", "StrictCUDAAttentionRuntime"] diff --git a/rl_engine/kernels/ops/cuda/ffn.py b/rl_engine/kernels/ops/cuda/ffn.py new file mode 100644 index 00000000..a4dd04f0 --- /dev/null +++ b/rl_engine/kernels/ops/cuda/ffn.py @@ -0,0 +1,833 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors +"""Bias-free gated FFN assembled from deterministic CUDA kernels.""" + +from __future__ import annotations + +import os +from collections.abc import Callable +from typing import Any + +import torch +from torch import Tensor + +from rl_engine.distributed.collectives import _COLLECTIVES as _SHARED_COLLECTIVES +from rl_engine.distributed.collectives import ( + collective_for_group, + deterministic_all_reduce_inplace, + deterministic_all_reduce_staged, + deterministic_staging_reserve, +) +from rl_engine.kernels.ops.base import _C, _EXT_AVAILABLE +from rl_engine.kernels.ops.cuda.matmul.det_gemm import ( + det_gemm_linear, + det_gemm_linear_input_gradient, + det_gemm_linear_weight_gradient, +) + +QWEN3_8B_HIDDEN_SIZE = 4096 +QWEN3_8B_INTERMEDIATE_SIZE = 12288 +BACKEND_ID = "rlkernel.ffn.qwen3.deterministic.v1" + +_DET_GEMM_SYMBOLS = ( + "det_gemm_fwd", + "det_gemm_fwd_rhs_transposed", + "det_gemm_db_transposed", +) +_SWIGLU_SYMBOLS = ( + "swiglu_forward", + "swiglu_backward", +) +_PACKED_SWIGLU_SYMBOLS = ( + "swiglu_packed_forward", + "swiglu_packed_backward", +) +_REQUIRED_SYMBOLS = _DET_GEMM_SYMBOLS + _SWIGLU_SYMBOLS +_COLLECTIVE_MIN_CAPACITY_BYTES = 64 * 1024 * 1024 +# Backward-compatible test hook; ownership lives in the shared communication layer. +_COLLECTIVES = _SHARED_COLLECTIVES +_PACKED_INFERENCE_OBSERVERS: list[Callable[[], None]] = [] + + +def register_packed_inference_observer(callback: Callable[[], None]) -> None: + """Arm one execution callback for the graph-captured rollout custom op.""" + + if not callable(callback): + raise TypeError("packed inference observer must be callable") + _PACKED_INFERENCE_OBSERVERS.append(callback) + + +def _notify_packed_inference_observers() -> None: + callbacks = tuple(_PACKED_INFERENCE_OBSERVERS) + _PACKED_INFERENCE_OBSERVERS.clear() + for callback in callbacks: + callback() + + +@torch.library.custom_op("rl_kernel::qwen3_ffn_packed_inference", mutates_args=()) +def _qwen3_ffn_packed_inference( + rmsnorm_output: Tensor, + fused_gate_up_weight: Tensor, + down_weight: Tensor, +) -> Tensor: + """Run the graph-safe compute portion of the strict rollout FFN.""" + + _notify_packed_inference_observers() + input_shape = rmsnorm_output.shape + hidden_2d = rmsnorm_output.reshape(-1, input_shape[-1]).contiguous() + gate_up = det_gemm_linear( + hidden_2d, + fused_gate_up_weight, + native_op=_C.det_gemm_fwd_rhs_transposed, + ) + activated = _C.swiglu_packed_forward(gate_up) + output = det_gemm_linear( + activated, + down_weight, + native_op=_C.det_gemm_fwd_rhs_transposed, + ) + return output.reshape(*input_shape[:-1], output.size(-1)) + + +@_qwen3_ffn_packed_inference.register_fake +def _qwen3_ffn_packed_inference_fake( + rmsnorm_output: Tensor, + fused_gate_up_weight: Tensor, + down_weight: Tensor, +) -> Tensor: + del fused_gate_up_weight + return rmsnorm_output.new_empty((*rmsnorm_output.shape[:-1], down_weight.shape[0])) + + +@torch.library.custom_op( + "rl_kernel::qwen3_ffn_packed_inference_to_staging", + mutates_args={"output"}, +) +def _qwen3_ffn_packed_inference_to_staging( + rmsnorm_output: Tensor, + fused_gate_up_weight: Tensor, + down_weight: Tensor, + output: Tensor, +) -> None: + """Run strict FFN compute with the down projection targeting IPC staging.""" + + _notify_packed_inference_observers() + hidden_2d = rmsnorm_output.reshape(-1, rmsnorm_output.shape[-1]).contiguous() + gate_up = det_gemm_linear( + hidden_2d, + fused_gate_up_weight, + native_op=_C.det_gemm_fwd_rhs_transposed, + ) + activated = _C.swiglu_packed_forward(gate_up) + det_gemm_linear( + activated, + down_weight, + native_op=_C.det_gemm_fwd_rhs_transposed, + out=output, + ) + + +@_qwen3_ffn_packed_inference_to_staging.register_fake +def _qwen3_ffn_packed_inference_to_staging_fake( + rmsnorm_output: Tensor, + fused_gate_up_weight: Tensor, + down_weight: Tensor, + output: Tensor, +) -> None: + del rmsnorm_output, fused_gate_up_weight, down_weight, output + + +def qwen3_ffn_packed_inference( + rmsnorm_output: Tensor, + fused_gate_up_weight: Tensor, + down_weight: Tensor, + *, + collective_handle: int = 0, + tp_world_size: int = 1, + collective: Any | None = None, +) -> Tensor: + """Inference-only packed FFN entry compatible with torch.compile.""" + + if tp_world_size <= 1: + return _qwen3_ffn_packed_inference( + rmsnorm_output, + fused_gate_up_weight, + down_weight, + ) + if collective_handle <= 0: + raise RuntimeError("packed rollout FFN requires a bound TP collective") + input_shape = rmsnorm_output.shape + output_shape_2d = ( + rmsnorm_output.numel() // input_shape[-1], + down_weight.shape[0], + ) + direct_output = ( + None + if collective is None + else collective.direct_staging_view(output_shape_2d, dtype=rmsnorm_output.dtype) + ) + if direct_output is not None: + deterministic_staging_reserve( + direct_output, + collective_handle=collective_handle, + ) + _qwen3_ffn_packed_inference_to_staging( + rmsnorm_output, + fused_gate_up_weight, + down_weight, + direct_output, + ) + reduced = deterministic_all_reduce_staged( + direct_output, + collective_handle=collective_handle, + ) + return reduced.reshape(*input_shape[:-1], down_weight.shape[0]) + output = _qwen3_ffn_packed_inference( + rmsnorm_output, + fused_gate_up_weight, + down_weight, + ) + return deterministic_all_reduce_inplace( + output, + collective_handle=collective_handle, + ) + + +def _require_ffn_kernels(*, disable_split_k: bool, packed_gate_up: bool = False) -> None: + required: tuple[str, ...] = _REQUIRED_SYMBOLS if disable_split_k else _SWIGLU_SYMBOLS + if packed_gate_up: + required = tuple(required) + tuple(_PACKED_SWIGLU_SYMBOLS) + missing = [name for name in required if not hasattr(_C, name)] + if not _EXT_AVAILABLE or _C is None or missing: + suffix = f" Missing symbols: {', '.join(missing)}." if missing else "" + needed = ( + "compiled deterministic GEMM and SwiGLU CUDA kernels" + if disable_split_k + else "compiled SwiGLU CUDA kernels" + ) + raise RuntimeError(f"qwen3_ffn requires the {needed}.{suffix}") + + +def _linear_fwd(a: Tensor, weight: Tensor, *, disable_split_k: bool) -> Tensor: + if disable_split_k: + return det_gemm_linear( + a, + weight, + native_op=_C.det_gemm_fwd_rhs_transposed, + ) + # cuBLASLt / CUTLASS: may use split-K. Detach so Autograd.Function owns backward. + with torch.no_grad(): + return torch.nn.functional.linear(a, weight) + + +def _linear_da(grad_output: Tensor, weight: Tensor, *, disable_split_k: bool) -> Tensor: + if disable_split_k: + return det_gemm_linear_input_gradient( + grad_output, + weight, + native_op=_C.det_gemm_fwd, + ) + with torch.no_grad(): + return torch.matmul(grad_output, weight) + + +def _linear_dw(a: Tensor, grad_output: Tensor, *, disable_split_k: bool) -> Tensor: + if disable_split_k: + return det_gemm_linear_weight_gradient( + a, + grad_output, + native_op=_C.det_gemm_db_transposed, + ) + with torch.no_grad(): + return torch.matmul(grad_output.t().contiguous(), a) + + +def _cp_sharded_linear_dw( + a: Tensor, + grad_output: Tensor, + *, + collective: Any, + disable_split_k: bool, +) -> Tensor: + """Compute disjoint output rows, then reconstruct the exact full wgrad.""" + + world_size = int(collective.world_size) + output_rows = int(grad_output.size(1)) + if output_rows % world_size != 0: + raise ValueError( + "CP-sharded weight-gradient rows must divide evenly across ranks, " + f"got {output_rows} rows and world_size={world_size}." + ) + rows_per_rank = output_rows // world_size + row_start = int(collective.rank) * rows_per_rank + grad_output_shard = grad_output.narrow(1, row_start, rows_per_rank).contiguous() + local_weight_gradient = _linear_dw( + a, + grad_output_shard, + disable_split_k=disable_split_k, + ) + return collective.all_gather(local_weight_gradient.contiguous()) + + +def _require_parallel_group(group: Any, name: str): + if group is None: + return None + + import torch.distributed as dist + + if not dist.is_available(): + raise RuntimeError(f"{name}-parallel FFN requires torch.distributed.") + if not dist.is_initialized(): + raise RuntimeError(f"{name}-parallel FFN requires an initialized process group.") + if dist.get_world_size(group=group) <= 1: + raise ValueError(f"{name}_group must contain at least two ranks.") + return dist + + +def _validate_ffn_inputs( + rmsnorm_output: Tensor, + gate_weight: Tensor, + up_weight: Tensor, + down_weight: Tensor, + fused_gate_up_weight: Tensor | None, +) -> None: + tensors = { + "rmsnorm_output": rmsnorm_output, + "gate_weight": gate_weight, + "up_weight": up_weight, + "down_weight": down_weight, + } + if fused_gate_up_weight is not None: + tensors["fused_gate_up_weight"] = fused_gate_up_weight + for name, tensor in tensors.items(): + if not isinstance(tensor, Tensor): + raise TypeError(f"{name} must be a torch.Tensor, got {type(tensor)!r}.") + + if rmsnorm_output.dim() < 1: + raise ValueError("rmsnorm_output must have at least one dimension.") + if rmsnorm_output.numel() == 0: + raise ValueError("rmsnorm_output must contain at least one token.") + for name, weight in ( + ("gate_weight", gate_weight), + ("up_weight", up_weight), + ("down_weight", down_weight), + ): + if weight.dim() != 2: + raise ValueError(f"{name} must be 2-D, got shape {tuple(weight.shape)}.") + + hidden_size = rmsnorm_output.size(-1) + intermediate_size = gate_weight.size(0) + expected_shapes = { + "gate_weight": (intermediate_size, hidden_size), + "up_weight": (intermediate_size, hidden_size), + "down_weight": (hidden_size, intermediate_size), + } + if fused_gate_up_weight is not None: + expected_shapes["fused_gate_up_weight"] = (2 * intermediate_size, hidden_size) + for name, expected in expected_shapes.items(): + actual = tuple(tensors[name].shape) + if actual != expected: + raise ValueError(f"{name} must have shape {expected}, got {actual}.") + + for name, tensor in tensors.items(): + if tensor.dtype != torch.bfloat16: + raise TypeError(f"{name} must have dtype bfloat16, got {tensor.dtype}.") + if not tensor.is_cuda: + raise RuntimeError(f"{name} must be on a CUDA device, got '{tensor.device}'.") + if tensor.device != rmsnorm_output.device: + raise RuntimeError( + f"all FFN inputs must be on {rmsnorm_output.device}, " + f"got {name} on {tensor.device}." + ) + + +def _collective_for_group(group: Any, *, min_size_bytes: int): + return collective_for_group( + group=group, + min_size_bytes=min_size_bytes, + minimum_capacity_bytes=_COLLECTIVE_MIN_CAPACITY_BYTES, + ) + + +def _all_gather_tokens(tensor: Tensor, collective: Any) -> Tensor: + return collective.all_gather(tensor.contiguous()) + + +def _all_gather_packed_tokens(*tensors: Tensor, collective: Any) -> tuple[Tensor, ...]: + """Gather same-row tensors with one handshake and no repacking copies.""" + + return collective.all_gather_many(tuple(tensor.contiguous() for tensor in tensors)) + + +def _reduce_scatter_tokens(tensor: Tensor, collective: Any) -> Tensor: + world_size = collective.world_size + if tensor.size(0) % world_size != 0: + raise ValueError( + "the gathered token count must be divisible by the tensor-parallel " + f"world size, got {tensor.size(0)} and {world_size}." + ) + return collective.reduce_scatter(tensor.contiguous()) + + +def _all_reduce_inplace(tensor: Tensor, collective: Any) -> Tensor: + return collective.all_reduce(tensor, out=tensor) + + +class _DeterministicFFNFunction(torch.autograd.Function): + @staticmethod + def forward( + ctx, + rmsnorm_output: Tensor, + gate_weight: Tensor, + up_weight: Tensor, + down_weight: Tensor, + fused_gate_up_weight: Tensor | None, + tp_group: Any, + cp_group: Any, + sequence_parallel: bool, + disable_split_k: bool, + ) -> Tensor: + tp_dist = _require_parallel_group(tp_group, "tensor") + _require_parallel_group(cp_group, "context") + if sequence_parallel and tp_dist is None: + raise ValueError("sequence_parallel requires a tensor-parallel group.") + + input_shape = rmsnorm_output.shape + rmsnorm_output_2d = rmsnorm_output.reshape(-1, input_shape[-1]).contiguous() + tp_world = tp_dist.get_world_size(group=tp_group) if tp_dist is not None else 1 + gemm_tokens = rmsnorm_output_2d.size(0) * (tp_world if sequence_parallel else 1) + element_size = rmsnorm_output_2d.element_size() + min_size_bytes = max( + gemm_tokens * rmsnorm_output_2d.size(1) * element_size, + gemm_tokens * gate_weight.size(0) * element_size, + gate_weight.numel() * element_size, + up_weight.numel() * element_size, + down_weight.numel() * element_size, + ) + if cp_group is not None: + packed_width = 2 * rmsnorm_output_2d.size(1) + 3 * gate_weight.size(0) + min_size_bytes = max( + min_size_bytes, + gemm_tokens * packed_width * element_size, + ) + # Create TP before CP so every rank follows the same group order. + tp_collective = _collective_for_group(tp_group, min_size_bytes=min_size_bytes) + cp_collective = _collective_for_group(cp_group, min_size_bytes=min_size_bytes) + + if sequence_parallel: + rmsnorm_output_2d = _all_gather_tokens(rmsnorm_output_2d, tp_collective) + + packed_gate_up = fused_gate_up_weight is not None and disable_split_k + if packed_gate_up: + assert fused_gate_up_weight is not None + gate_up = _linear_fwd( + rmsnorm_output_2d, + fused_gate_up_weight, + disable_split_k=True, + ) + activated = _C.swiglu_packed_forward(gate_up) + else: + gate = _linear_fwd( + rmsnorm_output_2d, + gate_weight, + disable_split_k=disable_split_k, + ) + up = _linear_fwd( + rmsnorm_output_2d, + up_weight, + disable_split_k=disable_split_k, + ) + activated = _C.swiglu_forward(gate, up) + output = _linear_fwd(activated, down_weight, disable_split_k=disable_split_k) + + if sequence_parallel: + output = _reduce_scatter_tokens(output, tp_collective) + elif tp_collective is not None: + output = _all_reduce_inplace(output, tp_collective) + + if packed_gate_up: + ctx.save_for_backward( + rmsnorm_output_2d, + gate_up, + activated, + gate_weight, + up_weight, + down_weight, + ) + else: + ctx.save_for_backward( + rmsnorm_output_2d, + gate, + up, + activated, + gate_weight, + up_weight, + down_weight, + ) + ctx.input_shape = input_shape + ctx.tp_collective = tp_collective + ctx.cp_collective = cp_collective + ctx.sequence_parallel = sequence_parallel + ctx.disable_split_k = disable_split_k + ctx.packed_gate_up = packed_gate_up + return output.reshape(*input_shape[:-1], output.size(-1)) + + @staticmethod + def backward(ctx, grad_output: Tensor): + if ctx.packed_gate_up: + ( + rmsnorm_output, + gate_up, + activated, + gate_weight, + up_weight, + down_weight, + ) = ctx.saved_tensors + else: + ( + rmsnorm_output, + gate, + up, + activated, + gate_weight, + up_weight, + down_weight, + ) = ctx.saved_tensors + tp_collective = ctx.tp_collective + cp_collective = ctx.cp_collective + disable_split_k = ctx.disable_split_k + grad_output = grad_output.reshape(-1, grad_output.size(-1)).contiguous() + if ctx.sequence_parallel: + grad_output = _all_gather_tokens(grad_output, tp_collective) + + # Down input-gradient shards concatenate across TP; no TP reduction. + grad_activated = _linear_da( + grad_output, + down_weight, + disable_split_k=disable_split_k, + ) + if ctx.packed_gate_up: + grad_gate, grad_up = _C.swiglu_packed_backward(grad_activated, gate_up) + else: + grad_gate, grad_up = _C.swiglu_backward(grad_activated, gate, up) + + # Weight gradients must see every CP token so the GEMM K-tree matches + # CP=1. These payloads become available before any weight-gradient GEMM, + # so one rank-ordered gather preserves the arithmetic contract while + # avoiding four redundant collective handshakes per layer. + if cp_collective is not None: + ( + activated_full, + grad_output_full, + rmsnorm_full, + grad_gate_full, + grad_up_full, + ) = _all_gather_packed_tokens( + activated, + grad_output, + rmsnorm_output, + grad_gate, + grad_up, + collective=cp_collective, + ) + grad_down_weight = _cp_sharded_linear_dw( + activated_full, + grad_output_full, + collective=cp_collective, + disable_split_k=disable_split_k, + ) + grad_gate_weight = _cp_sharded_linear_dw( + rmsnorm_full, + grad_gate_full, + collective=cp_collective, + disable_split_k=disable_split_k, + ) + grad_up_weight = _cp_sharded_linear_dw( + rmsnorm_full, + grad_up_full, + collective=cp_collective, + disable_split_k=disable_split_k, + ) + else: + grad_down_weight = _linear_dw( + activated, + grad_output, + disable_split_k=disable_split_k, + ) + grad_gate_weight = _linear_dw( + rmsnorm_output, + grad_gate, + disable_split_k=disable_split_k, + ) + grad_up_weight = _linear_dw( + rmsnorm_output, + grad_up, + disable_split_k=disable_split_k, + ) + + # Gate/Up input gradients reduce across TP, then add locally. + grad_rmsnorm_from_gate = _linear_da( + grad_gate, + gate_weight, + disable_split_k=disable_split_k, + ) + if ctx.sequence_parallel: + grad_rmsnorm_from_gate = _reduce_scatter_tokens( + grad_rmsnorm_from_gate, + tp_collective, + ) + elif tp_collective is not None: + grad_rmsnorm_from_gate = _all_reduce_inplace( + grad_rmsnorm_from_gate, + tp_collective, + ) + + grad_rmsnorm_from_up = _linear_da( + grad_up, + up_weight, + disable_split_k=disable_split_k, + ) + if ctx.sequence_parallel: + grad_rmsnorm_from_up = _reduce_scatter_tokens( + grad_rmsnorm_from_up, + tp_collective, + ) + elif tp_collective is not None: + grad_rmsnorm_from_up = _all_reduce_inplace( + grad_rmsnorm_from_up, + tp_collective, + ) + + grad_rmsnorm_output = grad_rmsnorm_from_gate.add_(grad_rmsnorm_from_up) + return ( + grad_rmsnorm_output.reshape(ctx.input_shape), + grad_gate_weight, + grad_up_weight, + grad_down_weight, + None, + None, + None, + None, + None, + ) + + +def qwen3_ffn( + rmsnorm_output: Tensor, + gate_weight: Tensor, + up_weight: Tensor, + down_weight: Tensor, + *, + fused_gate_up_weight: Tensor | None = None, + tp_group: Any = None, + cp_group: Any = None, + sequence_parallel: bool = False, + deterministic: bool | None = None, + disable_split_k: bool | None = None, +) -> Tensor: + """Apply a bias-free SiLU-gated FFN with deterministic backward kernels. + + Args: + rmsnorm_output: RMSNorm output, shape ``[..., H]``. + gate_weight: Gate projection weight in ``[out, in]`` layout, shape + ``[I_local, H]``. + up_weight: Up projection weight in ``[out, in]`` layout, shape + ``[I_local, H]``. + down_weight: Down projection weight in ``[out, in]`` layout, shape + ``[H, I_local]``. + fused_gate_up_weight: Optional existing framework weight in + ``[2 * I_local, H]`` layout. Strict CUDA execution consumes this + with one GEMM launch while returning gradients through the + separate gate/up views, so the framework parameter layout stays + unchanged. + tp_group: Optional tensor-parallel process group. Gate and Up are + column-parallel; Down is row-parallel. Reductions use the + deterministic fixed-tree collectives rather than NCCL. + cp_group: Optional context-parallel process group. Each rank owns + different token rows and the same local weight shards. Weight + gradients AllGather tokens along CP and run the full-token + ``det_gemm_db_transposed`` so they match CP=1 bitwise. + sequence_parallel: Whether ``rmsnorm_output`` and the returned output + are sharded on the flattened token dimension across ``tp_group``. + Token gather/scatter use the deterministic AllGather and + ReduceScatter. + deterministic: Select the RL-Kernel fixed-reduction GEMM when True + (default), or the production ``torch.matmul`` GEMM when False. + disable_split_k: Compatibility alias for ``deterministic``. New code + should use ``deterministic`` because Split-K is only one possible + implementation detail of the production GEMM. + + Returns: + FFN output with shape ``[..., H]``. + """ + if not isinstance(sequence_parallel, bool): + raise TypeError("sequence_parallel must be a bool.") + deterministic = _resolve_deterministic_mode(deterministic, disable_split_k) + _validate_ffn_inputs( + rmsnorm_output, + gate_weight, + up_weight, + down_weight, + fused_gate_up_weight, + ) + _require_ffn_kernels( + disable_split_k=deterministic, + packed_gate_up=fused_gate_up_weight is not None and deterministic, + ) + return _DeterministicFFNFunction.apply( + rmsnorm_output, + gate_weight, + up_weight, + down_weight, + fused_gate_up_weight, + tp_group, + cp_group, + sequence_parallel, + deterministic, + ) + + +def _resolve_deterministic_mode( + deterministic: bool | None, + disable_split_k: bool | None, +) -> bool: + if deterministic is not None and not isinstance(deterministic, bool): + raise TypeError("deterministic must be a bool or None.") + if disable_split_k is not None and not isinstance(disable_split_k, bool): + raise TypeError("disable_split_k must be a bool or None.") + if ( + deterministic is not None + and disable_split_k is not None + and deterministic != disable_split_k + ): + raise ValueError("deterministic and disable_split_k select conflicting FFN backends.") + if deterministic is not None: + return deterministic + if disable_split_k is not None: + return disable_split_k + return True + + +class Qwen3FFNOp: + """Instantiable Qwen3 FFN wrapper for semantic operator dispatch.""" + + op_class = "ffn" + is_batch_invariant = True + backend_id = BACKEND_ID + + def __init__(self) -> None: + # Keep graph-bound IPC resources alive independently of the module-level + # lookup cache. CUDA Graphs retain only the small opaque C++ handle. + self._packed_inference_collectives: dict[int, Any] = {} + + def prepare_packed_inference( + self, + fused_gate_up_weight: Tensor, + down_weight: Tensor, + *, + tp_group: Any, + ) -> tuple[int, int]: + """Create the rollout TP resource before Dynamo captures the model.""" + + if tp_group is None: + return 0, 1 + dist = _require_parallel_group(tp_group, "tensor") + if dist is None: + return 0, 1 + tp_world_size = int(dist.get_world_size(group=tp_group)) + if fused_gate_up_weight.size(0) % 2: + raise ValueError("fused gate/up weight must contain two equal shards") + element_size = fused_gate_up_weight.element_size() + min_size_bytes = ( + max( + fused_gate_up_weight.numel() // 2, + down_weight.numel(), + ) + * element_size + ) + collective = _collective_for_group( + tp_group, + min_size_bytes=min_size_bytes, + ) + max_capture = int(os.getenv("RL_KERNEL_VLLM_CUDAGRAPH_MAX_CAPTURE_SIZE", "0")) + if max_capture <= 0: + raise RuntimeError("packed rollout FFN requires a positive graph capture size") + collective.prepare_direct_staging_views( + ((batch, int(down_weight.shape[0])) for batch in range(1, max_capture + 1)), + dtype=down_weight.dtype, + ) + collective_handle = int(collective._handle) + self._packed_inference_collectives[collective_handle] = collective + return collective_handle, tp_world_size + + def packed_inference( + self, + rmsnorm_output: Tensor, + fused_gate_up_weight: Tensor, + down_weight: Tensor, + *, + collective_handle: int, + tp_world_size: int, + ) -> Tensor: + return qwen3_ffn_packed_inference( + rmsnorm_output, + fused_gate_up_weight, + down_weight, + collective_handle=collective_handle, + tp_world_size=tp_world_size, + collective=self._packed_inference_collectives.get(collective_handle), + ) + + def __call__( + self, + rmsnorm_output: Tensor, + gate_weight: Tensor, + up_weight: Tensor, + down_weight: Tensor, + *, + fused_gate_up_weight: Tensor | None = None, + tp_group: Any = None, + cp_group: Any = None, + sequence_parallel: bool = False, + deterministic: bool | None = None, + disable_split_k: bool | None = None, + ) -> Tensor: + return self.apply( + rmsnorm_output, + gate_weight, + up_weight, + down_weight, + fused_gate_up_weight=fused_gate_up_weight, + tp_group=tp_group, + cp_group=cp_group, + sequence_parallel=sequence_parallel, + deterministic=deterministic, + disable_split_k=disable_split_k, + ) + + def apply( + self, + rmsnorm_output: Tensor, + gate_weight: Tensor, + up_weight: Tensor, + down_weight: Tensor, + *, + fused_gate_up_weight: Tensor | None = None, + tp_group: Any = None, + cp_group: Any = None, + sequence_parallel: bool = False, + deterministic: bool | None = None, + disable_split_k: bool | None = None, + ) -> Tensor: + return qwen3_ffn( + rmsnorm_output, + gate_weight, + up_weight, + down_weight, + fused_gate_up_weight=fused_gate_up_weight, + tp_group=tp_group, + cp_group=cp_group, + sequence_parallel=sequence_parallel, + deterministic=deterministic, + disable_split_k=disable_split_k, + ) diff --git a/rl_engine/kernels/ops/cuda/loss/linear_logp.py b/rl_engine/kernels/ops/cuda/loss/linear_logp.py index 4df9fd09..e7e4a5b0 100644 --- a/rl_engine/kernels/ops/cuda/loss/linear_logp.py +++ b/rl_engine/kernels/ops/cuda/loss/linear_logp.py @@ -1143,7 +1143,7 @@ def _strict_tp_run( return logp, lse, hidden_2d, weight, target, logits, temp -def sm90_deterministic_logp_from_local_logits_tp( +def _strict_logp_from_local_logits_tp_run( local_logits: torch.Tensor, target_ids: torch.Tensor, *, @@ -1152,8 +1152,8 @@ def sm90_deterministic_logp_from_local_logits_tp( real_vocab_size: int = -1, temperature: Optional[torch.Tensor] = None, tp_group: Any, -) -> tuple[torch.Tensor, torch.Tensor]: - """Merge a deterministic TP LM-head result without recomputing its GEMM.""" +): + """Raw strict TP logp over a previously materialized deterministic LM head.""" required = "linear_logp_local_bf16_forward" if not (_EXT_AVAILABLE and hasattr(_C, required)): @@ -1179,6 +1179,7 @@ def sm90_deterministic_logp_from_local_logits_tp( _assert_global_targets_async(target, real_vocab) logits = local_logits + temp = local_logits.new_empty(0, dtype=torch.float32) if temperature is not None: temp = temperature.to(device=logits.device, dtype=torch.float32).reshape(-1) if temp.numel() == 1: @@ -1202,6 +1203,87 @@ def sm90_deterministic_logp_from_local_logits_tp( ) logp, lse = _merge_tp_local_logp(local_lse, local_target, tp_group=tp_group) lead_shape = target_ids.shape + return logp, lse, target, logits, temp, lead_shape + + +class _StrictLocalLogitsTPAutograd(torch.autograd.Function): + """Differentiate strict selected-logp through a reused TP LM-head result.""" + + @staticmethod + def forward( + ctx, + local_logits, + target_ids, + vocab_start_index, + global_vocab_size, + real_vocab_size, + temperature, + tp_group, + ): + logp, lse, target, logits, temp, lead_shape = _strict_logp_from_local_logits_tp_run( + local_logits, + target_ids, + vocab_start_index=vocab_start_index, + global_vocab_size=global_vocab_size, + real_vocab_size=real_vocab_size, + temperature=temperature, + tp_group=tp_group, + ) + ctx.save_for_backward(target, logits, lse, temp) + ctx.vocab_start = int(vocab_start_index) + ctx.input_shape = local_logits.shape + ctx.set_materialize_grads(False) + return logp.reshape(lead_shape), lse.reshape(lead_shape) + + @staticmethod + def backward(ctx, grad_logp, grad_lse): + target, logits, lse, temp = ctx.saved_tensors + if grad_logp is None and grad_lse is None: + return (None, None, None, None, None, None, None) + logp_grad = torch.zeros_like(lse) if grad_logp is None else grad_logp.reshape(-1).float() + dlogits = torch.empty_like(logits) + _C.linear_logp_logits_bf16_to_dlogits( + logits, dlogits, target, logp_grad, lse, ctx.vocab_start + ) + if grad_lse is not None: + probs = torch.exp(logits.float() - lse.reshape(-1, 1)) + dlogits = (dlogits.float() + probs * grad_lse.reshape(-1, 1).float()).to(torch.bfloat16) + if temp.numel(): + dlogits = (dlogits.float() / temp.reshape(-1, 1)).to(torch.bfloat16) + return dlogits.reshape(ctx.input_shape).contiguous(), None, None, None, None, None, None + + +def sm90_deterministic_logp_from_local_logits_tp( + local_logits: torch.Tensor, + target_ids: torch.Tensor, + *, + vocab_start_index: int, + global_vocab_size: int, + real_vocab_size: int = -1, + temperature: Optional[torch.Tensor] = None, + tp_group: Any, +) -> tuple[torch.Tensor, torch.Tensor]: + """Merge a deterministic TP LM-head result without recomputing its GEMM.""" + + if torch.is_grad_enabled() and local_logits.requires_grad: + return _StrictLocalLogitsTPAutograd.apply( + local_logits, + target_ids, + int(vocab_start_index), + int(global_vocab_size), + int(real_vocab_size), + temperature, + tp_group, + ) + logp, lse, _target, _logits, _temp, lead_shape = _strict_logp_from_local_logits_tp_run( + local_logits, + target_ids, + vocab_start_index=vocab_start_index, + global_vocab_size=global_vocab_size, + real_vocab_size=real_vocab_size, + temperature=temperature, + tp_group=tp_group, + ) return logp.reshape(lead_shape), lse.reshape(lead_shape) diff --git a/rl_engine/kernels/ops/cuda/matmul/det_gemm.py b/rl_engine/kernels/ops/cuda/matmul/det_gemm.py index ab64055d..7c80cdd7 100644 --- a/rl_engine/kernels/ops/cuda/matmul/det_gemm.py +++ b/rl_engine/kernels/ops/cuda/matmul/det_gemm.py @@ -139,13 +139,16 @@ def det_gemm_linear( weight: torch.Tensor, *, native_op: Callable[[torch.Tensor, torch.Tensor], torch.Tensor] | None = None, + out: torch.Tensor | None = None, ) -> torch.Tensor: """Apply a native [N,K] weight through the selected strict backend.""" if det_gemm_backend() == _CUBLASLT_BACKEND: _configure_cublaslt_nosplitk(a) _report_strict_route_once() - return torch.mm(a, weight.t()) + return torch.mm(a, weight.t(), out=out) if out is not None else torch.mm(a, weight.t()) + if out is not None: + raise RuntimeError("direct-output deterministic GEMM currently requires cublaslt_nosplitk") _require_sm90_backend() _report_strict_route_once() if native_op is not None: @@ -349,7 +352,13 @@ def __call__(self, a: torch.Tensor, b: torch.Tensor) -> torch.Tensor: ) return _DetGemmFn.apply(a.contiguous(), b.contiguous(), False) - def linear(self, a: torch.Tensor, weight: torch.Tensor) -> torch.Tensor: + def linear( + self, + a: torch.Tensor, + weight: torch.Tensor, + *, + out: torch.Tensor | None = None, + ) -> torch.Tensor: """Apply a native [N,K] linear weight without materializing weight.T.""" assert a.dtype == torch.bfloat16 and weight.dtype == torch.bfloat16, "BF16 only" assert a.is_cuda and weight.is_cuda, "Inputs must be on CUDA device" @@ -360,7 +369,13 @@ def linear(self, a: torch.Tensor, weight: torch.Tensor) -> torch.Tensor: ) if not self.has_hardware_op or any(not hasattr(_C, name) for name in required): raise RuntimeError("DetGemmOp.linear requires the rebuilt native-weight CUDA extension") - return _DetLinearFn.apply(a.contiguous(), weight.contiguous()) + a = a.contiguous() + weight = weight.contiguous() + if out is not None: + if torch.is_grad_enabled() and (a.requires_grad or weight.requires_grad): + raise RuntimeError("direct-output deterministic GEMM is inference-only") + return det_gemm_linear(a, weight, out=out) + return _DetLinearFn.apply(a, weight) def forward_fp32(self, a: torch.Tensor, b: torch.Tensor) -> torch.Tensor: assert a.dtype == torch.bfloat16 and b.dtype == torch.bfloat16, "BF16 only" diff --git a/rl_engine/kernels/ops/matmul/__init__.py b/rl_engine/kernels/ops/matmul/__init__.py new file mode 100644 index 00000000..bbca6640 --- /dev/null +++ b/rl_engine/kernels/ops/matmul/__init__.py @@ -0,0 +1,5 @@ +"""Platform-selected matrix multiplication operators.""" + +from .det_gemm import DetGemmOp, deterministic_gemm + +__all__ = ["DetGemmOp", "deterministic_gemm"] diff --git a/rl_engine/kernels/ops/matmul/det_gemm.py b/rl_engine/kernels/ops/matmul/det_gemm.py new file mode 100644 index 00000000..95131322 --- /dev/null +++ b/rl_engine/kernels/ops/matmul/det_gemm.py @@ -0,0 +1,39 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors +"""Import-time platform binding for strict deterministic GEMM.""" + +import torch + +if torch.version.hip is not None: + from rl_engine.kernels.ops.rocm.matmul.det_gemm import ( + DetGemmOp, + det_gemm_backend, + det_gemm_backend_id, + det_gemm_fallback_reason, + det_gemm_linear, + det_gemm_linear_input_gradient, + det_gemm_linear_weight_gradient, + deterministic_gemm, + ) +else: + from rl_engine.kernels.ops.cuda.matmul.det_gemm import ( + DetGemmOp, + det_gemm_backend, + det_gemm_backend_id, + det_gemm_fallback_reason, + det_gemm_linear, + det_gemm_linear_input_gradient, + det_gemm_linear_weight_gradient, + deterministic_gemm, + ) + +__all__ = [ + "DetGemmOp", + "det_gemm_backend", + "det_gemm_backend_id", + "det_gemm_fallback_reason", + "det_gemm_linear", + "det_gemm_linear_input_gradient", + "det_gemm_linear_weight_gradient", + "deterministic_gemm", +] diff --git a/rl_engine/kernels/ops/pytorch/ffn/ffn.py b/rl_engine/kernels/ops/pytorch/ffn/ffn.py index 453878f4..d7c495f5 100644 --- a/rl_engine/kernels/ops/pytorch/ffn/ffn.py +++ b/rl_engine/kernels/ops/pytorch/ffn/ffn.py @@ -4,6 +4,7 @@ from __future__ import annotations +import os from collections.abc import Callable from typing import Any @@ -11,9 +12,14 @@ from torch import Tensor from rl_engine.distributed.collectives import _COLLECTIVES as _SHARED_COLLECTIVES -from rl_engine.distributed.collectives import collective_for_group, deterministic_all_reduce_inplace +from rl_engine.distributed.collectives import ( + collective_for_group, + deterministic_all_reduce_inplace, + deterministic_all_reduce_staged, + deterministic_staging_reserve, +) from rl_engine.kernels.ops.base import _C, _EXT_AVAILABLE -from rl_engine.kernels.ops.cuda.matmul.det_gemm import ( +from rl_engine.kernels.ops.matmul.det_gemm import ( det_gemm_linear, det_gemm_linear_input_gradient, det_gemm_linear_weight_gradient, @@ -93,6 +99,44 @@ def _qwen3_ffn_packed_inference_fake( return rmsnorm_output.new_empty((*rmsnorm_output.shape[:-1], down_weight.shape[0])) +@torch.library.custom_op( + "rl_kernel::qwen3_ffn_packed_inference_to_staging", + mutates_args={"output"}, +) +def _qwen3_ffn_packed_inference_to_staging( + rmsnorm_output: Tensor, + fused_gate_up_weight: Tensor, + down_weight: Tensor, + output: Tensor, +) -> None: + """Run strict FFN compute with the down projection targeting IPC staging.""" + + _notify_packed_inference_observers() + hidden_2d = rmsnorm_output.reshape(-1, rmsnorm_output.shape[-1]).contiguous() + gate_up = det_gemm_linear( + hidden_2d, + fused_gate_up_weight, + native_op=_C.det_gemm_fwd_rhs_transposed, + ) + activated = _C.swiglu_packed_forward(gate_up) + det_gemm_linear( + activated, + down_weight, + native_op=_C.det_gemm_fwd_rhs_transposed, + out=output, + ) + + +@_qwen3_ffn_packed_inference_to_staging.register_fake +def _qwen3_ffn_packed_inference_to_staging_fake( + rmsnorm_output: Tensor, + fused_gate_up_weight: Tensor, + down_weight: Tensor, + output: Tensor, +) -> None: + del rmsnorm_output, fused_gate_up_weight, down_weight, output + + def qwen3_ffn_packed_inference( rmsnorm_output: Tensor, fused_gate_up_weight: Tensor, @@ -100,18 +144,58 @@ def qwen3_ffn_packed_inference( *, collective_handle: int = 0, tp_world_size: int = 1, + collective: Any | None = None, ) -> Tensor: """Inference-only packed FFN entry compatible with torch.compile.""" + if tp_world_size <= 1: + return _qwen3_ffn_packed_inference( + rmsnorm_output, + fused_gate_up_weight, + down_weight, + ) + if getattr(torch.version, "hip", None) is not None: + if collective is None: + raise RuntimeError("packed ROCm rollout FFN requires a bound TP collective") + output = _qwen3_ffn_packed_inference( + rmsnorm_output, + fused_gate_up_weight, + down_weight, + ) + return collective.all_reduce(output, out=output) + if collective_handle <= 0: + raise RuntimeError("packed rollout FFN requires a bound TP collective") + input_shape = rmsnorm_output.shape + output_shape_2d = ( + rmsnorm_output.numel() // input_shape[-1], + down_weight.shape[0], + ) + direct_output = ( + None + if collective is None + else collective.direct_staging_view(output_shape_2d, dtype=rmsnorm_output.dtype) + ) + if direct_output is not None: + deterministic_staging_reserve( + direct_output, + collective_handle=collective_handle, + ) + _qwen3_ffn_packed_inference_to_staging( + rmsnorm_output, + fused_gate_up_weight, + down_weight, + direct_output, + ) + reduced = deterministic_all_reduce_staged( + direct_output, + collective_handle=collective_handle, + ) + return reduced.reshape(*input_shape[:-1], down_weight.shape[0]) output = _qwen3_ffn_packed_inference( rmsnorm_output, fused_gate_up_weight, down_weight, ) - if tp_world_size <= 1: - return output - if collective_handle <= 0: - raise RuntimeError("packed rollout FFN requires a bound TP collective") return deterministic_all_reduce_inplace( output, collective_handle=collective_handle, @@ -634,11 +718,6 @@ def prepare_packed_inference( dist = _require_parallel_group(tp_group, "tensor") if dist is None: return 0, 1 - if getattr(torch.version, "hip", None) is not None: - raise RuntimeError( - "packed TP inference requires the native CUDA IPC collective and " - "is not available with the ROCm/RCCL transport" - ) tp_world_size = int(dist.get_world_size(group=tp_group)) if fused_gate_up_weight.size(0) % 2: raise ValueError("fused gate/up weight must contain two equal shards") @@ -654,10 +733,33 @@ def prepare_packed_inference( tp_group, min_size_bytes=min_size_bytes, ) + if getattr(torch.version, "hip", None) is not None: + collective_handle = int(collective._handle) + self._packed_inference_collectives[collective_handle] = collective + return collective_handle, tp_world_size + max_capture = int(os.getenv("RL_KERNEL_VLLM_CUDAGRAPH_MAX_CAPTURE_SIZE", "0")) + if max_capture <= 0: + raise RuntimeError("packed rollout FFN requires a positive graph capture size") + collective.prepare_direct_staging_views( + ((batch, int(down_weight.shape[0])) for batch in range(1, max_capture + 1)), + dtype=down_weight.dtype, + ) collective_handle = int(collective._handle) self._packed_inference_collectives[collective_handle] = collective return collective_handle, tp_world_size + def packed_inference_backend_id(self, collective_handle: int) -> str: + collective = self._packed_inference_collectives.get(collective_handle) + if collective is None: + raise RuntimeError("packed rollout FFN collective is not bound") + return str( + getattr( + collective, + "backend_id", + "deterministic_all_reduce.ipc_localized_fixed_tree.v1", + ) + ) + def packed_inference( self, rmsnorm_output: Tensor, @@ -673,6 +775,7 @@ def packed_inference( down_weight, collective_handle=collective_handle, tp_world_size=tp_world_size, + collective=self._packed_inference_collectives.get(collective_handle), ) def __call__( diff --git a/rl_engine/kernels/ops/rocm/attention/strict_runtime.py b/rl_engine/kernels/ops/rocm/attention/strict_runtime.py index 49a5b9a3..a1bf816d 100644 --- a/rl_engine/kernels/ops/rocm/attention/strict_runtime.py +++ b/rl_engine/kernels/ops/rocm/attention/strict_runtime.py @@ -27,7 +27,7 @@ from __future__ import annotations -from dataclasses import dataclass +from dataclasses import dataclass, replace from typing import Any import torch @@ -39,9 +39,10 @@ ) from rl_engine.kernels.ops.cuda.attention.cp_comm import ( AttentionCPBlockMetadata, + AttentionCPCommunicationUnavailable, AttentionCPCommunicationPlan, AttentionParallelSpec, - RCCLAGRSAttentionCPCommunication, + CUDAAGRSAttentionCPCommunication, ) from rl_engine.kernels.ops.cuda.attention.strict_runtime import StrictCUDAAttentionRuntime from rl_engine.kernels.ops.rocm.attention.flash_attn import StrictRocmAiterCKAttentionCore @@ -55,6 +56,37 @@ _validate_global_positions = StrictCUDAAttentionRuntime._validate_global_positions +class RCCLAGRSAttentionCPCommunication(CUDAAGRSAttentionCPCommunication): + """ROCm adapter over the shared deterministic fixed-tree collective.""" + + backend_id = "rccl_ag_rs" + collective_label = "self-owned RCCL AG/RS" + supports_autograd = True + transport_only = True + supports_async_overlap = False + supports_compute_communication_fusion = False + + def _dist(self): + import torch.distributed as dist + + if not dist.is_available() or not dist.is_initialized(): + raise AttentionCPCommunicationUnavailable( + "self-owned RCCL AG/RS requires initialized torch.distributed" + ) + return dist + + def _validate_cuda_plan(self, plan: AttentionCPCommunicationPlan) -> None: + if plan.backend != "rccl_ag_rs" or plan.status != "implemented": + raise AttentionCPCommunicationUnavailable( + "self-owned RCCL AG/RS requires an implemented rccl_ag_rs plan" + ) + replace(plan, backend="cuda_ag_rs").validate() + if torch.version.hip is None or not torch.cuda.is_available(): + raise AttentionCPCommunicationUnavailable( + "self-owned RCCL AG/RS requires an available ROCm device" + ) + + @dataclass(frozen=True) class StrictRocmAttentionResult: out: torch.Tensor diff --git a/rl_engine/kernels/ops/rocm/loss/vocab_parallel_logp.py b/rl_engine/kernels/ops/rocm/loss/vocab_parallel_logp.py index bfabe8b1..a1f3578f 100644 --- a/rl_engine/kernels/ops/rocm/loss/vocab_parallel_logp.py +++ b/rl_engine/kernels/ops/rocm/loss/vocab_parallel_logp.py @@ -13,9 +13,10 @@ * ``hip_deterministic_logp_backward`` produces ``grad_logits`` for the selected logprob and LSE outputs in one fused pass from the saved input shard. -``apply`` runs through the shared :func:`apply_with_kernels` path; -``apply_with_entropy`` keeps the shared autograd path (with the HIP tile -kernel) because the entropy gradient needs the full probability tensor anyway. +``apply`` keeps the fused ROCm forward/backward local while reusing the shared +contract validation, rank-ordered transport, and fixed tile merge helpers. +``apply_with_entropy`` keeps the inherited reference autograd path because the +entropy gradient needs the full probability tensor anyway. """ from __future__ import annotations @@ -24,11 +25,17 @@ import torch -from rl_engine.kernels.logprob_contract import LogprobContract +from rl_engine.kernels.logprob_contract import LogprobContract, LogprobContractError from rl_engine.kernels.ops.pytorch.loss.vocab_parallel_logp import ( DEFAULT_NUM_VOCAB_TILES, VocabParallelLogprobOp, - apply_with_kernels, + _gather_target_logit, + _gather_tile_stats, + _merge_tile_partials, + _preflight_cross_rank_agreement, + _tile_size, + _validate_active_targets, + _validate_invocation, ) BACKEND_ID = "rocm-vocab-parallel-logp-ws2" @@ -69,13 +76,138 @@ def backward( ) +class _RocmVocabParallelLogprobFunction(torch.autograd.Function): + """ROCm tile statistics and backward with the shared WS2 merge contract.""" + + @staticmethod + def forward(ctx, local_logits, target_1d, active_mask, contract, tp_group, tile): + sharding = contract.sharding + shard = local_logits.contiguous() + local_tiles = sharding.local_vocab_size // tile + if local_tiles <= 0: + raise RuntimeError("native tile stats require at least one local vocab tile") + local_m, local_s = _HipKernels.tile_stats( + shard, + sharding.local_vocab_start, + sharding.real_vocab_size, + local_tiles, + ) + tile_counts = [ + (end - start) // tile for start, end in sharding.vocab_shard_bounds + ] + m_all, s_all = _gather_tile_stats( + local_m.contiguous(), + local_s.contiguous(), + contract, + tp_group, + tile_counts, + ) + safe_target = torch.where(active_mask, target_1d, torch.zeros_like(target_1d)) + target_logit = _gather_target_logit( + shard, safe_target, contract, tp_group + ).float() + lse = _merge_tile_partials(m_all, s_all) + selected_logp = torch.where( + active_mask, target_logit - lse, torch.zeros_like(lse) + ) + + ctx.save_for_backward(shard, lse, safe_target, active_mask) + ctx.local_vocab_start = sharding.local_vocab_start + ctx.real_vocab_size = sharding.real_vocab_size + ctx.set_materialize_grads(False) + return selected_logp, lse + + @staticmethod + def backward(ctx, grad_logp, grad_lse): + if not ctx.needs_input_grad[0] or (grad_logp is None and grad_lse is None): + return None, None, None, None, None, None + shard, lse, safe_target, active_mask = ctx.saved_tensors + rows, local_vocab = shard.shape + start = ctx.local_vocab_start + if grad_logp is not None: + coef_logp = ( + torch.where(active_mask, grad_logp, torch.zeros_like(grad_logp)) + .float() + .contiguous() + ) + owns = (safe_target >= start) & (safe_target < start + local_vocab) + target_local = torch.where( + owns & active_mask, + safe_target - start, + torch.full_like(safe_target, -1), + ).contiguous() + else: + coef_logp = lse.new_zeros((rows,)) + target_local = torch.full( + (rows,), -1, dtype=torch.long, device=shard.device + ) + has_lse_grad = grad_lse is not None + coef_lse = ( + grad_lse.float().contiguous() + if has_lse_grad + else lse.new_zeros((rows,)) + ) + grad = _HipKernels.backward( + shard, + lse.contiguous(), + coef_logp, + coef_lse, + target_local, + start, + ctx.real_vocab_size, + has_lse_grad, + ) + return grad, None, None, None, None, None + + +def _apply_with_kernels( + local_logits: torch.Tensor, + target_ids: torch.Tensor, + *, + contract: LogprobContract, + tp_group: Any, + num_vocab_tiles: int, + validate: bool, +) -> tuple[torch.Tensor, torch.Tensor]: + if not isinstance(contract, LogprobContract): + raise LogprobContractError("contract must be a LogprobContract") + tile = _tile_size(contract, num_vocab_tiles) + _validate_invocation(local_logits, target_ids, contract, tp_group) + + target_1d = target_ids.reshape(-1).to( + device=local_logits.device, dtype=torch.long + ) + active_mask = torch.tensor( + contract.mask.active_mask, + dtype=torch.bool, + device=local_logits.device, + ) + if validate: + _validate_active_targets( + target_1d, active_mask, contract.sharding.real_vocab_size + ) + if contract.sharding.tp_world_size > 1: + _preflight_cross_rank_agreement( + contract, tp_group, num_vocab_tiles, True + ) + + selected_logp, lse = _RocmVocabParallelLogprobFunction.apply( + local_logits, target_1d, active_mask, contract, tp_group, tile + ) + if validate and bool((~torch.isfinite(lse) & active_mask).any().item()): + raise LogprobContractError( + "non-finite logsumexp on an active row: logits over the real " + "vocabulary must be finite for every active token" + ) + return selected_logp, lse + + class RocmVocabParallelLogprobOp(VocabParallelLogprobOp): """Contract-preserving ROCm implementation with HIP local reductions.""" op_class = "logprob" is_batch_invariant = True backend_id = BACKEND_ID - use_native_tile_stats = True def apply( self, @@ -103,14 +235,13 @@ def apply( f"{BACKEND_ID} requires rl_engine._C built with a ROCm toolchain " "(hip_deterministic_logp_* symbols are missing); it does not fall back" ) - return apply_with_kernels( + return _apply_with_kernels( local_logits, target_ids, contract=contract, tp_group=tp_group, num_vocab_tiles=num_vocab_tiles, validate=validate, - kernels=_HipKernels, ) diff --git a/rl_engine/kernels/ops/rocm/matmul/__init__.py b/rl_engine/kernels/ops/rocm/matmul/__init__.py new file mode 100644 index 00000000..f0c5f23d --- /dev/null +++ b/rl_engine/kernels/ops/rocm/matmul/__init__.py @@ -0,0 +1,5 @@ +"""ROCm matrix multiplication operators.""" + +from .det_gemm import DetGemmOp, RocmDetGemmOp, deterministic_gemm + +__all__ = ["DetGemmOp", "RocmDetGemmOp", "deterministic_gemm"] diff --git a/rl_engine/kernels/ops/rocm/matmul/det_gemm.py b/rl_engine/kernels/ops/rocm/matmul/det_gemm.py new file mode 100644 index 00000000..6fed72b0 --- /dev/null +++ b/rl_engine/kernels/ops/rocm/matmul/det_gemm.py @@ -0,0 +1,220 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors +"""Strict deterministic GEMM facade for ROCm.""" + +from __future__ import annotations + +import os +from collections.abc import Callable +from threading import Lock + +import torch + +from rl_engine.kernels.ops.backward_runtime import record_backward +from rl_engine.kernels.ops.triton.matmul.det_gemm import ( + TritonDetGemmOp, + _triton_gemm_fp32, + _triton_tree_gemm, + deterministic_gemm_triton, +) +from rl_engine.runtime_mode import rl_kernel_mode, route_report_enabled + +_BACKEND_ENV = "RL_KERNEL_DET_GEMM_BACKEND" +_AUTO_BACKEND = "auto" +_TRITON_BACKEND = "triton" +_ROUTE_REPORTED = False +_ROUTE_REPORT_LOCK = Lock() + + +def _requested_det_gemm_backend() -> str: + value = os.getenv(_BACKEND_ENV, _AUTO_BACKEND).strip().lower() + value = {"rocm": _TRITON_BACKEND}.get(value, value) + if value not in {_AUTO_BACKEND, _TRITON_BACKEND}: + raise RuntimeError( + f"{_BACKEND_ENV} must be '{_AUTO_BACKEND}' or " + f"'{_TRITON_BACKEND}' on ROCm, got {value!r}" + ) + return value + + +_REQUESTED_BACKEND = _requested_det_gemm_backend() + + +def det_gemm_backend() -> str: + """Return the strict ROCm GEMM implementation.""" + + return _TRITON_BACKEND + + +def det_gemm_fallback_reason() -> str | None: + return None + + +def det_gemm_backend_id() -> str: + return "rlkernel.det_gemm.triton_tree_rocm.v1" + + +def _report_strict_route_once() -> None: + global _ROUTE_REPORTED + if torch._dynamo.is_compiling() or not route_report_enabled(): + return + with _ROUTE_REPORT_LOCK: + if _ROUTE_REPORTED: + return + _ROUTE_REPORTED = True + print( + f"[RL-Kernel][route] mode={rl_kernel_mode().value} module=gemm " + f"requested={_REQUESTED_BACKEND} " + f"actual={det_gemm_backend_id()} fallback=false", + flush=True, + ) + + +def det_gemm_linear( + a: torch.Tensor, + weight: torch.Tensor, + *, + native_op: Callable[[torch.Tensor, torch.Tensor], torch.Tensor] | None = None, + out: torch.Tensor | None = None, +) -> torch.Tensor: + """Apply a native [N,K] weight through the strict ROCm backend.""" + + del native_op + return _triton_tree_gemm(a, weight.t().contiguous(), out=out) + + +def det_gemm_linear_input_gradient( + grad_output: torch.Tensor, + weight: torch.Tensor, + *, + native_op: Callable[[torch.Tensor, torch.Tensor], torch.Tensor] | None = None, +) -> torch.Tensor: + """Compute ``dX = dY @ weight`` through the strict ROCm backend.""" + + del native_op + return deterministic_gemm_triton(grad_output, weight) + + +def det_gemm_linear_weight_gradient( + a: torch.Tensor, + grad_output: torch.Tensor, + *, + native_op: Callable[[torch.Tensor, torch.Tensor], torch.Tensor] | None = None, +) -> torch.Tensor: + """Compute ``dWeight = dY.T @ X`` through the strict ROCm backend.""" + + del native_op + return _triton_tree_gemm( + a.t(), + grad_output, + transpose_output=True, + preserve_a_strides=True, + ) + + +class _DetLinearFn(torch.autograd.Function): + @staticmethod + def forward(ctx, a, weight): + ctx.save_for_backward(a, weight) + return det_gemm_linear(a, weight) + + @staticmethod + def backward(ctx, grad_out): + a, weight = ctx.saved_tensors + grad_out = grad_out.contiguous() + if grad_out.dtype != torch.bfloat16: + grad_out = grad_out.to(torch.bfloat16) + da = ( + det_gemm_linear_input_gradient(grad_out, weight) + if ctx.needs_input_grad[0] + else None + ) + dweight = ( + det_gemm_linear_weight_gradient(a, grad_out) + if ctx.needs_input_grad[1] + else None + ) + record_backward( + "det_gemm", + kernel_id=det_gemm_backend_id(), + impl="strict_det_gemm", + family="rocm", + ) + return da, dweight + + +class RocmDetGemmOp: + """Batch-invariant deterministic GEMM backed by the ROCm Triton tree.""" + + def __init__(self): + det_gemm_backend() + self._triton = TritonDetGemmOp() + self.op = self._triton + self.has_hardware_op = True + _report_strict_route_once() + + def __call__(self, a: torch.Tensor, b: torch.Tensor) -> torch.Tensor: + assert a.dtype == torch.bfloat16 and b.dtype == torch.bfloat16, "BF16 only" + assert a.is_cuda and b.is_cuda, "Inputs must be on ROCm device" + return deterministic_gemm_triton(a.contiguous(), b.contiguous()) + + def linear( + self, + a: torch.Tensor, + weight: torch.Tensor, + *, + out: torch.Tensor | None = None, + ) -> torch.Tensor: + """Apply a native [N,K] linear weight without changing the GEMM tree.""" + + assert a.dtype == torch.bfloat16 and weight.dtype == torch.bfloat16, "BF16 only" + assert a.is_cuda and weight.is_cuda, "Inputs must be on ROCm device" + a = a.contiguous() + weight = weight.contiguous() + if out is not None: + if torch.is_grad_enabled() and (a.requires_grad or weight.requires_grad): + raise RuntimeError("direct-output deterministic GEMM is inference-only") + return det_gemm_linear(a, weight, out=out) + return _DetLinearFn.apply(a, weight) + + def forward_fp32(self, a: torch.Tensor, b: torch.Tensor) -> torch.Tensor: + assert a.dtype == torch.bfloat16 and b.dtype == torch.bfloat16, "BF16 only" + assert a.is_cuda and b.is_cuda, "Inputs must be on ROCm device" + return _triton_gemm_fp32(a, b) + + def forward_accum_fp32(self, a: torch.Tensor, b: torch.Tensor) -> torch.Tensor: + if a.dtype not in (torch.bfloat16, torch.float32) or b.dtype not in ( + torch.bfloat16, + torch.float32, + ): + raise TypeError("FP32-accumulation GEMM requires BF16 or FP32 inputs") + assert a.is_cuda and b.is_cuda, "Inputs must be on ROCm device" + return _triton_gemm_fp32(a, b).to(a.dtype) + + def parameter_vjp_contributions_fp32(self, *, a, b, grad_output): + del b + rows_a = a.float() + rows_g = grad_output.float() + return {"b": rows_a[:, :, None] * rows_g[:, None, :]} + + +DetGemmOp = RocmDetGemmOp + + +def deterministic_gemm(a: torch.Tensor, b: torch.Tensor) -> torch.Tensor: + """Functional strict ROCm GEMM entry.""" + + return deterministic_gemm_triton(a, b) + + +__all__ = [ + "DetGemmOp", + "RocmDetGemmOp", + "det_gemm_backend", + "det_gemm_backend_id", + "det_gemm_fallback_reason", + "det_gemm_linear", + "det_gemm_linear_input_gradient", + "det_gemm_linear_weight_gradient", + "deterministic_gemm", +] diff --git a/rl_engine/kernels/ops/rocm/rotary_embedding/__init__.py b/rl_engine/kernels/ops/rocm/rotary_embedding/__init__.py new file mode 100644 index 00000000..d2d6879e --- /dev/null +++ b/rl_engine/kernels/ops/rocm/rotary_embedding/__init__.py @@ -0,0 +1,5 @@ +"""ROCm rotary embedding operators.""" + +from .rope import RocmDeterministicRoPEOp + +__all__ = ["RocmDeterministicRoPEOp"] diff --git a/rl_engine/kernels/ops/rocm/rotary_embedding/rope.py b/rl_engine/kernels/ops/rocm/rotary_embedding/rope.py new file mode 100644 index 00000000..6164523f --- /dev/null +++ b/rl_engine/kernels/ops/rocm/rotary_embedding/rope.py @@ -0,0 +1,77 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors +"""Precompiled deterministic ROCm RoPE matching the shared reference layout.""" + +from __future__ import annotations + +import torch +from torch import Tensor + +from rl_engine.kernels.ops.base import _C, _EXT_AVAILABLE +from rl_engine.kernels.ops.cuda.rotary_embedding.rope import _restore_rope, _rope_table + + +class _RocmRoPEFunction(torch.autograd.Function): + @staticmethod + def forward(ctx, x: Tensor, positions: Tensor, theta: float) -> Tensor: + x_2d, cos, sin = _rope_table(x, positions, theta) + ctx.save_for_backward(cos, sin) + ctx.x_shape = tuple(x.shape) + ctx.pos_dim = positions.dim() + out_2d = _C.deterministic_rope_apply_rocm(x_2d, cos, sin, 1.0) + return _restore_rope(out_2d, x, positions) + + @staticmethod + def backward(ctx, grad_out: Tensor): + cos, sin = ctx.saved_tensors + grad_x = None + if ctx.needs_input_grad[0]: + if ctx.pos_dim == 2 and len(ctx.x_shape) == 4: + g_2d = grad_out.permute(1, 0, 2, 3).contiguous().reshape(-1, ctx.x_shape[-1]) + out_2d = _C.deterministic_rope_apply_rocm(g_2d, cos, sin, -1.0) + heads, batch, seq, dim = ( + ctx.x_shape[1], + ctx.x_shape[0], + ctx.x_shape[2], + ctx.x_shape[3], + ) + grad_x = ( + out_2d.reshape(heads, batch, seq, dim) + .permute(1, 0, 2, 3) + .contiguous() + ) + else: + g_2d = grad_out.contiguous().reshape(-1, grad_out.shape[-1]) + grad_x = _C.deterministic_rope_apply_rocm(g_2d, cos, sin, -1.0).reshape( + grad_out.shape + ) + return grad_x, None, None + + +class RocmDeterministicRoPEOp: + """Precompiled HIP RoPE path shared by ROCm training and rollout.""" + + backend_id = "rlkernel.rocm.deterministic_rope" + op_class = "elementwise" + fallback = False + + def __init__(self) -> None: + if torch.version.hip is None: + raise RuntimeError("RocmDeterministicRoPEOp requires a ROCm PyTorch build") + if not _EXT_AVAILABLE or not hasattr(_C, "deterministic_rope_apply_rocm"): + raise RuntimeError( + "ROCm deterministic RoPE is unavailable; rebuild rl_engine._C for ROCm" + ) + + def __call__(self, x: Tensor, positions: Tensor, *, theta: float = 1_000_000.0) -> Tensor: + return self.forward(x, positions, theta=theta) + + def forward(self, x: Tensor, positions: Tensor, *, theta: float = 1_000_000.0) -> Tensor: + if not x.is_cuda: + raise RuntimeError("ROCm deterministic RoPE requires a GPU tensor") + if x.dtype not in (torch.float16, torch.bfloat16): + raise ValueError("ROCm deterministic RoPE requires FP16 or BF16") + return _RocmRoPEFunction.apply(x, positions, theta) + + +__all__ = ["RocmDeterministicRoPEOp"] diff --git a/rl_engine/kernels/ops/triton/matmul/det_gemm.py b/rl_engine/kernels/ops/triton/matmul/det_gemm.py index 078167b0..68db1379 100644 --- a/rl_engine/kernels/ops/triton/matmul/det_gemm.py +++ b/rl_engine/kernels/ops/triton/matmul/det_gemm.py @@ -35,6 +35,11 @@ # Pinned. NOT autotuned (autotune picks per-shape configs -> breaks invariance). _BLOCK_M, _BLOCK_N, _BLOCK_K = 64, 64, 32 +# Keep the intermediate tree workspace bounded. A VLLM warmup can present a +# very large flattened token dimension; allocating ``node_count * M * N`` in +# one shot would otherwise create a multi-terabyte virtual tensor on ROCm. +# Chunking over M preserves the exact K-tree and BF16 rounding contract. +_MAX_TREE_WORKSPACE_ELEMENTS = 128 * 1024 * 1024 _TREE_PLANS: dict[tuple[int, int], "_DeviceTreePlan"] = {} _TREE_PLAN_LOCK = threading.Lock() @@ -277,9 +282,12 @@ def _det_gemm_tree_leaf_kernel( pid_n = tl.program_id(2) offs_m = pid_m * BLOCK_M + tl.arange(0, BLOCK_M) offs_n = pid_n * BLOCK_N + tl.arange(0, BLOCK_N) - leaf_start = tl.load(leaf_starts_ptr + leaf) - leaf_length = tl.load(leaf_lengths_ptr + leaf) - leaf_node = tl.load(leaf_nodes_ptr + leaf) + # Qwen prefill workspaces can exceed 2**31 elements. ROCm Triton + # otherwise keeps the index expression in i32 and wraps addresses, + # causing GPU memory faults for large M*N trees. + leaf_start = tl.load(leaf_starts_ptr + leaf).to(tl.int64) + leaf_length = tl.load(leaf_lengths_ptr + leaf).to(tl.int64) + leaf_node = tl.load(leaf_nodes_ptr + leaf).to(tl.int64) acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32) # Keep the leaf's ascending scalar FMA order identical to the native # gfx942 correctness kernel. A tl.dot/MFMA leaf is topology-stable but @@ -298,7 +306,11 @@ def _det_gemm_tree_leaf_kernel( other=0.0, ).to(tl.float32) acc += a[:, None] * b[None, :] - output_offsets = leaf_node * M * N + offs_m[:, None] * N + offs_n[None, :] + output_offsets = ( + leaf_node * (M * N) + + offs_m[:, None].to(tl.int64) * N + + offs_n[None, :].to(tl.int64) + ) output_mask = (offs_m[:, None] < M) & (offs_n[None, :] < N) tl.store( workspace_ptr + output_offsets, @@ -306,6 +318,10 @@ def _det_gemm_tree_leaf_kernel( mask=output_mask, ) + # Preserve the WS1 manifest's historical kernel symbol while retaining the + # more descriptive implementation name used by the optimized tree path. + _det_gemm_kernel = _det_gemm_tree_leaf_kernel + @triton.jit def _det_gemm_tree_reduce_kernel( workspace_ptr, @@ -318,12 +334,12 @@ def _det_gemm_tree_reduce_kernel( ): operation = tl.program_id(0) block = tl.program_id(1) - offsets = block * BLOCK + tl.arange(0, BLOCK) + offsets = (block * BLOCK + tl.arange(0, BLOCK)).to(tl.int64) elements = M * N mask = offsets < elements - lower_node = tl.load(lower_nodes_ptr + operation) - upper_node = tl.load(upper_nodes_ptr + operation) - output_node = tl.load(output_nodes_ptr + operation) + lower_node = tl.load(lower_nodes_ptr + operation).to(tl.int64) + upper_node = tl.load(upper_nodes_ptr + operation).to(tl.int64) + output_node = tl.load(output_nodes_ptr + operation).to(tl.int64) lower = tl.load( workspace_ptr + lower_node * elements + offsets, mask=mask, @@ -349,7 +365,7 @@ def _copy_tree_root_kernel( elements, BLOCK: tl.constexpr, ): - offsets = tl.program_id(0) * BLOCK + tl.arange(0, BLOCK) + offsets = (tl.program_id(0) * BLOCK + tl.arange(0, BLOCK)).to(tl.int64) mask = offsets < elements values = tl.load(workspace_ptr + root * elements + offsets, mask=mask) tl.store(output_ptr + offsets, values, mask=mask) @@ -364,8 +380,8 @@ def _copy_tree_root_transposed_kernel( BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr, ): - offsets_m = tl.program_id(0) * BLOCK_M + tl.arange(0, BLOCK_M) - offsets_n = tl.program_id(1) * BLOCK_N + tl.arange(0, BLOCK_N) + offsets_m = (tl.program_id(0) * BLOCK_M + tl.arange(0, BLOCK_M)).to(tl.int64) + offsets_n = (tl.program_id(1) * BLOCK_N + tl.arange(0, BLOCK_N)).to(tl.int64) elements = M * N mask = (offsets_m[:, None] < M) & (offsets_n[None, :] < N) values = tl.load( @@ -457,6 +473,21 @@ def _triton_gemm_fp32(a: torch.Tensor, b: torch.Tensor) -> torch.Tensor: return c +def _triton_gemm( + a: torch.Tensor, + b: torch.Tensor, + *, + output_dtype: torch.dtype | None = None, +) -> torch.Tensor: + """Compatibility entry point for canonical Triton linear operators.""" + + if output_dtype is None or output_dtype == torch.bfloat16: + return _triton_tree_gemm(a, b) + if output_dtype == torch.float32: + return _triton_gemm_fp32(a, b) + raise TypeError(f"unsupported deterministic Triton GEMM output dtype: {output_dtype}") + + def _triton_tree_gemm( a: torch.Tensor, b: torch.Tensor, @@ -484,6 +515,24 @@ def _triton_tree_gemm( b = b.contiguous() m_size, k_size = a.shape n_size = b.size(1) + plan = _device_tree_plan(k_size, a.device) + workspace_elements = plan.host.node_count * m_size * n_size + if workspace_elements > _MAX_TREE_WORKSPACE_ELEMENTS and out is None: + rows_per_chunk = max( + 1, + _MAX_TREE_WORKSPACE_ELEMENTS // (plan.host.node_count * n_size), + ) + chunks = [] + for start in range(0, m_size, rows_per_chunk): + stop = min(m_size, start + rows_per_chunk) + chunk = _triton_tree_gemm( + a[start:stop], + b, + transpose_output=transpose_output, + preserve_a_strides=preserve_a_strides, + ) + chunks.append(chunk) + return torch.cat(chunks, dim=1 if transpose_output else 0) result_shape = (n_size, m_size) if transpose_output else (m_size, n_size) if out is None: result = torch.empty(result_shape, dtype=torch.bfloat16, device=a.device) @@ -504,7 +553,6 @@ def _triton_tree_gemm( if out.requires_grad: raise ValueError("Triton tree GEMM output buffer must not require gradients") result = out - plan = _device_tree_plan(k_size, a.device) workspace = torch.empty( (plan.host.node_count, m_size, n_size), dtype=torch.bfloat16, @@ -525,7 +573,7 @@ def _triton_tree_gemm( if leaf_config.n_fastest else (len(plan.host.leaf_nodes), tiles_m, tiles_n) ) - _det_gemm_tree_leaf_kernel[leaf_grid]( + _det_gemm_kernel[leaf_grid]( a, b, workspace, diff --git a/rl_engine/kernels/registry.py b/rl_engine/kernels/registry.py index 23286d94..7070728a 100644 --- a/rl_engine/kernels/registry.py +++ b/rl_engine/kernels/registry.py @@ -301,7 +301,7 @@ def _default_semantic_descriptors() -> tuple[OperatorBackendDescriptor, ...]: semantic_op="ffn", backend_id="rlkernel.ffn.qwen3.deterministic.v1", supported_targets=frozenset({"rollout", "training"}), - supported_devices=frozenset({"cuda"}), + supported_devices=frozenset({"cuda", "rocm"}), supported_dtypes=frozenset({"bfloat16"}), supported_topologies={"*": "*"}, determinism_or_alignment_properties={ diff --git a/tests/distributed/test_transport_deterministic_collective.py b/tests/distributed/test_transport_deterministic_collective.py index 5c9ffb84..23621ee0 100644 --- a/tests/distributed/test_transport_deterministic_collective.py +++ b/tests/distributed/test_transport_deterministic_collective.py @@ -10,6 +10,7 @@ import rl_engine.distributed as distributed import rl_engine.distributed.collectives as collectives +import rl_engine.distributed.rocm_collectives as rocm_collectives from rl_engine.distributed import ( RCCLDeterministicCollective, TorchDistributedDeterministicCollective, @@ -121,7 +122,7 @@ def _make_collective( peer_signatures=peer_signatures, peer_capacities=peer_capacities, ) - monkeypatch.setattr(collectives, "dist", fake_dist) + monkeypatch.setattr(rocm_collectives, "dist", fake_dist) collective = TorchDistributedDeterministicCollective( group=object(), device="cpu", @@ -332,7 +333,7 @@ def test_reduce_scatter_many_rejects_oversized_packed_input( peers, max_size_bytes=32, ) - monkeypatch.setattr(collectives, "_PACKED_REDUCE_SCATTER_MAX_BYTES", 1024) + monkeypatch.setattr(rocm_collectives, "_PACKED_REDUCE_SCATTER_MAX_BYTES", 1024) with pytest.raises(ValueError, match="packed input requires"): collective.reduce_scatter_many((peers[0], peers[0])) @@ -344,7 +345,7 @@ def test_reduce_scatter_many_uses_separate_calls_for_large_payloads( ) -> None: peers = [torch.ones(4, 2, dtype=torch.float32) for _ in range(2)] collective, fake_dist = _make_collective(monkeypatch, peers) - monkeypatch.setattr(collectives, "_PACKED_REDUCE_SCATTER_MAX_BYTES", 1) + monkeypatch.setattr(rocm_collectives, "_PACKED_REDUCE_SCATTER_MAX_BYTES", 1) outputs = collective.reduce_scatter_many((peers[0], peers[0])) @@ -422,14 +423,14 @@ def test_unsupported_world_size_is_rejected( world_size: int, ) -> None: fake_dist = _FakeDistributed([torch.ones(1)] * world_size) - monkeypatch.setattr(collectives, "dist", fake_dist) + monkeypatch.setattr(rocm_collectives, "dist", fake_dist) with pytest.raises(ValueError, match="world_size in"): TorchDistributedDeterministicCollective(group=object(), device="cpu") def test_rccl_class_requires_rocm_build(monkeypatch: pytest.MonkeyPatch) -> None: - monkeypatch.setattr(collectives.torch.version, "hip", None, raising=False) + monkeypatch.setattr(rocm_collectives.torch.version, "hip", None, raising=False) with pytest.raises(RuntimeError, match="ROCm PyTorch build"): RCCLDeterministicCollective(group=object(), device="cuda:0") @@ -437,10 +438,10 @@ def test_rccl_class_requires_rocm_build(monkeypatch: pytest.MonkeyPatch) -> None def test_rccl_class_requires_nccl_process_group(monkeypatch: pytest.MonkeyPatch) -> None: fake_dist = _FakeDistributed([torch.ones(1)], backend="gloo") - monkeypatch.setattr(collectives, "dist", fake_dist) - monkeypatch.setattr(collectives.torch.version, "hip", "6.3", raising=False) - monkeypatch.setattr(collectives.torch.cuda, "is_available", lambda: True) - monkeypatch.setattr(collectives.torch.cuda, "current_device", lambda: 0) + monkeypatch.setattr(rocm_collectives, "dist", fake_dist) + monkeypatch.setattr(rocm_collectives.torch.version, "hip", "6.3", raising=False) + monkeypatch.setattr(rocm_collectives.torch.cuda, "is_available", lambda: True) + monkeypatch.setattr(rocm_collectives.torch.cuda, "current_device", lambda: 0) with pytest.raises(RuntimeError, match="NCCL process-group API"): RCCLDeterministicCollective(group=object(), device="cuda:0") @@ -449,8 +450,8 @@ def test_rccl_class_requires_nccl_process_group(monkeypatch: pytest.MonkeyPatch) def test_rccl_class_rejects_cpu_before_process_group_exchange( monkeypatch: pytest.MonkeyPatch, ) -> None: - monkeypatch.setattr(collectives.torch.version, "hip", "6.3", raising=False) - monkeypatch.setattr(collectives.torch.cuda, "is_available", lambda: True) + monkeypatch.setattr(rocm_collectives.torch.version, "hip", "6.3", raising=False) + monkeypatch.setattr(rocm_collectives.torch.cuda, "is_available", lambda: True) with pytest.raises(ValueError, match="ROCm device"): RCCLDeterministicCollective(group=object(), device="cpu") diff --git a/tests/test_framework_runtime_adapters.py b/tests/test_framework_runtime_adapters.py index 5e140b2a..34d92e48 100644 --- a/tests/test_framework_runtime_adapters.py +++ b/tests/test_framework_runtime_adapters.py @@ -1,777 +1,934 @@ -# SPDX-License-Identifier: Apache-2.0 -# Copyright (c) 2026 RL-Kernel Contributors - -from __future__ import annotations - -import ast -import os -import sys -from collections import namedtuple -from dataclasses import dataclass -from pathlib import Path -from types import ModuleType, SimpleNamespace - -import pytest -import torch - -import rl_engine.integrations.framework_operators as framework_operators -from rl_engine.integrations.ablation import ( - IntegrationPlan, - configure_integration_environment, - integration_plan_from_environment, -) -from rl_engine.integrations.framework_operators import ( - MegatronAttentionOperator, - SemanticOperatorHandle, - VllmAttentionOperator, - VllmLogpOperator, - _megatron_zigzag_layout, - _packed_local_sequence_layout, - _vllm_kv_cache_views, -) +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors + +from __future__ import annotations + +import ast +import os +import sys +from collections import namedtuple +from dataclasses import dataclass +from pathlib import Path +from types import ModuleType, SimpleNamespace + +import pytest +import torch + +import rl_engine.integrations.framework_operators as framework_operators +from rl_engine.integrations.ablation import ( + IntegrationPlan, + configure_integration_environment, + integration_plan_from_environment, +) +from rl_engine.integrations.framework_operators import ( + MegatronAttentionOperator, + SemanticOperatorHandle, + VllmAttentionOperator, + VllmLogpOperator, + _megatron_zigzag_layout, + _packed_local_sequence_layout, + _vllm_kv_cache_views, +) from rl_engine.integrations.megatron_runtime import ( + _deterministic_reduce_from_tensor_model_parallel_region, + _install_torch_dist_object_compatibility, _patch_strict_attention_projections, install_megatron_integration, ) -from rl_engine.integrations.runtime import FrameworkOperatorIntegration -from rl_engine.integrations.state import clear_active_integration -from rl_engine.integrations.vllm_runtime import ( - _patch_qwen3_strict_model, - _register_attention_backend, - configure_vllm_environment, -) -from rl_engine.kernels.attention_contract import ( - STRICT_ATTENTION_FA4_SCHEDULE_ID, - STRICT_ATTENTION_PRODUCTION_CORE_ID, - AttentionContract, - AttentionDType, - AttentionMode, - AttentionRole, - ReductionSpec, - ShardingSpec, +from rl_engine.integrations.runtime import FrameworkOperatorIntegration +from rl_engine.integrations.state import clear_active_integration +from rl_engine.integrations.vllm_runtime import ( + _patch_qwen3_strict_model, + _register_attention_backend, + configure_vllm_environment, +) +from rl_engine.kernels.attention_contract import ( + STRICT_ATTENTION_FA4_SCHEDULE_ID, + STRICT_ATTENTION_PRODUCTION_CORE_ID, + AttentionContract, + AttentionDType, + AttentionMode, + AttentionRole, + ReductionSpec, + ShardingSpec, +) +from rl_engine.kernels.ops.cuda.attention.strict_runtime import ( + StrictCUDAAttentionRuntime, ) -from rl_engine.kernels.ops.cuda.attention.strict_runtime import StrictCUDAAttentionRuntime - - -def test_framework_adapters_do_not_construct_registered_kernels_directly(): - source_path = ( - Path(__file__).parents[1] / "rl_engine" / "integrations" / "framework_operators.py" - ) - tree = ast.parse(source_path.read_text(encoding="utf-8")) - forbidden = { - "AttentionAblationOp", - "DeterministicCPAttentionReferenceOp", - "Qwen3FFNOp", - "StrictCUDAAttentionRuntime", - "VocabParallelLogprobOp", - } - constructed = { - node.func.id - for node in ast.walk(tree) - if isinstance(node, ast.Call) and isinstance(node.func, ast.Name) - } - - assert constructed.isdisjoint(forbidden) - - -def test_semantic_handle_uses_operator_bridge_and_exposes_instance_provenance(): - handle = SemanticOperatorHandle( - target="training", - semantic_op="attention", - backend_id="rlkernel.attention.deterministic.v1", - ) - tensor = torch.zeros(1, 1, 1, 8) - - first = handle.get( - tensor, - topology={ - "world_size": 1, - "tensor_parallel_size": 1, - "context_parallel_size": 1, - }, - ) - second = handle.get( - tensor, - topology={ - "world_size": 1, - "tensor_parallel_size": 1, - "context_parallel_size": 1, - }, - ) - - assert first is second - assert first.backend_id == "rlkernel.attention.deterministic.v1" - assert handle.provenance is not None - assert handle.provenance["semantic_op"] == "attention" - assert handle.provenance["backend_id"] == first.backend_id - - -def test_plan_environment_is_shared_by_both_framework_installers(monkeypatch, tmp_path): - for variable in ( - "RL_KERNEL_ATTENTION_CASE", - "RL_KERNEL_FFN_CASE", - "RL_KERNEL_LOGP_CASE", - "RL_KERNEL_READBACK_DIR", - ): - monkeypatch.setenv(variable, "previous") - plan = IntegrationPlan.from_case_ids(attention="P/R", ffn="R/P", logp="R/R") - configure_integration_environment(plan, readback_dir=str(tmp_path)) - - assert integration_plan_from_environment() == plan - assert Path(os.environ["RL_KERNEL_READBACK_DIR"]) == tmp_path -def test_vllm_rlkernel_attention_overrides_selected_flash_attn_backend( - monkeypatch, -): - monkeypatch.delenv("VLLM_ATTENTION_BACKEND", raising=False) - plan = IntegrationPlan.from_case_ids(attention="P/R") - - configure_vllm_environment(plan) - - assert os.environ["VLLM_ATTENTION_BACKEND"] == "FLASH_ATTN" - - -def test_vllm_rocm_attention_selects_aiter_metadata_backend(monkeypatch): - monkeypatch.delenv("VLLM_ATTENTION_BACKEND", raising=False) - monkeypatch.setattr(torch.version, "hip", "7.1") - plan = IntegrationPlan.from_case_ids(attention="P/R") - - configure_vllm_environment(plan) - - assert os.environ["VLLM_ATTENTION_BACKEND"] == "ROCM_AITER_FA" - - -def test_vllm_rocm_registration_wraps_aiter_backend(monkeypatch): - selected = [] - - class BackendEnum: - ROCM_AITER_FA = "ROCM_AITER_FA" - FLASH_ATTN = "FLASH_ATTN" - - class AiterImpl: - def __init__(self, *args, **kwargs): - del args, kwargs - - def forward(self, *args, **kwargs): - return args, kwargs - - class AiterBuilder: - pass - - class AiterBackend: - pass - - registry = ModuleType("vllm.v1.attention.backends.registry") - registry.AttentionBackendEnum = BackendEnum - registry.register_backend = lambda backend, path: selected.append((backend, path)) - aiter = ModuleType("vllm.v1.attention.backends.rocm_aiter_fa") - aiter.AiterFlashAttentionBackend = AiterBackend - aiter.AiterFlashAttentionImpl = AiterImpl - aiter.AiterFlashAttentionMetadataBuilder = AiterBuilder - monkeypatch.setitem(sys.modules, registry.__name__, registry) - monkeypatch.setitem(sys.modules, aiter.__name__, aiter) - monkeypatch.setattr(torch.version, "hip", "7.1") - - plan = IntegrationPlan.from_case_ids(attention="P/R") - - class Integration: - def __init__(self): - self.plan = plan - self.installed = {} - self.hooks = [] +def test_torch_dist_object_compatibility_deserializes_scalar_bytes_io(monkeypatch): + strategy_name = "megatron.core.dist_checkpointing.strategies.torch" + strategy = ModuleType(strategy_name) + calls = [] - def install_operator(self, module, operator): - self.installed[module] = operator + def replace(state_dict, flat_mapping, rename_mapping): + calls.append((state_dict, flat_mapping, rename_mapping)) + return state_dict - def record_installed_hook(self, module, hook): - self.hooks.append((module, hook)) + strategy._replace_sharded_keys_with_state_dict_keys = replace + monkeypatch.setitem(sys.modules, strategy_name, strategy) - integration = Integration() - _register_attention_backend(integration) + _install_torch_dist_object_compatibility() + installed = strategy._replace_sharded_keys_with_state_dict_keys + _install_torch_dist_object_compatibility() + payload = __import__("io").BytesIO() + torch.save([{"recipe": "checkpoint object"}], payload) - assert selected == [ - ( - BackendEnum.ROCM_AITER_FA, - "rl_engine.integrations.vllm_runtime.RlKernelAttentionBackend", - ) - ] - assert "attention" in integration.installed - assert integration.hooks == [ - ( - "attention", - "rl_engine.integrations.vllm_runtime.RlKernelAttentionBackend", - ) + assert strategy._replace_sharded_keys_with_state_dict_keys is installed + assert installed({"state": payload}, "flat", "rename") == { + "state": [{"recipe": "checkpoint object"}] + } + assert calls == [ + ({"state": [{"recipe": "checkpoint object"}]}, "flat", "rename") ] -def test_megatron_install_is_idempotent_in_one_actor(): - class Attention: - def forward(self, value): - return value - - class FFN: - def forward(self, value): - return value - - plan = IntegrationPlan.from_case_ids() - clear_active_integration("megatron") - try: - first = install_megatron_integration( - plan, - attention_classes=(Attention,), - ffn_classes=(FFN,), - ) - second = install_megatron_integration( - plan, - attention_classes=(Attention,), - ffn_classes=(FFN,), - ) - assert first is second - finally: - clear_active_integration("megatron") - - -def test_megatron_zigzag_positions_preserve_global_cp_ownership(): - rank_zero = _megatron_zigzag_layout(4, cp_rank=0, cp_world_size=2) - rank_one = _megatron_zigzag_layout(4, cp_rank=1, cp_world_size=2) - - assert rank_zero == ((0, 1, 6, 7), (0, 3), (0, 6), (0, 2, 4)) - assert rank_one == ((2, 3, 4, 5), (1, 2), (2, 4), (0, 2, 4)) - assert sorted(rank_zero[0] + rank_one[0]) == list(range(8)) - - -def test_packed_layout_recovers_local_offsets_from_global_cu_seqlens(): - packed = SimpleNamespace( - qkv_format="thd", - cu_seqlens_q=torch.tensor([0, 8, 16], dtype=torch.int32), - cu_seqlens_kv=torch.tensor([0, 8, 16], dtype=torch.int32), - ) - - local_offsets, global_lengths = _packed_local_sequence_layout( - packed, - cp_world_size=2, - local_query_tokens=8, - local_kv_tokens=8, - ) - - assert local_offsets == (0, 4, 8) - assert global_lengths == (8, 8) - - -def test_megatron_packed_attention_runs_each_sequence_in_thd_order(monkeypatch): - calls: list[dict[str, object]] = [] - - class Operator: - def bind_accelerator_runtime(self, tensor, *, process_group=None): - assert tensor is query - assert process_group == "cp-group" - - def __call__(self, q, k, v, **kwargs): - del k, v - calls.append(kwargs) - return SimpleNamespace( - out=q.clone(), - provenance={ - "actual_backend": "rlkernel.cuda.attention.fa4_ag_rs.v1", - "core_rows": [{"actual_backend": "rlkernel.cuda.attention.fa4.v1"}], - }, - ) - - operator = Operator() - - class Handle: - provenance = {} - - def get(self, tensor, *, topology): - assert tensor.shape == (8, 2, 4) - assert topology["context_parallel_size"] == 2 - return operator - - parallel_state = SimpleNamespace( - get_context_parallel_world_size=lambda: 2, - get_context_parallel_rank=lambda: 0, - get_tensor_model_parallel_world_size=lambda: 2, - get_tensor_model_parallel_rank=lambda: 0, - get_context_parallel_group=lambda: "cp-group", - ) +def test_framework_adapters_do_not_construct_registered_kernels_directly(): + source_path = ( + Path(__file__).parents[1] + / "rl_engine" + / "integrations" + / "framework_operators.py" + ) + tree = ast.parse(source_path.read_text(encoding="utf-8")) + forbidden = { + "AttentionAblationOp", + "DeterministicCPAttentionReferenceOp", + "Qwen3FFNOp", + "StrictCUDAAttentionRuntime", + "VocabParallelLogprobOp", + } + constructed = { + node.func.id + for node in ast.walk(tree) + if isinstance(node, ast.Call) and isinstance(node.func, ast.Name) + } + + assert constructed.isdisjoint(forbidden) + + +def test_semantic_handle_uses_operator_bridge_and_exposes_instance_provenance(): + handle = SemanticOperatorHandle( + target="training", + semantic_op="attention", + backend_id="rlkernel.attention.deterministic.v1", + ) + tensor = torch.zeros(1, 1, 1, 8) + + first = handle.get( + tensor, + topology={ + "world_size": 1, + "tensor_parallel_size": 1, + "context_parallel_size": 1, + }, + ) + second = handle.get( + tensor, + topology={ + "world_size": 1, + "tensor_parallel_size": 1, + "context_parallel_size": 1, + }, + ) + + assert first is second + assert first.backend_id == "rlkernel.attention.deterministic.v1" + assert handle.provenance is not None + assert handle.provenance["semantic_op"] == "attention" + assert handle.provenance["backend_id"] == first.backend_id + + +def test_plan_environment_is_shared_by_both_framework_installers(monkeypatch, tmp_path): + for variable in ( + "RL_KERNEL_ATTENTION_CASE", + "RL_KERNEL_FFN_CASE", + "RL_KERNEL_LOGP_CASE", + "RL_KERNEL_READBACK_DIR", + ): + monkeypatch.setenv(variable, "previous") + plan = IntegrationPlan.from_case_ids(attention="P/R", ffn="R/P", logp="R/R") + configure_integration_environment(plan, readback_dir=str(tmp_path)) + + assert integration_plan_from_environment() == plan + assert Path(os.environ["RL_KERNEL_READBACK_DIR"]) == tmp_path + + +def test_vllm_rlkernel_attention_overrides_selected_flash_attn_backend( + monkeypatch, +): + monkeypatch.delenv("VLLM_ATTENTION_BACKEND", raising=False) + plan = IntegrationPlan.from_case_ids(attention="P/R") + + configure_vllm_environment(plan) + + expected = "ROCM_AITER_FA" if torch.version.hip is not None else "FLASH_ATTN" + assert os.environ["VLLM_ATTENTION_BACKEND"] == expected + + +def test_vllm_rocm_attention_selects_aiter_metadata_backend(monkeypatch): + monkeypatch.delenv("VLLM_ATTENTION_BACKEND", raising=False) + monkeypatch.setattr(torch.version, "hip", "7.1") + plan = IntegrationPlan.from_case_ids(attention="P/R") + + configure_vllm_environment(plan) + + assert os.environ["VLLM_ATTENTION_BACKEND"] == "ROCM_AITER_FA" + + +def test_vllm_rocm_registration_wraps_aiter_backend(monkeypatch): + selected = [] + + class BackendEnum: + ROCM_AITER_FA = "ROCM_AITER_FA" + FLASH_ATTN = "FLASH_ATTN" + + class AiterImpl: + def __init__(self, *args, **kwargs): + del args, kwargs + + def forward(self, *args, **kwargs): + return args, kwargs + + class AiterBuilder: + pass + + class AiterBackend: + pass + + registry = ModuleType("vllm.v1.attention.backends.registry") + registry.AttentionBackendEnum = BackendEnum + registry.register_backend = lambda backend, path: selected.append((backend, path)) + aiter = ModuleType("vllm.v1.attention.backends.rocm_aiter_fa") + aiter.AiterFlashAttentionBackend = AiterBackend + aiter.AiterFlashAttentionImpl = AiterImpl + aiter.AiterFlashAttentionMetadataBuilder = AiterBuilder + monkeypatch.setitem(sys.modules, registry.__name__, registry) + monkeypatch.setitem(sys.modules, aiter.__name__, aiter) + monkeypatch.setattr(torch.version, "hip", "7.1") + + plan = IntegrationPlan.from_case_ids(attention="P/R") + + class Integration: + def __init__(self): + self.plan = plan + self.installed = {} + self.hooks = [] + + def install_operator(self, module, operator): + self.installed[module] = operator + + def record_installed_hook(self, module, hook): + self.hooks.append((module, hook)) + + integration = Integration() + _register_attention_backend(integration) + + assert selected == [ + ( + BackendEnum.ROCM_AITER_FA, + "rl_engine.integrations.vllm_runtime.RlKernelAttentionBackend", + ) + ] + assert "attention" in integration.installed + assert integration.hooks == [ + ( + "attention", + "rl_engine.integrations.vllm_runtime.RlKernelAttentionBackend", + ) + ] + + +def test_megatron_install_is_idempotent_in_one_actor(): + class Attention: + def forward(self, value): + return value + + class FFN: + def forward(self, value): + return value + + plan = IntegrationPlan.from_case_ids() + clear_active_integration("megatron") + try: + first = install_megatron_integration( + plan, + attention_classes=(Attention,), + ffn_classes=(FFN,), + ) + second = install_megatron_integration( + plan, + attention_classes=(Attention,), + ffn_classes=(FFN,), + ) + assert first is second + finally: + clear_active_integration("megatron") + + +def test_megatron_zigzag_positions_preserve_global_cp_ownership(): + rank_zero = _megatron_zigzag_layout(4, cp_rank=0, cp_world_size=2) + rank_one = _megatron_zigzag_layout(4, cp_rank=1, cp_world_size=2) + + assert rank_zero == ((0, 1, 6, 7), (0, 3), (0, 6), (0, 2, 4)) + assert rank_one == ((2, 3, 4, 5), (1, 2), (2, 4), (0, 2, 4)) + assert sorted(rank_zero[0] + rank_one[0]) == list(range(8)) + + +def test_packed_layout_recovers_local_offsets_from_global_cu_seqlens(): + packed = SimpleNamespace( + qkv_format="thd", + cu_seqlens_q=torch.tensor([0, 8, 16], dtype=torch.int32), + cu_seqlens_kv=torch.tensor([0, 8, 16], dtype=torch.int32), + ) + + local_offsets, global_lengths = _packed_local_sequence_layout( + packed, + cp_world_size=2, + local_query_tokens=8, + local_kv_tokens=8, + ) + + assert local_offsets == (0, 4, 8) + assert global_lengths == (8, 8) + + +def test_megatron_packed_attention_runs_each_sequence_in_thd_order(monkeypatch): + calls: list[dict[str, object]] = [] + + class Operator: + def bind_accelerator_runtime(self, tensor, *, process_group=None): + assert tensor is query + assert process_group == "cp-group" + + def __call__(self, q, k, v, **kwargs): + del k, v + calls.append(kwargs) + return SimpleNamespace( + out=q.clone(), + provenance={ + "actual_backend": "rlkernel.cuda.attention.fa4_ag_rs.v1", + "core_rows": [{"actual_backend": "rlkernel.cuda.attention.fa4.v1"}], + }, + ) + + operator = Operator() + + class Handle: + provenance = {} + + def get(self, tensor, *, topology): + assert tensor.shape == (8, 2, 4) + assert topology["context_parallel_size"] == 2 + return operator + + parallel_state = SimpleNamespace( + get_context_parallel_world_size=lambda: 2, + get_context_parallel_rank=lambda: 0, + get_tensor_model_parallel_world_size=lambda: 2, + get_tensor_model_parallel_rank=lambda: 0, + get_context_parallel_group=lambda: "cp-group", + ) monkeypatch.setattr(framework_operators, "_megatron_parallel_state", lambda: parallel_state) monkeypatch.setattr( framework_operators, "_require_attention_accelerator", lambda tensor: "cuda", ) - packed = SimpleNamespace( - qkv_format="thd", - cu_seqlens_q=torch.tensor([0, 8, 16], dtype=torch.int32), - cu_seqlens_kv=torch.tensor([0, 8, 16], dtype=torch.int32), - ) - query = torch.zeros(8, 2, 4, dtype=torch.bfloat16) - key = torch.zeros(8, 1, 4, dtype=torch.bfloat16) - - output = MegatronAttentionOperator(handle=Handle())( - SimpleNamespace(softmax_scale=0.5), - query, - key, - key, - None, - packed_seq_params=packed, - num_splits=1, - ) - - assert output.shape == (8, 8) - assert len(calls) == 1 - assert [call["contract"].sharding.global_sequence_length for call in calls] == [ - 8, - ] - assert calls[0]["query_position_ids"].tolist() == [ - [0, 1, 6, 7], - [0, 1, 6, 7], - ] + packed = SimpleNamespace( + qkv_format="thd", + cu_seqlens_q=torch.tensor([0, 8, 16], dtype=torch.int32), + cu_seqlens_kv=torch.tensor([0, 8, 16], dtype=torch.int32), + ) + query = torch.zeros(8, 2, 4, dtype=torch.bfloat16) + key = torch.zeros(8, 1, 4, dtype=torch.bfloat16) + + output = MegatronAttentionOperator(handle=Handle())( + SimpleNamespace(softmax_scale=0.5), + query, + key, + key, + None, + packed_seq_params=packed, + num_splits=1, + ) + + assert output.shape == (8, 8) + assert len(calls) == 1 + assert [call["contract"].sharding.global_sequence_length for call in calls] == [ + 8, + ] + assert calls[0]["query_position_ids"].tolist() == [ + [0, 1, 6, 7], + [0, 1, 6, 7], + ] def test_megatron_attention_binds_rocm_core_and_schedule(monkeypatch): - calls = [] - - class Operator: - def bind_accelerator_runtime(self, tensor, *, process_group=None): - calls.append((tensor, process_group)) - - def __call__(self, q, k, v, **kwargs): - del k, v - calls.append(kwargs) - return SimpleNamespace( - out=q.clone(), - provenance={ - "actual_backend": "rlkernel.rocm.attention.aiter_ck_ag_rs.v1", - "fallback": False, - }, - ) - - operator = Operator() - - class Handle: - provenance = {} - - def get(self, tensor, *, topology): - assert topology["context_parallel_size"] == 1 - return operator - - parallel_state = SimpleNamespace( - get_context_parallel_world_size=lambda: 1, - get_context_parallel_rank=lambda: 0, - get_tensor_model_parallel_world_size=lambda: 2, - get_tensor_model_parallel_rank=lambda: 0, - ) - monkeypatch.setattr(framework_operators, "_megatron_parallel_state", lambda: parallel_state) - monkeypatch.setattr( - framework_operators, - "_require_attention_accelerator", - lambda tensor: "rocm", - ) - query = torch.zeros(4, 1, 2, 8, dtype=torch.bfloat16) - key = torch.zeros(4, 1, 1, 8, dtype=torch.bfloat16) - adapter = MegatronAttentionOperator(handle=Handle()) - - output = adapter(SimpleNamespace(softmax_scale=0.25), query, key, key, None) - - assert output.shape == (4, 1, 16) - assert calls[0] == (query, None) - config = calls[1]["config"] - assert config.strict_core_id == "rlkernel.attention.rocm.aiter_ck_dense_mha.v1" - assert config.strict_schedule == "single_batch_aiter_ck_dense_mha_no_splitkv" - assert calls[1]["communication_backend"] == "none" - assert adapter.provenance["execution"]["runtime_platform"] == "rocm" - - -def test_vllm_attention_routes_paged_cache_to_rocm_runtime(monkeypatch): - runtime_calls = [] - - class Runtime: - def forward_paged_with_lse(self, q, k, v, **kwargs): - runtime_calls.append((q, k, v, kwargs)) - return SimpleNamespace( - out=q.clone(), - lse=torch.zeros(q.shape[:-1], dtype=torch.float32), - provenance={ - "actual_backend": "rlkernel.rocm.attention.aiter_ck_ag_rs.v1", - "fallback": False, - }, - ) - - class Operator: - def bind_accelerator_runtime(self, tensor, *, process_group=None): - assert process_group is None - return Runtime() - - class Handle: - provenance = {} - - def get(self, tensor, *, topology): - assert topology["context_parallel_size"] == 1 - return Operator() - - monkeypatch.setattr( - framework_operators, - "_require_attention_accelerator", - lambda tensor: "rocm", - ) - query = torch.zeros(1, 2, 8, dtype=torch.bfloat16) - kv_cache = torch.zeros(2, 1, 4, 16, dtype=torch.bfloat16) - metadata = SimpleNamespace( - block_table=torch.tensor([[0]], dtype=torch.int32), - query_start_loc=torch.tensor([0, 1], dtype=torch.int32), - seq_lens=torch.tensor([1], dtype=torch.int32), - num_actual_tokens=1, - max_seq_len=1, - ) - impl = SimpleNamespace(head_size=8, num_heads=2, num_kv_heads=1, scale=8**-0.5) - adapter = VllmAttentionOperator(handle=Handle()) - - output = adapter(impl, object(), query, query, query, kv_cache, metadata) - - assert output.shape == (1, 16) - assert len(runtime_calls) == 1 - assert runtime_calls[0][0].shape == (1, 2, 1, 8) - assert adapter.provenance["execution"]["runtime_platform"] == "rocm" - assert ( - adapter.provenance["execution"]["materialization"] - == "logical_paged_kv_to_aiter_ck_dense" - ) - - + calls = [] + + class Operator: + def bind_accelerator_runtime(self, tensor, *, process_group=None): + calls.append((tensor, process_group)) + + def __call__(self, q, k, v, **kwargs): + del k, v + calls.append(kwargs) + return SimpleNamespace( + out=q.clone(), + provenance={ + "actual_backend": "rlkernel.rocm.attention.aiter_ck_ag_rs.v1", + "fallback": False, + }, + ) + + operator = Operator() + + class Handle: + provenance = {} + + def get(self, tensor, *, topology): + assert topology["context_parallel_size"] == 1 + return operator + + parallel_state = SimpleNamespace( + get_context_parallel_world_size=lambda: 1, + get_context_parallel_rank=lambda: 0, + get_tensor_model_parallel_world_size=lambda: 2, + get_tensor_model_parallel_rank=lambda: 0, + ) + monkeypatch.setattr(framework_operators, "_megatron_parallel_state", lambda: parallel_state) + monkeypatch.setattr( + framework_operators, + "_require_attention_accelerator", + lambda tensor: "rocm", + ) + query = torch.zeros(4, 1, 2, 8, dtype=torch.bfloat16) + key = torch.zeros(4, 1, 1, 8, dtype=torch.bfloat16) + adapter = MegatronAttentionOperator(handle=Handle()) + + output = adapter(SimpleNamespace(softmax_scale=0.25), query, key, key, None) + + assert output.shape == (4, 1, 16) + assert calls[0] == (query, None) + config = calls[1]["config"] + assert config.strict_core_id == "rlkernel.attention.rocm.aiter_ck_dense_mha.v1" + assert config.strict_schedule == "single_batch_aiter_ck_dense_mha_no_splitkv" + assert calls[1]["communication_backend"] == "none" + assert adapter.provenance["execution"]["runtime_platform"] == "rocm" + + +def test_vllm_attention_routes_paged_cache_to_rocm_runtime(monkeypatch): + runtime_calls = [] + + class Runtime: + def forward_paged_with_lse(self, q, k, v, **kwargs): + runtime_calls.append((q, k, v, kwargs)) + return SimpleNamespace( + out=q.clone(), + lse=torch.zeros(q.shape[:-1], dtype=torch.float32), + provenance={ + "actual_backend": "rlkernel.rocm.attention.aiter_ck_ag_rs.v1", + "fallback": False, + }, + ) + + class Operator: + def bind_accelerator_runtime(self, tensor, *, process_group=None): + assert process_group is None + return Runtime() + + class Handle: + provenance = {} + + def get(self, tensor, *, topology): + assert topology["context_parallel_size"] == 1 + return Operator() + + monkeypatch.setattr( + framework_operators, + "_require_attention_accelerator", + lambda tensor: "rocm", + ) + query = torch.zeros(1, 2, 8, dtype=torch.bfloat16) + kv_cache = torch.zeros(2, 1, 4, 16, dtype=torch.bfloat16) + metadata = SimpleNamespace( + block_table=torch.tensor([[0]], dtype=torch.int32), + query_start_loc=torch.tensor([0, 1], dtype=torch.int32), + seq_lens=torch.tensor([1], dtype=torch.int32), + num_actual_tokens=1, + max_seq_len=1, + ) + impl = SimpleNamespace(head_size=8, num_heads=2, num_kv_heads=1, scale=8**-0.5) + adapter = VllmAttentionOperator(handle=Handle()) + + output = adapter(impl, object(), query, query, query, kv_cache, metadata) + + assert output.shape == (1, 16) + assert len(runtime_calls) == 1 + assert runtime_calls[0][0].shape == (1, 2, 1, 8) + assert adapter.provenance["execution"]["runtime_platform"] == "rocm" + assert ( + adapter.provenance["execution"]["materialization"] + == "logical_paged_kv_to_aiter_ck_dense" + ) + + def test_vllm_current_flash_attention_kv_cache_layout_is_materialized(): cache = torch.arange(2 * 3 * 4 * 10).reshape(2, 3, 4, 10) - - key, value = _vllm_kv_cache_views(cache, head_size=5) - - assert key.shape == (2, 4, 3, 5) - assert value.shape == (2, 4, 3, 5) - assert torch.equal(key, cache.transpose(1, 2)[..., :5]) + + key, value = _vllm_kv_cache_views(cache, head_size=5) + + assert key.shape == (2, 4, 3, 5) + assert value.shape == (2, 4, 3, 5) + assert torch.equal(key, cache.transpose(1, 2)[..., :5]) assert torch.equal(value, cache.transpose(1, 2)[..., 5:]) -def test_vllm_rocm_kv_cache_pair_axis_is_materialized(): - cache = torch.arange(3 * 2 * 4 * 1 * 5).reshape(3, 2, 4, 1, 5) +def test_vllm_rocm_native_kv_cache_with_two_heads_is_not_treated_as_pair_axis(): + cache = torch.arange(3 * 2 * 4 * 10).reshape(3, 2, 4, 10) key, value = _vllm_kv_cache_views( cache, head_size=5, - num_kv_heads=1, + num_kv_heads=2, platform="rocm", ) - assert key.shape == (3, 4, 1, 5) - assert value.shape == (3, 4, 1, 5) - assert torch.equal(key, cache[:, 0]) - assert torch.equal(value, cache[:, 1]) - - -def test_vllm_rocm_lhbnc_kv_cache_is_normalized_to_block_major(): - cache = torch.arange(2 * 1 * 3 * 4 * 5).reshape(2, 1, 3, 4, 5) - - key, value = _vllm_kv_cache_views(cache, head_size=5, num_kv_heads=1) - - assert key.shape == (3, 4, 1, 5) - assert value.shape == (3, 4, 1, 5) - assert torch.equal(key, cache[0].permute(1, 2, 0, 3)) - assert torch.equal(value, cache[1].permute(1, 2, 0, 3)) + assert key.shape == (3, 4, 2, 5) + assert value.shape == (3, 4, 2, 5) + assert torch.equal(key, cache.transpose(1, 2)[..., :5]) + assert torch.equal(value, cache.transpose(1, 2)[..., 5:]) +def test_vllm_rocm_kv_cache_pair_axis_is_materialized(): + cache = torch.arange(3 * 2 * 4 * 1 * 5).reshape(3, 2, 4, 1, 5) + + key, value = _vllm_kv_cache_views( + cache, + head_size=5, + num_kv_heads=1, + platform="rocm", + ) + + assert key.shape == (3, 4, 1, 5) + assert value.shape == (3, 4, 1, 5) + assert torch.equal(key, cache[:, 0]) + assert torch.equal(value, cache[:, 1]) + + +def test_vllm_rocm_lhbnc_kv_cache_is_normalized_to_block_major(): + cache = torch.arange(2 * 1 * 3 * 4 * 5).reshape(2, 1, 3, 4, 5) + + key, value = _vllm_kv_cache_views(cache, head_size=5, num_kv_heads=1) + + assert key.shape == (3, 4, 1, 5) + assert value.shape == (3, 4, 1, 5) + assert torch.equal(key, cache[0].permute(1, 2, 0, 3)) + assert torch.equal(value, cache[1].permute(1, 2, 0, 3)) + + def test_vllm_rocm_flattened_kv_cache_is_unpacked(): - cache = torch.arange(3 * 2 * 4 * 10).reshape(3, 2, 4, 10) + cache = torch.arange(3 * 2 * 4 * 15).reshape(3, 2, 4, 15) key, value = _vllm_kv_cache_views( cache, head_size=5, - num_kv_heads=2, + num_kv_heads=3, platform="rocm", ) - assert key.shape == (3, 4, 2, 5) - assert value.shape == (3, 4, 2, 5) - assert torch.equal(key.flatten(2), cache[:, 0]) - assert torch.equal(value.flatten(2), cache[:, 1]) - - -def test_megatron_strict_attention_projections_install_without_debug_environment( - monkeypatch, -): - monkeypatch.delenv("RL_KERNEL_MODEL_DEBUG_DIR", raising=False) - - class ColumnLinear: - def __init__(self): - self.allreduce_dgrad = True - - def _forward_impl(self, input, weight, *args, **kwargs): - del args, kwargs - return input.new_full((*input.shape[:-1], weight.shape[0]), -1) - - class RowLinear: - def _forward_impl(self, input, weight, *args, **kwargs): - del args, kwargs - return input.new_full((*input.shape[:-1], weight.shape[0]), -1) - - class SelfAttention: - def __init__(self): - self.linear_qkv = ColumnLinear() - self.linear_proj = RowLinear() - - _patch_strict_attention_projections( - self_attention_cls=SelfAttention, - column_linear_cls=ColumnLinear, - row_linear_cls=RowLinear, - det_gemm=lambda lhs, rhs: lhs @ rhs, - ) - attention = SelfAttention() - value = torch.tensor([[1.0, 2.0]]) - weight = torch.tensor([[1.0, 0.0], [0.0, 1.0], [1.0, 1.0]]) - - qkv = attention.linear_qkv._forward_impl(value, weight, bias=None) - projection = attention.linear_proj._forward_impl(value, weight, bias=None) - native = ColumnLinear()._forward_impl(value, weight, bias=None) - - assert torch.equal(qkv, torch.tensor([[1.0, 2.0, 3.0]])) - assert torch.equal(projection, qkv) - assert torch.equal(native, torch.full((1, 3), -1.0)) - assert attention.linear_qkv.allreduce_dgrad is False - - -def test_vllm_qwen3_strict_model_installs_without_debug_environment(monkeypatch): - monkeypatch.delenv("RL_KERNEL_MODEL_DEBUG_DIR", raising=False) - - class RMSNorm: - def __init__(self): - self.variance_size_override = None - self.has_weight = True - self.hidden_size = 2 - self.weight = torch.tensor([1.5, 0.5]) - self.variance_epsilon = 1e-6 - - def forward_cuda(self, x, residual=None): - del residual - return x.new_full(x.shape, -1) - - def forward_native(self, x, residual=None): - del residual - return x.new_full(x.shape, -2) - - class LinearLayer: - def __init__(self): - self.weight = torch.tensor([[1.0, 0.0], [0.0, 1.0], [1.0, 1.0]]) - - class LinearMethod: - def apply(self, layer, x, bias=None): - del bias - return x.new_full((*x.shape[:-1], layer.weight.shape[0]), -1) - - class Attention: - def __init__(self): - self.qkv_proj = LinearLayer() - self.o_proj = LinearLayer() - - _patch_qwen3_strict_model( - rms_norm_cls=RMSNorm, - linear_method_cls=LinearMethod, - attention_cls=Attention, - det_gemm=lambda lhs, rhs: lhs @ rhs, - ) - attention = Attention() - method = LinearMethod() - value = torch.tensor([[1.0, 2.0]]) - norm = RMSNorm() - - assert torch.equal( - method.apply(attention.qkv_proj, value), - torch.tensor([[1.0, 2.0, 3.0]]), - ) - assert torch.equal( - method.apply(LinearLayer(), value), - torch.full((1, 3), -1.0), - ) - assert torch.equal( - norm.forward_cuda(value), - torch.nn.functional.rms_norm(value, (2,), norm.weight, 1e-6), - ) - - -def test_vllm_logp_replaces_every_duplicate_sampled_token_column(): - logprobs_type = namedtuple( - "LogprobsTensors", - ("logprob_token_ids", "logprobs", "selected_token_ranks"), - ) - - @dataclass(frozen=True) - class SamplerResult: - sampled_token_ids: torch.Tensor - logprobs_tensors: object - - operator = VllmLogpOperator(lambda *_args, **_kwargs: None) - result = SamplerResult( - sampled_token_ids=torch.tensor([[7], [8]]), - logprobs_tensors=logprobs_type( - logprob_token_ids=torch.tensor([[7, 7], [8, 9]]), - logprobs=torch.tensor([[-0.1, -0.1], [-0.2, -0.3]]), - selected_token_ranks=torch.tensor([1, 2]), - ), - ) - - updated = operator._replace_sampled_value( - result, - token_ids=torch.tensor([7, 8]), - selected=torch.tensor([-1.25, -2.5]), - provenance={}, - ) - - assert torch.equal( - updated.logprobs_tensors.logprobs, - torch.tensor([[-1.25, -1.25], [-2.5, -0.3]]), - ) - - -def _cp1_contract(tokens: int) -> AttentionContract: - return AttentionContract( - role=AttentionRole.TRAIN, - mode=AttentionMode.PREFILL, - dtype=AttentionDType.BF16, - batch_size=1, - query_sequence_length=tokens, - head_dim=4, - causal=True, - causal_offsets=(0,), - sharding=ShardingSpec( - tp_rank=0, - tp_world_size=1, - cp_rank=0, - cp_world_size=1, - global_q_heads=2, - global_kv_heads=1, - local_q_head_start=0, - local_q_heads=2, - local_kv_head_start=0, - local_kv_heads=1, - global_sequence_length=tokens, - local_sequence_length=tokens, - global_block_indices=(0,), - global_block_token_starts=(0,), - local_block_offsets=(0, tokens), - ), - reduction=ReductionSpec(), - ) - - -def test_strict_cuda_runtime_pins_training_to_single_query_prefixes(monkeypatch): - calls: list[tuple[int, int, bool]] = [] - - class Core: - core_id = STRICT_ATTENTION_PRODUCTION_CORE_ID - strict_schedule = STRICT_ATTENTION_FA4_SCHEDULE_ID - - def forward_bshd_with_lse(self, q, k, v, *, causal, **kwargs): - del v, kwargs - calls.append((q.size(1), k.size(1), causal)) - return SimpleNamespace( - out=q.clone(), - lse=torch.zeros((q.size(0), q.size(2), q.size(1)), dtype=torch.float32), - provenance={"actual_backend": "fake.cuda.fa4"}, - ) - - monkeypatch.setattr( - StrictCUDAAttentionRuntime, "_require_nvidia_cuda", lambda self, tensor: None - ) - runtime = StrictCUDAAttentionRuntime(core=Core(), communication=object()) - q = torch.zeros(1, 2, 4, 4, dtype=torch.bfloat16) - k = torch.zeros(1, 1, 4, 4, dtype=torch.bfloat16) - positions = torch.arange(4).unsqueeze(0) - - result = runtime.forward_with_lse( - q, - k, - k, - contract=_cp1_contract(4), - causal=True, - scale=0.5, - cp_world_size=1, - query_position_ids=positions, - key_position_ids=positions, - ) - - assert calls == [(4, 4, True)] + assert key.shape == (3, 4, 3, 5) + assert value.shape == (3, 4, 3, 5) + assert torch.equal(key.flatten(2), cache[:, 0]) + assert torch.equal(value.flatten(2), cache[:, 1]) + + +def test_megatron_strict_attention_projections_install_without_debug_environment( + monkeypatch, +): + monkeypatch.delenv("RL_KERNEL_MODEL_DEBUG_DIR", raising=False) + + class ColumnLinear: + def __init__(self): + self.allreduce_dgrad = True + + def _forward_impl(self, input, weight, *args, **kwargs): + del args, kwargs + return input.new_full((*input.shape[:-1], weight.shape[0]), -1) + + class RowLinear: + def _forward_impl(self, input, weight, *args, **kwargs): + del args, kwargs + return input.new_full((*input.shape[:-1], weight.shape[0]), -1) + + class SelfAttention: + def __init__(self): + self.linear_qkv = ColumnLinear() + self.linear_proj = RowLinear() + + _patch_strict_attention_projections( + self_attention_cls=SelfAttention, + column_linear_cls=ColumnLinear, + row_linear_cls=RowLinear, + det_gemm=lambda lhs, rhs: lhs @ rhs, + ) + attention = SelfAttention() + value = torch.tensor([[1.0, 2.0]]) + weight = torch.tensor([[1.0, 0.0], [0.0, 1.0], [1.0, 1.0]]) + + qkv = attention.linear_qkv._forward_impl(value, weight, bias=None) + projection = attention.linear_proj._forward_impl(value, weight, bias=None) + native = ColumnLinear()._forward_impl(value, weight, bias=None) + + assert torch.equal(qkv, torch.tensor([[1.0, 2.0, 3.0]])) + assert torch.equal(projection, qkv) + assert torch.equal(native, torch.full((1, 3), -1.0)) + assert attention.linear_qkv.allreduce_dgrad is False + + +def test_megatron_te_attention_projection_uses_injected_strict_tp_reduce(): + calls: list[torch.Tensor] = [] + + class ColumnLinear: + def __init__(self): + self.layer_norm_weight = torch.ones(2) + self.weight = torch.tensor([[1.0, 0.0], [0.0, 1.0]]) + + def _forward_impl(self, input, weight, *args, **kwargs): + del args, kwargs + return input @ weight.t() + + class RowLinear: + def __init__(self): + self.weight = torch.tensor([[1.0, 0.0], [0.0, 1.0]]) + + def _forward_impl(self, input, weight, *args, **kwargs): + del args, kwargs + return input @ weight.t() + + class SelfAttention: + def __init__(self): + self.linear_qkv = ColumnLinear() + self.linear_proj = RowLinear() + + def reduce_from_tp(value: torch.Tensor) -> torch.Tensor: + calls.append(value) + return value * 4 + + _patch_strict_attention_projections( + self_attention_cls=SelfAttention, + column_linear_cls=ColumnLinear, + row_linear_cls=RowLinear, + det_gemm=lambda lhs, rhs: lhs @ rhs, + copy_to_tp=lambda value: value, + reduce_from_tp=reduce_from_tp, + ) + attention = SelfAttention() + value = torch.tensor([[1.0, 2.0]]) + + output, bias = attention.linear_proj.forward(value) + + assert bias is None + assert len(calls) == 1 + assert torch.equal(calls[0], value) + assert torch.equal(output, value * 4) + + +def test_megatron_deterministic_tp_reduce_keeps_identity_backward(monkeypatch): + class Collective: + def all_reduce(self, value): + return value * 4 + + monkeypatch.setattr( + "rl_engine.distributed.collectives.collective_for_group", + lambda group, min_size_bytes: Collective(), + ) + value = torch.tensor([1.0, 2.0], requires_grad=True) + + output = _deterministic_reduce_from_tensor_model_parallel_region(value, object()) + output.sum().backward() + + assert torch.equal(output, value.detach() * 4) + assert torch.equal(value.grad, torch.ones_like(value)) + + +def test_vllm_qwen3_strict_model_installs_without_debug_environment(monkeypatch): + monkeypatch.delenv("RL_KERNEL_MODEL_DEBUG_DIR", raising=False) + + class RMSNorm: + def __init__(self): + self.variance_size_override = None + self.has_weight = True + self.hidden_size = 2 + self.weight = torch.tensor([1.5, 0.5]) + self.variance_epsilon = 1e-6 + + def forward_cuda(self, x, residual=None): + del residual + return x.new_full(x.shape, -1) + + def forward_native(self, x, residual=None): + del residual + return x.new_full(x.shape, -2) + + class LinearLayer: + def __init__(self): + self.weight = torch.tensor([[1.0, 0.0], [0.0, 1.0], [1.0, 1.0]]) + + class LinearMethod: + def apply(self, layer, x, bias=None): + del bias + return x.new_full((*x.shape[:-1], layer.weight.shape[0]), -1) + + class Attention: + def __init__(self): + self.qkv_proj = LinearLayer() + self.o_proj = LinearLayer() + + _patch_qwen3_strict_model( + rms_norm_cls=RMSNorm, + linear_method_cls=LinearMethod, + attention_cls=Attention, + det_gemm=lambda lhs, rhs: lhs @ rhs, + ) + attention = Attention() + method = LinearMethod() + value = torch.tensor([[1.0, 2.0]]) + norm = RMSNorm() + + assert torch.equal( + method.apply(attention.qkv_proj, value), + torch.tensor([[1.0, 2.0, 3.0]]), + ) + assert torch.equal( + method.apply(LinearLayer(), value), + torch.full((1, 3), -1.0), + ) + assert torch.equal( + norm.forward_cuda(value), + torch.nn.functional.rms_norm(value, (2,), norm.weight, 1e-6), + ) + + +def test_vllm_logp_replaces_every_duplicate_sampled_token_column(): + logprobs_type = namedtuple( + "LogprobsTensors", + ("logprob_token_ids", "logprobs", "selected_token_ranks"), + ) + + @dataclass(frozen=True) + class SamplerResult: + sampled_token_ids: torch.Tensor + logprobs_tensors: object + + operator = VllmLogpOperator(lambda *_args, **_kwargs: None) + result = SamplerResult( + sampled_token_ids=torch.tensor([[7], [8]]), + logprobs_tensors=logprobs_type( + logprob_token_ids=torch.tensor([[7, 7], [8, 9]]), + logprobs=torch.tensor([[-0.1, -0.1], [-0.2, -0.3]]), + selected_token_ranks=torch.tensor([1, 2]), + ), + ) + + updated = operator._replace_sampled_value( + result, + token_ids=torch.tensor([7, 8]), + selected=torch.tensor([-1.25, -2.5]), + provenance={}, + ) + + assert torch.equal( + updated.logprobs_tensors.logprobs, + torch.tensor([[-1.25, -1.25], [-2.5, -0.3]]), + ) + + +def _cp1_contract(tokens: int) -> AttentionContract: + return AttentionContract( + role=AttentionRole.TRAIN, + mode=AttentionMode.PREFILL, + dtype=AttentionDType.BF16, + batch_size=1, + query_sequence_length=tokens, + head_dim=4, + causal=True, + causal_offsets=(0,), + sharding=ShardingSpec( + tp_rank=0, + tp_world_size=1, + cp_rank=0, + cp_world_size=1, + global_q_heads=2, + global_kv_heads=1, + local_q_head_start=0, + local_q_heads=2, + local_kv_head_start=0, + local_kv_heads=1, + global_sequence_length=tokens, + local_sequence_length=tokens, + global_block_indices=(0,), + global_block_token_starts=(0,), + local_block_offsets=(0, tokens), + ), + reduction=ReductionSpec(), + ) + + +def test_strict_cuda_runtime_pins_training_to_single_query_prefixes(monkeypatch): + calls: list[tuple[int, int, bool]] = [] + + class Core: + core_id = STRICT_ATTENTION_PRODUCTION_CORE_ID + strict_schedule = STRICT_ATTENTION_FA4_SCHEDULE_ID + + def forward_bshd_with_lse(self, q, k, v, *, causal, **kwargs): + del v, kwargs + calls.append((q.size(1), k.size(1), causal)) + return SimpleNamespace( + out=q.clone(), + lse=torch.zeros((q.size(0), q.size(2), q.size(1)), dtype=torch.float32), + provenance={"actual_backend": "fake.cuda.fa4"}, + ) + + monkeypatch.setattr( + StrictCUDAAttentionRuntime, "_require_nvidia_cuda", lambda self, tensor: None + ) + runtime = StrictCUDAAttentionRuntime(core=Core(), communication=object()) + q = torch.zeros(1, 2, 4, 4, dtype=torch.bfloat16) + k = torch.zeros(1, 1, 4, 4, dtype=torch.bfloat16) + positions = torch.arange(4).unsqueeze(0) + + result = runtime.forward_with_lse( + q, + k, + k, + contract=_cp1_contract(4), + causal=True, + scale=0.5, + cp_world_size=1, + query_position_ids=positions, + key_position_ids=positions, + ) + + assert calls == [(4, 4, True)] assert result.provenance["query_schedule"] == "full_sequence_causal_single_launch" - - -class _ReadbackOperator: - backend_id = "rlkernel.attention.test" - - def __init__(self, provenance): - self.provenance = provenance - - def __call__(self, value): - return value - - -@pytest.mark.parametrize( - ("provenance", "match"), - [ - ({"runtime_platform": "cpu"}, "non-cuda"), - ({"runtime_platform": "cuda", "actual_backend": "triton.attention"}, "triton"), - ({"runtime_platform": "cuda", "triton_used": True}, "triton"), - ], -) -def test_strict_readback_rejects_non_cuda_and_triton(provenance, match): - plan = IntegrationPlan.from_case_ids(attention="R/R") - integration = FrameworkOperatorIntegration( - framework="megatron", - target="training", - plan=plan, - rl_kernel_operators={"attention": _ReadbackOperator(provenance)}, - ) - integration.record_installed_hook("attention", "test.attention") - integration.execute("attention", lambda value: value, "x") - - with pytest.raises(RuntimeError, match=match): - integration.assert_strict_ready() - - -def test_strict_readback_accepts_cuda_without_triton(): + + +class _ReadbackOperator: + backend_id = "rlkernel.attention.test" + + def __init__(self, provenance): + self.provenance = provenance + + def __call__(self, value): + return value + + +@pytest.mark.parametrize( + ("provenance", "match"), + [ + ({"runtime_platform": "cpu"}, "non-cuda"), + ({"runtime_platform": "cuda", "actual_backend": "triton.attention"}, "triton"), + ({"runtime_platform": "cuda", "triton_used": True}, "triton"), + ], +) +def test_strict_readback_rejects_non_cuda_and_triton(provenance, match): + plan = IntegrationPlan.from_case_ids(attention="R/R") + integration = FrameworkOperatorIntegration( + framework="megatron", + target="training", + plan=plan, + rl_kernel_operators={"attention": _ReadbackOperator(provenance)}, + ) + integration.record_installed_hook("attention", "test.attention") + integration.execute("attention", lambda value: value, "x") + + with pytest.raises(RuntimeError, match=match): + integration.assert_strict_ready() + + +def test_strict_readback_accepts_cuda_without_triton(): + plan = IntegrationPlan.from_case_ids(attention="R/R") + integration = FrameworkOperatorIntegration( + framework="megatron", + target="training", + plan=plan, + rl_kernel_operators={ + "attention": _ReadbackOperator( + { + "runtime_platform": "cuda", + "actual_backend": "rlkernel.cuda.fa4", + "triton_used": False, + } + ) + }, + ) + integration.record_installed_hook("attention", "test.attention") + integration.execute("attention", lambda value: value, "x") + + integration.assert_strict_ready() + + +def test_strict_readback_accepts_rocm_without_triton(): plan = IntegrationPlan.from_case_ids(attention="R/R") integration = FrameworkOperatorIntegration( framework="megatron", target="training", plan=plan, rl_kernel_operators={ - "attention": _ReadbackOperator( - { - "runtime_platform": "cuda", - "actual_backend": "rlkernel.cuda.fa4", - "triton_used": False, - } - ) - }, - ) - integration.record_installed_hook("attention", "test.attention") - integration.execute("attention", lambda value: value, "x") - + "attention": _ReadbackOperator( + { + "runtime_platform": "rocm", + "actual_backend": "rlkernel.rocm.attention.aiter_ck_ag_rs.v1", + "triton_used": False, + } + ) + }, + ) + integration.record_installed_hook("attention", "test.attention") + integration.execute("attention", lambda value: value, "x") + integration.assert_strict_ready() -def test_strict_readback_accepts_rocm_without_triton(): - plan = IntegrationPlan.from_case_ids(attention="R/R") +def test_production_readback_infers_platform_from_real_execution_tensors(): + plan = IntegrationPlan.from_case_ids(attention="P/P") integration = FrameworkOperatorIntegration( framework="megatron", target="training", plan=plan, - rl_kernel_operators={ - "attention": _ReadbackOperator( - { - "runtime_platform": "rocm", - "actual_backend": "rlkernel.rocm.attention.aiter_ck_ag_rs.v1", - "triton_used": False, - } - ) - }, - ) - integration.record_installed_hook("attention", "test.attention") - integration.execute("attention", lambda value: value, "x") - - integration.assert_strict_ready() + rl_kernel_operators={}, + ) + value = torch.zeros(2) + + integration.execute("attention", lambda tensor: tensor + 1, value) + + readback = integration.readback()["operators"]["attention"] + assert readback["implementation"] == "production" + assert readback["provenance"]["runtime_platform"] == "cpu" + + +def test_production_readback_uses_structural_result_provenance(): + plan = IntegrationPlan.from_case_ids(logp="P/P") + integration = FrameworkOperatorIntegration( + framework="megatron", + target="training", + plan=plan, + rl_kernel_operators={}, + ) + request = SimpleNamespace(logits=torch.zeros(2, 4), target_ids=torch.zeros(2)) + + def native(actual_request): + return SimpleNamespace( + logp=actual_request.logits[:, :1], + provenance={"actual_backend": "production.logp.test"}, + ) + + integration.execute("logp", native, request) + + provenance = integration.readback()["operators"]["logp"]["provenance"] + assert provenance["actual_backend"] == "production.logp.test" + assert provenance["runtime_platform"] == "cpu" diff --git a/tests/test_rocm_packed_ffn.py b/tests/test_rocm_packed_ffn.py new file mode 100644 index 00000000..70dcf47e --- /dev/null +++ b/tests/test_rocm_packed_ffn.py @@ -0,0 +1,139 @@ +# SPDX-License-Identifier: Apache-2.0 + +from __future__ import annotations + +import pytest +import torch + +import rl_engine.kernels.ops.pytorch.ffn.ffn as ffn_module +from rl_engine.integrations.framework_operators import VllmFFNOperator +from rl_engine.kernels.ops.matmul.det_gemm import ( + DetGemmOp, + det_gemm_linear_weight_gradient, +) +from rl_engine.kernels.registry import _default_semantic_descriptors + + +pytestmark = pytest.mark.skipif( + getattr(torch.version, "hip", None) is None, + reason="ROCm packed FFN tests require a ROCm PyTorch build", +) + + +class _FakeDist: + @staticmethod + def get_world_size(*, group): + del group + return 4 + + +class _FakeCollective: + _handle = 17 + backend_id = "rocm_ipc_fixed_tree" + + def __init__(self) -> None: + self.reduced = None + + def all_reduce(self, value, *, out): + assert out is value + self.reduced = value + return out + + +def test_ffn_descriptor_advertises_rocm_support(): + descriptor = next( + item + for item in _default_semantic_descriptors() + if item.backend_id == "rlkernel.ffn.qwen3.deterministic.v1" + ) + assert descriptor.supported_devices == frozenset({"cuda", "rocm"}) + + +def test_rocm_prepare_binds_eager_fixed_tree_without_cuda_graph_staging(monkeypatch): + collective = _FakeCollective() + monkeypatch.setattr(ffn_module, "_require_parallel_group", lambda *_args: _FakeDist()) + monkeypatch.setattr(ffn_module, "_collective_for_group", lambda *_args, **_kwargs: collective) + monkeypatch.delenv("RL_KERNEL_VLLM_CUDAGRAPH_MAX_CAPTURE_SIZE", raising=False) + + operator = ffn_module.Qwen3FFNOp() + handle, world_size = operator.prepare_packed_inference( + torch.empty((16, 8), dtype=torch.bfloat16), + torch.empty((8, 8), dtype=torch.bfloat16), + tp_group=object(), + ) + + assert (handle, world_size) == (17, 4) + assert operator.packed_inference_backend_id(handle) == "rocm_ipc_fixed_tree" + + +def test_rocm_packed_forward_reduces_in_place_with_bound_collective(monkeypatch): + collective = _FakeCollective() + expected = torch.arange(16, dtype=torch.bfloat16).reshape(2, 8) + monkeypatch.setattr( + ffn_module, + "_qwen3_ffn_packed_inference", + lambda *_args: expected, + ) + + actual = ffn_module.qwen3_ffn_packed_inference( + torch.empty((2, 8), dtype=torch.bfloat16), + torch.empty((16, 8), dtype=torch.bfloat16), + torch.empty((8, 8), dtype=torch.bfloat16), + collective_handle=17, + tp_world_size=4, + collective=collective, + ) + + assert actual is expected + assert collective.reduced is expected + + +def test_rocm_weight_gradient_uses_parameter_layout_and_is_bitwise_stable(): + inputs = torch.randn((16, 8), device="cuda", dtype=torch.bfloat16) + grad_output = torch.randn((16, 12), device="cuda", dtype=torch.bfloat16) + + first = det_gemm_linear_weight_gradient(inputs, grad_output) + second = det_gemm_linear_weight_gradient(inputs, grad_output) + + assert first.shape == (12, 8) + assert torch.equal(first, second) + assert torch.equal(first, torch.mm(grad_output.t(), inputs)) + + +def test_rocm_det_linear_preserves_autograd_and_bitwise_gradients(): + inputs = torch.randn( + (16, 8), device="cuda", dtype=torch.bfloat16, requires_grad=True + ) + weight = torch.randn( + (12, 8), device="cuda", dtype=torch.bfloat16, requires_grad=True + ) + grad_output = torch.randn((16, 12), device="cuda", dtype=torch.bfloat16) + op = DetGemmOp() + + first_output = op.linear(inputs, weight) + assert first_output.requires_grad + first_output.backward(grad_output) + first_input_grad = inputs.grad.detach().clone() + first_weight_grad = weight.grad.detach().clone() + + inputs.grad = None + weight.grad = None + second_output = op.linear(inputs, weight) + second_output.backward(grad_output) + + assert torch.equal(first_output, second_output) + assert inputs.grad is not None and inputs.grad.shape == inputs.shape + assert weight.grad is not None and weight.grad.shape == weight.shape + assert torch.equal(first_input_grad, inputs.grad) + assert torch.equal(first_weight_grad, weight.grad) + + +def test_vllm_ffn_provenance_reports_rocm_triton_and_fixed_tree(): + operator = VllmFFNOperator() + operator._set_runtime_provenance(4, "rocm_ipc_fixed_tree") + + execution = operator.provenance["execution"] + assert execution["runtime_platform"] == "rocm" + assert execution["actual_backend"] == "rlkernel.rocm.det_gemm_swiglu" + assert execution["deterministic_all_reduce_backend"] == "rocm_ipc_fixed_tree" + assert execution["triton_used"] is True diff --git a/tests/test_vime_linear_logp_provider.py b/tests/test_vime_linear_logp_provider.py index 681aad67..d8071a6f 100644 --- a/tests/test_vime_linear_logp_provider.py +++ b/tests/test_vime_linear_logp_provider.py @@ -12,6 +12,12 @@ from rl_engine.integrations import framework_operators from rl_engine.integrations.framework_operators import MegatronLogpOperator +from rl_engine.integrations.ablation import IntegrationPlan +from rl_engine.integrations.megatron import MegatronIntegration +from rl_engine.integrations.state import ( + clear_active_integration, + set_active_integration, +) from rl_engine.integrations.vime.linear_logp_provider import ( LinearLogpProviderUnavailable, LinearLogpResult, @@ -19,6 +25,11 @@ ) +@pytest.fixture(autouse=True) +def _select_rlkernel_logp_route(monkeypatch): + monkeypatch.setenv("RL_KERNEL_LOGP_CASE", "R/R") + + def _request(*, cp_rank: int = 0, with_entropy: bool = False, keep_mask=None): logits = torch.tensor( [[0.25, -0.5, 1.0, 0.1, -0.3, 0.6, -0.7, 0.4] for _ in range(3)], @@ -95,7 +106,9 @@ def test_provider_entropy_preserves_vime_semantics_and_autograd(): result = provider(request) reference_logits = request.logits.detach().clone().requires_grad_(True) log_probs = torch.log_softmax(reference_logits[:, :7], dim=-1) - reference_logp = log_probs[torch.arange(reference_logits.size(0)), request.target_ids] + reference_logp = log_probs[ + torch.arange(reference_logits.size(0)), request.target_ids + ] reference_entropy = -(log_probs.exp() * log_probs).sum(dim=-1) torch.testing.assert_close(result.logp.squeeze(-1), reference_logp) @@ -121,9 +134,16 @@ def __call__(self, hidden, weight, target_ids, bias, **_kwargs): torch.arange(target_ids.size(0)), target_ids ] + def from_local_logits(self, local_logits, target_ids, **_kwargs): + return torch.log_softmax(local_logits[:, :7], dim=-1)[ + torch.arange(target_ids.size(0)), target_ids + ] + import rl_engine.integrations.vime.linear_logp_provider as provider_module - monkeypatch.setattr(provider_module, "_default_strict_linear_logp", lambda: FakeLinearLogp()) + monkeypatch.setattr( + provider_module, "_default_strict_linear_logp", lambda: FakeLinearLogp() + ) request = _structural_request() result = provider(request) @@ -148,7 +168,9 @@ def provider(actual_request, *, linear_logp): ) wrapper = SimpleNamespace(backend_id="fake-linear-logp", provenance={}) - monkeypatch.setattr(framework_operators, "_require_nvidia_cuda", lambda *_args: None) + monkeypatch.setattr( + framework_operators, "_require_nvidia_cuda", lambda *_args: None + ) result = MegatronLogpOperator(provider, linear_logp=wrapper)(request) assert observed["context"] is request.context @@ -169,3 +191,20 @@ def test_provider_rejects_local_vocab_metadata_that_cannot_describe_tp_ownership with pytest.raises(LinearLogpProviderUnavailable, match="cover padded_vocab_size"): provider(request) + + +def test_production_route_rejects_rlkernel_provider_configuration(monkeypatch): + monkeypatch.setenv("RL_KERNEL_LOGP_CASE", "P/P") + integration = MegatronIntegration( + IntegrationPlan.from_case_ids(logp="P/P"), + rl_kernel_operators={}, + ) + clear_active_integration("megatron") + set_active_integration("megatron", integration) + try: + with pytest.raises(RuntimeError, match="omit --linear-logp-provider"): + provider(_request()) + finally: + clear_active_integration("megatron") + + assert "logp" not in integration.readback()["operators"] diff --git a/tests/test_vime_rocm_attention_topology.py b/tests/test_vime_rocm_attention_topology.py new file mode 100644 index 00000000..aafe384b --- /dev/null +++ b/tests/test_vime_rocm_attention_topology.py @@ -0,0 +1,130 @@ +# SPDX-License-Identifier: Apache-2.0 + +from __future__ import annotations + +from pathlib import Path + +import pytest +import torch + +from examples.vime_rocm_attention_ablation.run import MatrixConfig, build_plan +from rl_engine.kernels.ops.cuda.attention.cp_comm import ( + AttentionCPCommunicationPlan, + AttentionParallelSpec, +) +from rl_engine.kernels.ops.rocm.attention.strict_runtime import ( + RCCLAGRSAttentionCPCommunication, +) + +ROOT = Path(__file__).parents[1] + + +def _config(tmp_path: Path, **overrides) -> MatrixConfig: + values = { + "vime_root": tmp_path / "vime", + "rl_kernel_root": tmp_path / "rl-kernel", + "megatron_root": tmp_path / "megatron", + "model_root": tmp_path / "model", + "reference_checkpoint": tmp_path / "checkpoint", + "prompt_data": tmp_path / "prompts.jsonl", + "run_dir": tmp_path / "run", + "launcher": tmp_path / "launch.sh", + } + values.update(overrides) + return MatrixConfig(**values) + + +def test_default_topology_matches_pr377_colocated_tp4_cp2(tmp_path): + config = _config(tmp_path) + config.validate(require_paths=False) + + parameters = build_plan(config)["parameters"] + assert parameters["training"] == { + "num_gpus": 8, + "tensor_parallel_size": 4, + "context_parallel_size": 2, + "pipeline_parallel_size": 1, + "sequence_parallel": False, + "dtype": "bf16", + "attention_backend": "flash", + "attention_dropout": 0.0, + "hidden_dropout": 0.0, + } + assert parameters["rollout"]["num_gpus"] == 8 + assert parameters["rollout"]["engine_count"] == 2 + assert parameters["rollout"]["tensor_parallel_size"] == 4 + assert parameters["rollout"]["router_policy"] == "round_robin" + assert parameters["batch"]["rollout_batch_size"] == 2 + assert parameters["batch"]["samples_per_prompt"] == 1 + assert parameters["batch"]["global_batch_size"] == 2 + assert parameters["placement"] == { + "colocate": True, + "offload_train": False, + "offload_rollout": True, + } + + environment = build_plan(config)["arms"][0]["environment"] + assert environment["RLK_ABLATION_COLOCATE"] == "1" + assert environment["RLK_ABLATION_ROUTER_POLICY"] == "round_robin" + assert environment["RL_KERNEL_FFN_CASE"] == "R/R" + assert environment["RL_KERNEL_LOGP_CASE"] == "R/R" + assert environment["RL_KERNEL_VLLM_REAL_VOCAB_SIZE"] == "151936" + assert environment["RL_KERNEL_VLLM_PADDED_VOCAB_SIZE"] == "152064" + + +def test_colocated_topology_requires_training_to_cover_all_gpus(tmp_path): + config = _config(tmp_path, tensor_parallel_size=2) + with pytest.raises(ValueError, match="must use all visible GPUs"): + config.validate(require_paths=False) + + +def test_router_requires_at_least_one_request_per_engine(tmp_path): + config = _config( + tmp_path, + rollout_batch_size=1, + samples_per_prompt=2, + ) + with pytest.raises(ValueError, match="one request per rollout engine"): + config.validate(require_paths=False) + + +def test_rocm_cp_adapter_accepts_rccl_plan_without_widening_cuda_contract(monkeypatch): + plan = AttentionCPCommunicationPlan( + parallel=AttentionParallelSpec( + tp_world_size=4, + tp_rank=0, + cp_world_size=2, + cp_rank=0, + ), + backend="rccl_ag_rs", + status="implemented", + ) + communication = object.__new__(RCCLAGRSAttentionCPCommunication) + monkeypatch.setattr(torch.version, "hip", "test") + monkeypatch.setattr(torch.cuda, "is_available", lambda: True) + + communication._validate_cuda_plan(plan) + + assert plan.backend == "rccl_ag_rs" + + +def test_dashboard_cannot_overlap_ray_worker_port_range(tmp_path): + config = _config(tmp_path, ray_dashboard_port=18265) + with pytest.raises(ValueError, match="worker range"): + config.validate(require_paths=False) + + +def test_launcher_uses_pr377_torch_dist_actor_load_without_reference_model(): + launcher = ( + ROOT / "examples" / "vime_rocm_attention_ablation" / "launch_arm.sh" + ).read_text(encoding="utf-8") + + assert '--load "${RLK_ABLATION_REFERENCE_CHECKPOINT}"' in launcher + assert "--megatron-to-hf-mode" not in launcher + assert "--use-kl-loss" not in launcher + assert "--kl-loss-coef" not in launcher + assert "--linear-logp-provider" in launcher + assert "rl_engine.integrations.vime.linear_logp_provider.provider" in launcher + assert "--linear-logp-provider-mode strict" in launcher + assert '"${RL_KERNEL_FFN_CASE:-}" != "R/R"' in launcher + assert '"${RL_KERNEL_LOGP_CASE:-}" != "R/R"' in launcher diff --git a/tests/test_vime_tp4_example.py b/tests/test_vime_tp4_example.py new file mode 100644 index 00000000..6bed6327 --- /dev/null +++ b/tests/test_vime_tp4_example.py @@ -0,0 +1,214 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors + +import os +import subprocess +from dataclasses import asdict +from pathlib import Path + +from examples.vime_qwen3_8b_tp4_cp2_200.run_arm import ( + ARMS, + MEGATRON_ATTENTION_BACKEND, + RL_KERNEL_LINEAR_LOGP_PROVIDER, + _linear_logp_provider_args, +) +from examples.vime_qwen3_8b_tp4_cp2_200.run_supplement_suite import specs +from examples.vime_qwen3_8b_tp4_cp2_200.validate_run import ( + VIME_NATIVE_LINEAR_LOGP_MARKER, + _validate_readbacks, +) + + +def _production_operator(framework: str, target: str, module: str) -> dict: + return { + "framework": framework, + "target": target, + "module": module, + "case_id": "P/P", + "implementation": "production", + "backend_id": f"{framework}.production.{module}", + "call_count": 1, + "provenance": {"runtime_platform": "cuda"}, + } + + +def _production_readbacks() -> list[dict]: + megatron_modules = ("attention", "ffn") + vllm_modules = ("attention", "ffn", "logp") + return [ + { + "framework": "megatron", + "target": "training", + "installed_hooks": {module: module for module in megatron_modules}, + "operators": { + module: _production_operator("megatron", "training", module) + for module in megatron_modules + }, + "fallbacks": [], + }, + { + "framework": "vllm", + "target": "rollout", + "installed_hooks": {module: module for module in vllm_modules}, + "operators": { + module: _production_operator("vllm", "rollout", module) for module in vllm_modules + }, + "fallbacks": [], + }, + ] + + +def test_tp4_formal_matrix_pins_the_vime_qwen3_attention_backend(): + assert MEGATRON_ATTENTION_BACKEND == "fused" + + +def test_launcher_forces_cuda_graph_without_a_logp_provider(): + root = Path(__file__).parents[1] + launcher = root / "examples" / "vime_qwen3_8b_tp4_cp2_200" / "aligned_python_entrypoint.sh" + env = os.environ.copy() + env.update( + RL_KERNEL_REAL_PYTHON="/bin/echo", + RL_KERNEL_ROOT=str(root), + RL_KERNEL_VLLM_CUDAGRAPH_MAX_CAPTURE_SIZE="8", + ) + + result = subprocess.run( + [ + str(launcher), + "train.py", + "--rollout-batch-size", + "1", + "--n-samples-per-prompt", + "8", + ], + check=True, + capture_output=True, + text=True, + env=env, + ) + + assert "--linear-logp-provider" not in result.stdout + assert "--vllm-optimization-level 0" in result.stdout + assert '"cudagraph_mode":"FULL_DECODE_ONLY"' in result.stdout + assert '"cudagraph_capture_sizes":[1,2,3,4,5,6,7,8]' in result.stdout + assert "required vLLM full-decode CUDA Graph capture sizes" in result.stderr + + +def test_production_arms_do_not_install_the_rlkernel_logp_provider(): + assert _linear_logp_provider_args(ARMS["G00"]) == () + assert _linear_logp_provider_args(ARMS["G10"]) == () + + +def test_rlkernel_arms_install_the_strict_logp_provider(): + expected = ( + "--linear-logp-provider", + RL_KERNEL_LINEAR_LOGP_PROVIDER, + "--linear-logp-provider-mode", + "strict", + ) + + assert _linear_logp_provider_args(ARMS["G01"]) == expected + assert _linear_logp_provider_args(ARMS["G11"]) == expected + + +def test_module_ablation_matrix_changes_only_operator_routes(): + expected = { + "M000": ("P/P", "P/P", "P/P"), + "M100": ("R/R", "P/P", "P/P"), + "M010": ("P/P", "R/R", "P/P"), + "M001": ("P/P", "P/P", "R/R"), + "M110": ("R/R", "R/R", "P/P"), + "M101": ("R/R", "P/P", "R/R"), + "M011": ("P/P", "R/R", "R/R"), + "M111": ("R/R", "R/R", "R/R"), + } + + for group, cases in expected.items(): + arm = ARMS[group] + assert not arm.framework_use_rollout_logprobs + assert (arm.attention_case, arm.ffn_case, arm.logp_case) == cases + + +def test_module_ablation_logp_provider_follows_training_route(): + provider_groups = {"M001", "M101", "M011", "M111"} + expected_module_groups = { + "M000", + "M100", + "M010", + "M001", + "M110", + "M101", + "M011", + "M111", + } + for group in expected_module_groups: + has_provider = bool(_linear_logp_provider_args(ARMS[group])) + assert has_provider == (group in provider_groups) + assert expected_module_groups == {key for key in ARMS if key.startswith("M")} + + +def test_supplement_suite_uses_short_module_and_three_seed_precision_designs(): + module = specs("module") + assert len(module) == 8 + assert all(rounds == 8 and seed == 1234 for _, rounds, seed in module) + + precision = specs("precision") + assert len(precision) == 12 + assert {group for group, _, _ in precision} == {"G00", "G10", "G01", "G11"} + assert {seed for _, _, seed in precision} == {1234, 2345, 3456} + assert all(rounds == 8 for _, rounds, _ in precision) + + +def test_validator_accepts_native_vime_logp_evidence_for_production_arm(): + report = _validate_readbacks( + _production_readbacks(), + asdict(ARMS["G10"]), + VIME_NATIVE_LINEAR_LOGP_MARKER, + ) + + assert report["passed"] + training_logp = report["frameworks"]["megatron/training"]["modules"]["logp"] + assert training_logp["native_marker_present"] + assert training_logp["call_count"] == 0 + + +def test_validator_rejects_provider_readback_on_production_megatron_logp(): + readbacks = _production_readbacks() + contaminated = _production_operator("megatron", "training", "logp") + contaminated["backend_id"] = "pytorch-vocab-parallel-logp-ws2" + contaminated["provenance"] = { + "runtime_platform": "cuda", + "actual_backend": "rlkernel.linear_logp.bitwise.v1", + "deterministic_linear_logp": True, + "execution": {"strict_backend": True}, + } + readbacks[0]["installed_hooks"]["logp"] = "rlkernel-provider" + readbacks[0]["operators"]["logp"] = contaminated + + report = _validate_readbacks( + readbacks, + asdict(ARMS["G10"]), + VIME_NATIVE_LINEAR_LOGP_MARKER, + ) + + assert not report["passed"] + assert any("unexpectedly entered provider readback" in error for error in report["errors"]) + + +def test_validator_rejects_production_label_over_rlkernel_actual_backend(): + readbacks = _production_readbacks() + readbacks[1]["operators"]["logp"]["provenance"] = { + "runtime_platform": "cuda", + "actual_backend": "rlkernel.linear_logp.bitwise.v1", + } + + report = _validate_readbacks( + readbacks, + asdict(ARMS["G10"]), + VIME_NATIVE_LINEAR_LOGP_MARKER, + ) + + assert not report["passed"] + assert any( + "production route executed an RL-Kernel backend" in error for error in report["errors"] + )