release v0.5.2: native Windows/MSVC support + upstream PR notes

Bump 0.5.1 -> 0.5.2. Adds UPSTREAM_PR_NOTES.md documenting the MSVC
portability fixes (enum signedness, edition-2024 unsafe, gated amdgcn
GPU-sort kernel) for a future PR to RustNSparks/rocm-rs.

Cargo.toml

@@ -1,6 +1,6 @@
 [package]
 name = "rocm-rs"
-version = "0.5.1"
+version = "0.5.2"
 edition = "2024"
 description = "Rust bindings for AMD ROCm libraries"
 license = "MIT"

UPSTREAM_PR_NOTES.md

@@ -0,0 +1,72 @@
+# Upstream PR notes — native Windows / MSVC support for rocm-rs
+
+This fork (`hanzoai/rocm-rs`, tagged **v0.5.2**) makes rocm-rs build and run on
+**native Windows (MSVC)** with ROCm 7.13.0 (TheRock / `repo.amd.com` gfx1151
+wheels), in addition to Linux. Everything here is a candidate for a PR back to
+`RustNSparks/rocm-rs`. Captured now so the PR is easy to assemble later.
+
+## Why
+Stock 0.5.1 builds on Linux but fails on native Windows/MSVC with **200+
+compile errors**. All of them are platform-ABI differences, not logic bugs.
+
+## Root causes & fixes (4 commits on top of 0.5.1)
+
+### 1. C-enum signedness — ~50× `error[E0308] expected u32, found i32`
+C enums with only non-negative values are `unsigned int`/`u32` under GCC/Clang
+(Linux) but **signed `int`/`i32` under MSVC**. bindgen follows the ABI, so every
+`<lib>_status` / enum alias (`rocblas_status`, `rocrand_status`,
+`rocrand_rng_type`, `rocrand_ordering`, `rocfft_status_e`,
+`rocfft_transform_type`, …) is `u32` on Linux and `i32` on Windows. The crate
+hardcodes literal `u32` at FFI return types, `Error::new`, and enum-arg params.
+
+**Fix:** reference the bindgen-generated type instead of a literal `u32`
+(`ffi::rocblas_status`, `bindings::rocrand_rng_type`, …). **No-op on Linux**
+(the alias *is* `u32` there); correct on MSVC. Files: `rocblas/macros.rs`,
+`rocrand/{error,generator,mod}.rs`, `rocfft/{error,plan,description}.rs`,
+`rocarray/random.rs`. → clean, low-risk, clearly upstreamable.
+
+### 2. Edition-2024 `unsafe_op_in_unsafe_fn` — 17× `error[E0133]`
+The crate is `edition = "2024"`, where `unsafe_op_in_unsafe_fn` is deny-by-
+default. This Linux toolchain emitted it as a *warning* while MSVC enforced it
+as a hard *error* (toolchain lint-config difference). Thin `pub unsafe fn`
+rocBLAS forwarders call their unsafe trait method with no `unsafe {}` block.
+
+**Fix:** wrap the calls in `unsafe { }`. Identical on both platforms. Files:
+`rocblas/level{1,2,3}.rs`. → clean, clearly upstreamable (and arguably a
+correctness fix the Linux build should also enforce).
+
+### 3. amdgcn GPU-sort kernel — `error[E0463]: can't find crate for core`
+`hip/memory_ext/sorting.rs` embeds an amdgcn kernel via the external
+`rocm_kernel_macros::amdgpu_kernel_finalize!` proc-macro, which shells out a
+nested `cargo build` for `amdgcn-amd-amdhsa` (nightly + `-Zbuild-std`,
+rust-src). That nested build **inherits the parent rustc environment** and
+drops off the nightly toolchain on non-Linux hosts → core can't be built.
+
+**Fix (fork-local, pragmatic):** gate the embedded kernel behind a **default-on
+`gpu-sort` feature** with an empty-bytes fallback. The crate then builds
+everywhere; GPU sort becomes a runtime no-op when the feature is off (LLM
+inference never calls it). Files: `hip/memory_ext/sorting.rs`, `Cargo.toml`.
+
+**Real upstream fix to discuss:** the bug is in `rocm_kernel_macros`'s nested
+build invocation — it only does `.env_remove("RUSTUP_TOOLCHAIN")`; it should
+also remove `RUSTC`, `CARGO`, `CARGO_BUILD_TARGET`, `RUSTFLAGS`,
+`CARGO_ENCODED_RUSTFLAGS` so the nested `cargo` resolves the kernel's own
+`rust-toolchain.toml` (nightly) + `.cargo/config.toml` (`build-std`) cleanly on
+every host. The feature gate is the safe interim; the macro fix is the proper
+cure. Suggest opening that as a separate issue/PR on `rocm_kernel_macros`.
+
+## Verification
+- **Linux:** `cargo check --no-default-features --features macros` *and* default
+  features both reach `Finished` — proving every type change is a no-op there.
+- **Windows (MSVC, ROCm 7.13.0, gfx1151 / Radeon 8060S):** builds clean; the
+  downstream LLM engine builds and runs inference at **158.7 tok/s** decode
+  (Qwen3-0.6B-Q8), ~36% faster than the same engine under WSL2 (no `/dev/dxg`
+  per-dispatch overhead).
+
+## PR packaging
+1. Commits #1 (enum types) and #2 (unsafe blocks) → one PR, "native Windows /
+   MSVC build support". Low risk, no Linux behavior change.
+2. Commit #3 (gpu-sort gate) → either fold in with a note, or hold pending the
+   `rocm_kernel_macros` env fix. Maintainer's call.
+3. Do **not** include regenerated `bindings.rs` (per-toolchain build artifacts;
+   bindgen rewrites them against the locally-installed ROCm headers).