Optimize GPTQ-Pro kernels for Ampere decode and prefill

[groxaxo]

Summary

Introduces specialized Ampere execution paths while retaining the existing general-shape kernel as a correctness fallback.

CUDA runtime

• Dedicated coalesced FP16-dequant/FP32-accumulating GEMV kernel for M <= 4 decode workloads.
• Four-warp CTA Tensor Core kernel for aligned GEMMs with a double-buffered cp.async pipeline.
• Coalesced 16-byte staging for activations, native GPTQ int32 qweight, and scales.
• LOP3-assisted INT4-to-FP16 fragment conversion.
• Vectorized FP16 output stores.
• Explicit auto, gemv, ampere, and legacy modes for validation and benchmarking.
• Validator-backed legacy path retained for unusual compatible edge shapes.

Memory and runtime

• All CUDA paths consume the checkpoint's native int32 qweight directly.
• Removes the persistent pair-packed duplicate of quantized model weights and the module-initialization repacking transform.
• Retains the historical byte-view conversion helper only for compatibility/inspection tooling; inference does not call it.
• Normalizes qweight/scale contiguity once during post_init().
• Fixes group_size=-1 handling in runtime g_idx validation and kernel invocation.
• Versions the extension ABI as gptqmodel_gptq_pro_kernels_v2 so stale cached/prebuilt binaries cannot be loaded with the new binding.
• Adds GPTQMODEL_GPTQ_PRO_KERNEL as an expert/debug dispatch override.

Validation and measurement

• Replaces the standalone validator with native-qweight coverage for fragment decode, forced GEMV, forced Ampere, AUTO dispatch, M/N tails, and legacy odd-shape fallback.
• Adds standalone CUDA compilation for sm_80, sm_86, sm_87, and compute_87 PTX.
• Adds compilation/import of the actual PyTorch C++/CUDA binding.
• Adds a native-qweight raw-kernel benchmark with forced-mode timing, TF32-disabled PyTorch references, and numerical checks.
• Adds CPU tests for qweight layout, dispatch normalization, no duplicate runtime buffer, Ampere source contracts, and compile-matrix coverage.

Validation status

• PASS: targeted Ruff formatting/lint and Python syntax checks.
• PASS: standalone CUDA validator compilation for sm_80, sm_86, sm_87, and forward-compatible PTX.
• PASS: clean compilation and import of the real PyTorch 2.8.0/CUDA 12.6 C++/CUDA extension binding.
• PASS: isolated CPU qweight/dispatch/source-contract tests.
• NOT RUN HERE: execution of the numerical CUDA validator and performance benchmark on a physical Ampere GPU. No measured speedup is claimed solely from source implementation or compile success.
• The repository-wide Ruff workflow has a pre-existing unrelated lint backlog; this PR uses a targeted clean-file gate for its modified Python surface.