Fix CUDA 13 build error in gqa_unfused_attention.cu

Merged
TI
6mo ago
+1-1

Description

Fix CUDA 13 build failure introduced by PR #28198 (commit 997c4798).

Root cause: gqa_unfused_attention.cu directly includes <cub/cub.cuh>, which on CUDA 13.x transitively pulls in CCCL's tcgen05_ld.h. That header uses __out as a parameter name in inline PTX assembly, but on Windows MSVC the SAL annotation macro #define __out expands it, turning __out[0] into [0] — causing a parse error.

Fix: Changed #include <cub/cub.cuh> to #include "core/providers/cuda/cu_inc/cub.cuh" — the existing ORT wrapper that #undef __out before including CUB.

Motivation and Context

The CUDA 13 packaging pipeline (py-cuda13-packaging-pipeline.yml) has been failing since PR #28198 was merged, with errors like:

E:/_work/_temp/v13.0/include/cccl/cuda/__ptx/instructions/generated/tcgen05_ld.h(20): error : expected an identifier
      asm("tcgen05.ld.sync.aligned.16x64b.x1.b32 {%0}, [%1];" : "=r"( [0]) : "r"(__taddr) : "memory");