Fix HIP FA launch bounds and Windows CPU packaging

Parenthesize the type-aware vec FlashAttention launch-bound expression so HIP's variadic __launch_bounds__ macro does not split the template argument list. Replace the Windows CPU release package's hard-coded VS2022 Enterprise OpenMP runtime lookup with a Visual Studio install-root search that works on newer runner images.

Author: Anbeeld

.github/workflows/release.yml

@@ -801,15 +801,8 @@ jobs:
           $ErrorActionPreference = "Stop"
           $bin = ".\build\bin\Release"
           Copy-Item ".\LICENSE" $bin
-          $redistRoot = "C:\Program Files\Microsoft Visual Studio\2022\Enterprise\VC\Redist\MSVC"
-          $omp = Get-ChildItem $redistRoot -Recurse -Filter "libomp140.x86_64.dll" |
-            Where-Object { $_.FullName -like "*\debug_nonredist\x64\Microsoft.VC143.OpenMP.LLVM\libomp140.x86_64.dll" } |
-            Sort-Object FullName -Descending |
-            Select-Object -First 1
-          if (-not $omp) {
-            throw "VC143 x64 libomp140.x86_64.dll not found under $redistRoot"
-          }
-          Copy-Item $omp.FullName $bin -Force
+          $omp = & ".\scripts\find-msvc-openmp-runtime.ps1" -Architecture x64
+          Copy-Item -LiteralPath $omp -Destination $bin -Force
           python scripts\verify-windows-package.py $bin
           & "$bin\llama-server.exe" --version
           & "$bin\llama-cli.exe" --version

ggml/src/ggml-cuda/fattn-vec.cuh

@@ -28,7 +28,7 @@ static constexpr __host__ __device__ int ggml_cuda_fattn_vec_get_min_blocks() {
 #pragma clang diagnostic ignored "-Wpass-failed"
 #endif // __clang__
 template<int D, int ncols, ggml_type type_K, ggml_type type_V, bool use_logit_softcap> // D == head size
-__launch_bounds__(ggml_cuda_fattn_vec_get_nthreads_device(), ggml_cuda_fattn_vec_get_min_blocks<type_K, type_V>())
+__launch_bounds__(ggml_cuda_fattn_vec_get_nthreads_device(), (ggml_cuda_fattn_vec_get_min_blocks<type_K, type_V>()))
 static __global__ void flash_attn_ext_vec(
         const char * Q_ptr,
         const char * K_ptr,

scripts/find-msvc-openmp-runtime.ps1

@@ -0,0 +1,88 @@
+param(
+    [ValidateSet("x64")]
+    [string] $Architecture = "x64"
+)
+
+$ErrorActionPreference = "Stop"
+
+function Get-VersionOrZero($File) {
+    try {
+        return [Version] $File.VersionInfo.FileVersion
+    } catch {
+        return [Version] "0.0.0.0"
+    }
+}
+
+$programFilesX86 = [Environment]::GetEnvironmentVariable("ProgramFiles(x86)")
+$programRoots = @($env:ProgramFiles, $programFilesX86) | Where-Object { $_ }
+
+$vsRoots = @()
+$vswhereCandidates = foreach ($programRoot in $programRoots) {
+    Join-Path $programRoot "Microsoft Visual Studio\Installer\vswhere.exe"
+}
+$vswhereCandidates = $vswhereCandidates |
+    Where-Object { Test-Path -LiteralPath $_ } |
+    Sort-Object -Unique
+
+$vswhereArgSets = @(
+    ,@("-all", "-products", "*", "-requires", "Microsoft.VisualStudio.Component.VC.Redist.14.Latest", "-property", "installationPath"),
+    ,@("-all", "-products", "*", "-requires", "Microsoft.VisualStudio.Component.VC.Tools.x86.x64", "-property", "installationPath"),
+    ,@("-all", "-products", "*", "-property", "installationPath")
+)
+
+foreach ($vswhere in $vswhereCandidates) {
+    foreach ($argSet in $vswhereArgSets) {
+        $vsRoots += & $vswhere @argSet 2>$null
+    }
+}
+
+foreach ($programRoot in $programRoots) {
+    $visualStudioRoot = Join-Path $programRoot "Microsoft Visual Studio"
+    if (-not (Test-Path -LiteralPath $visualStudioRoot)) {
+        continue
+    }
+
+    foreach ($versionRoot in Get-ChildItem -LiteralPath $visualStudioRoot -Directory -ErrorAction SilentlyContinue) {
+        foreach ($editionRoot in Get-ChildItem -LiteralPath $versionRoot.FullName -Directory -ErrorAction SilentlyContinue) {
+            $vsRoots += $editionRoot.FullName
+        }
+    }
+}
+
+$vsRoots = $vsRoots |
+    Where-Object { $_ -and (Test-Path -LiteralPath $_) } |
+    Sort-Object -Unique
+
+$searchRoots = foreach ($vsRoot in $vsRoots) {
+    foreach ($subdir in @("VC\Redist\MSVC", "VC\Tools\MSVC")) {
+        $path = Join-Path $vsRoot $subdir
+        if (Test-Path -LiteralPath $path) {
+            $path
+        }
+    }
+}
+$searchRoots = $searchRoots | Sort-Object -Unique
+
+$candidates = foreach ($searchRoot in $searchRoots) {
+    Get-ChildItem -LiteralPath $searchRoot -Recurse -File -Filter "libomp140.x86_64.dll" -ErrorAction SilentlyContinue |
+        Where-Object {
+            $_.FullName -match "\\$Architecture\\" -and (
+                $_.FullName -match "\\Microsoft\.VC\d+\.OpenMP\.LLVM\\" -or
+                $_.FullName -match "\\bin\\Host(?:x86|x64)\\$Architecture\\"
+            )
+        }
+}
+
+$runtime = $candidates |
+    Sort-Object `
+        @{ Expression = { Get-VersionOrZero $_ } },
+        @{ Expression = { $_.FullName } } `
+        -Descending |
+    Select-Object -First 1
+
+if (-not $runtime) {
+    $searched = if ($searchRoots) { $searchRoots -join "; " } else { "<none>" }
+    throw "MSVC $Architecture libomp140.x86_64.dll not found. Searched: $searched"
+}
+
+$runtime.FullName

tests/test-cuda-fattn-vec-policy.cpp

@@ -71,8 +71,10 @@ int main(int argc, char ** argv) {
     ok &= expect(helpers.find("? 2 : 1") != std::string::npos,
         "VEC FA min blocks must be 2 only for Turbo/TCQ pairs and 1 otherwise");
 
-    ok &= expect(kernel.find("ggml_cuda_fattn_vec_get_min_blocks<type_K, type_V>()") != std::string::npos,
+    ok &= expect(kernel.find("ggml_cuda_fattn_vec_get_nthreads_device(), (ggml_cuda_fattn_vec_get_min_blocks<type_K, type_V>())") != std::string::npos,
         "VEC FA launch_bounds must use the K/V type-aware min-block policy");
+    ok &= expect(kernel.find("ggml_cuda_fattn_vec_get_nthreads_device(), ggml_cuda_fattn_vec_get_min_blocks<type_K, type_V>()") == std::string::npos,
+        "VEC FA launch_bounds min-block expression must be parenthesized for HIP's variadic macro parser");
     ok &= expect(kernel.find("ggml_cuda_fattn_vec_get_nthreads_device(), 2") == std::string::npos,
         "VEC FA launch_bounds must not force all cache types to minBlocksPerSM=2");