diff --git a/.ci/scripts/install_nvidia_driver_windows.ps1 b/.ci/scripts/install_nvidia_driver_windows.ps1 new file mode 100644 index 00000000000..5b18ceac284 --- /dev/null +++ b/.ci/scripts/install_nvidia_driver_windows.ps1 @@ -0,0 +1,28 @@ +# Copyright (c) Meta Platforms, Inc. and affiliates. +# All rights reserved. +# +# This source code is licensed under the BSD-style license found in the +# LICENSE file in the root directory of this source tree. + +# Installs the NVIDIA display driver on a Windows GPU runner. The runner images carry the +# GPU but no driver, so without this nvcuda.dll is missing and every CUDA call fails with +# cudaErrorInsufficientDriver. The same driver, from the same place, that PyTorch's Windows +# CUDA smoke tests install first (.ci/pytorch/windows/internal/driver_update.bat in +# pytorch/pytorch), so both projects test against one driver. +$ErrorActionPreference = "Stop" +$version = "580.88" +$installer = Join-Path $env:RUNNER_TEMP "$version-data-center-tesla-desktop-win10-win11-64bit-dch-international.exe" +$url = "https://ossci-windows.s3.amazonaws.com/$(Split-Path $installer -Leaf)" + +$nvcuda = Join-Path $env:SystemRoot "System32\nvcuda.dll" +if (Test-Path $nvcuda) { + Write-Host "NVIDIA driver already present: $((Get-Item $nvcuda).VersionInfo.FileVersion)" + exit 0 +} +curl.exe --retry 3 -fsSL $url --output $installer +if ($LASTEXITCODE -ne 0) { throw "downloading the NVIDIA driver from $url failed" } +$process = Start-Process -FilePath $installer -ArgumentList "-s", "-noreboot" -Wait -PassThru +Remove-Item $installer -ErrorAction SilentlyContinue +if ($process.ExitCode -ne 0) { throw "the NVIDIA driver installer exited with $($process.ExitCode)" } +if (-not (Test-Path $nvcuda)) { throw "the NVIDIA driver installed but $nvcuda is missing" } +Write-Host "NVIDIA driver $version installed: $((Get-Item $nvcuda).VersionInfo.FileVersion)" diff --git a/.ci/scripts/test_cuda_runtime_windows.ps1 b/.ci/scripts/test_cuda_runtime_windows.ps1 new file mode 100644 index 00000000000..9e27aabf848 --- /dev/null +++ b/.ci/scripts/test_cuda_runtime_windows.ps1 @@ -0,0 +1,44 @@ +# Copyright (c) Meta Platforms, Inc. and affiliates. +# All rights reserved. +# +# This source code is licensed under the BSD-style license found in the +# LICENSE file in the root directory of this source tree. + +# Builds and runs the CUDA backend's runtime C++ tests on a Windows GPU runner, the +# counterpart of the unittest-cuda-runtime job in cuda.yml. The Windows runners carry +# data-center GPUs in TCC mode, which have no memory pools, so this is also what covers the +# allocator's synchronous fallback; the Linux job covers the pool path. +$ErrorActionPreference = "Stop" + +# Windows PowerShell 5.1 does not stop on a failing native command. +function Invoke-Native { + param([Parameter(Mandatory = $true)][scriptblock]$Command) + & $Command + if ($LASTEXITCODE -ne 0) { + throw "exit code ${LASTEXITCODE}: $Command" + } +} + +& "$PSScriptRoot\install_nvidia_driver_windows.ps1" + +$cudaNvcc = Join-Path $env:CUDA_HOME "bin\nvcc.exe" +if (-not (Test-Path $cudaNvcc)) { throw "CUDA compiler not found at '$cudaNvcc'" } +$env:CUDACXX = $cudaNvcc +$env:PATH = "$env:CUDA_HOME\bin\x64;$env:CUDA_HOME\bin;$env:PATH" + +$tests = @( + "test_cuda_allocator", + "test_cuda_mutable_state", + "test_cuda_weight_cache", + "test_cuda_guard", + "test_cuda_stream_guard" +) +$numCores = [Math]::Max([Environment]::ProcessorCount - 1, 1) +Invoke-Native { + cmake --preset llm-release-cuda -DEXECUTORCH_BUILD_TESTS=ON -DCMAKE_CXX_STANDARD=20 ` + -T "cuda=$env:CUDA_HOME" "-DCMAKE_CUDA_COMPILER=$cudaNvcc" "-DCUDAToolkit_ROOT=$env:CUDA_HOME" +} +Invoke-Native { cmake --build cmake-out --config Release -j $numCores --target $tests } +foreach ($test in $tests) { + Invoke-Native { ctest --test-dir cmake-out -C Release -R "^$test`$" --output-on-failure -V } +} diff --git a/.github/workflows/cuda-windows.yml b/.github/workflows/cuda-windows.yml index 27d76258c85..ae0068d475c 100644 --- a/.github/workflows/cuda-windows.yml +++ b/.github/workflows/cuda-windows.yml @@ -223,6 +223,40 @@ jobs: .ci/scripts/test_model_e2e_windows.ps1 -Device cuda-windows -HfModel '${{ matrix.model_repo }}/${{ matrix.model_name }}' -QuantName '${{ matrix.quant }}' -ModelDir \$artifactDir -ExpectedCudaVersion '13.0' }" + # The counterpart of cuda.yml's unittest-cuda-runtime on Windows. The Windows GPU runners + # carry data-center GPUs in TCC mode, which have no memory pools, so this also covers the + # allocator's synchronous fallback that the Linux job cannot reach. + unittest-cuda-runtime-windows: + name: unittest-cuda-runtime-windows + needs: [changed-files, run-decision] + if: | + contains(needs.changed-files.outputs.changed-files, 'backends/cuda') || + contains(needs.changed-files.outputs.changed-files, 'backends/aoti') || + contains(needs.changed-files.outputs.changed-files, 'extension/cuda') || + contains(needs.changed-files.outputs.changed-files, '.github/workflows/cuda-windows.yml') || + contains(needs.changed-files.outputs.changed-files, '.ci/scripts/test_cuda_runtime_windows.ps1') || + contains(needs.changed-files.outputs.changed-files, '.ci/scripts/install_nvidia_driver_windows.ps1') || + needs.run-decision.outputs.is-full-run == 'true' + uses: pytorch/test-infra/.github/workflows/windows_job.yml@main + with: + timeout: 120 + runner: windows.g5.4xlarge.nvidia.gpu + gpu-arch-type: cuda + gpu-arch-version: "13.0" + ref: ${{ github.event_name == 'pull_request' && github.event.pull_request.head.sha || github.sha }} + script: | + git config --global http.sslBackend openssl + git submodule update --init --recursive + conda init powershell + powershell -Command "& { + Set-PSDebug -Trace 1 + \$ErrorActionPreference = 'Stop' + \$env:CUDA_HOME = 'C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v13.0' + \$env:CUDA_PATH = \$env:CUDA_HOME + .ci/scripts/setup-windows.ps1 + .ci/scripts/test_cuda_runtime_windows.ps1 + }" + delete-model-cuda-windows-artifacts: name: delete-model-cuda-windows-artifacts # The exports exist only to hand the models to the Windows job above, about diff --git a/backends/cuda/runtime/cuda_allocator.cpp b/backends/cuda/runtime/cuda_allocator.cpp index 6d046764953..a48d7d6d438 100644 --- a/backends/cuda/runtime/cuda_allocator.cpp +++ b/backends/cuda/runtime/cuda_allocator.cpp @@ -121,6 +121,38 @@ cudaMemPool_t mem_pool_for(int device) { return pool; } +// Whether `device` supports the stream-ordered allocator. Devices without +// memory pools, such as data-center GPUs in TCC mode on Windows, reject +// cudaMallocAsync and cudaMemPoolCreate with cudaErrorNotSupported, so their +// allocations take the synchronous path instead. Cached because it is asked on +// every allocation. An attribute query that fails is treated as supported, +// which keeps the stream-ordered path and its own error reporting. +bool memory_pools_supported(int device) { + static std::mutex mutex; + static std::unordered_map cache; + const std::lock_guard lock(mutex); + const auto it = cache.find(device); + if (it != cache.end()) { + return it->second; + } + int value = 0; + const cudaError_t err = + cudaDeviceGetAttribute(&value, cudaDevAttrMemoryPoolsSupported, device); + if (err != cudaSuccess) { + (void)cudaGetLastError(); + } + const bool supported = err != cudaSuccess || value != 0; + if (!supported) { + ET_LOG( + Info, + "CUDA device %d does not support memory pools; stream-ordered " + "allocations use cudaMalloc and cudaFree instead.", + device); + } + cache.emplace(device, supported); + return supported; +} + #endif // !EXECUTORCH_USE_HIP Error copy_impl( @@ -426,6 +458,13 @@ Result CudaAllocator::allocate_async( (void)cudaGetLastError(); stream_device = -1; } + const int target = device >= 0 ? device : stream_device; + if (target >= 0 && !memory_pools_supported(target)) { + // Without memory pools there is no stream-ordered allocation to make. A + // synchronous allocation is usable from any stream at once, so the caller's + // ordering still holds; deallocate_async frees it to match. + return CudaAllocator::instance().allocate(nbytes, index); + } cudaMemPool_t pool = (device >= 0 && device == stream_device) ? mem_pool_for(device) : nullptr; if (device >= 0) { @@ -460,6 +499,23 @@ void CudaAllocator::deallocate_async( return; } +#if !defined(EXECUTORCH_USE_HIP) + // Decided by the device the memory lives on, the same answer allocate_async + // reached for it. cudaFree waits for all work on the device first, so work + // still using the block on `stream` finishes before it is released. + cudaPointerAttributes attributes{}; + if (cudaPointerGetAttributes(&attributes, ptr) == cudaSuccess) { + if (attributes.type == cudaMemoryTypeDevice && attributes.device >= 0 && + !memory_pools_supported(attributes.device)) { + CudaAllocator::instance().deallocate( + ptr, static_cast(attributes.device)); + return; + } + } else { + (void)cudaGetLastError(); + } +#endif + cudaError_t err = cudaFreeAsync(ptr, stream); if (err != cudaSuccess) { ET_LOG( @@ -535,7 +591,11 @@ void CudaAllocator::release_cached_memory(DeviceIndex index) { // Device scoped, unlike everything else in this function: it releases // unused graph memory cached by every user of the device, so another // library in this process pays to map its own graph allocations again. - // Nothing breaks, since only unused blocks go. + // Nothing breaks, since only unused blocks go. Graph memory is served by + // memory pools, so a device without them has none to release. + if (!memory_pools_supported(device)) { + continue; + } const cudaError_t graph_err = cudaDeviceGraphMemTrim(device); if (graph_err != cudaSuccess) { ET_LOG( diff --git a/backends/cuda/runtime/shims/rand.cu b/backends/cuda/runtime/shims/rand.cu index b18d8f79079..8fe75768743 100644 --- a/backends/cuda/runtime/shims/rand.cu +++ b/backends/cuda/runtime/shims/rand.cu @@ -11,6 +11,7 @@ #include #include #include +#include #include #include @@ -60,7 +61,14 @@ static std::once_flag g_rng_init_flag; // from any thread are no-ops thanks to std::call_once. void ensure_rng_init(cudaStream_t stream) { std::call_once(g_rng_init_flag, [&]() { - cudaMallocAsync(&d_rng, sizeof(RngState), stream); + // Through the backend allocator, which falls back to cudaMalloc on a device + // without memory pools, where cudaMallocAsync is not supported. + auto rng = CudaAllocator::allocate_async(sizeof(RngState), -1, stream); + if (!rng.ok()) { + ET_LOG(Error, "rand: allocating the RNG state failed"); + return; + } + d_rng = static_cast(rng.get()); RngState h; h.seed = static_cast(time(nullptr)); h.counter = 0; diff --git a/backends/cuda/runtime/shims/sort.cu b/backends/cuda/runtime/shims/sort.cu index 36312df64ba..3c578690e9a 100644 --- a/backends/cuda/runtime/shims/sort.cu +++ b/backends/cuda/runtime/shims/sort.cu @@ -17,6 +17,7 @@ #include #include +#include #include #include #include @@ -111,21 +112,22 @@ void launch_permute( } // Stream-ordered scratch for thrust. With par_nosync this keeps each slice sort -// from blocking on cudaMalloc/cudaFree and synchronizing the stream. +// from blocking on cudaMalloc/cudaFree and synchronizing the stream. Through the +// backend allocator, which falls back to those on a device without memory pools. struct StreamOrderedAllocator { using value_type = char; char* allocate(std::ptrdiff_t bytes) { - void* ptr = nullptr; - if (cudaMallocAsync(&ptr, static_cast(bytes), stream) != - cudaSuccess) { + auto ptr = + CudaAllocator::allocate_async(static_cast(bytes), -1, stream); + if (!ptr.ok()) { throw std::bad_alloc(); } - return static_cast(ptr); + return static_cast(ptr.get()); } void deallocate(char* ptr, size_t) { - (void)cudaFreeAsync(ptr, stream); + CudaAllocator::deallocate_async(ptr, -1, stream); } cudaStream_t stream; @@ -318,14 +320,21 @@ AOTITorchError aoti_torch_cuda_sort_stable( inner_size *= input_sizes[d]; } - ET_CUDA_CHECK_OR_RETURN_ERROR(cudaMallocAsync( - &temp_values_buf, - static_cast(total_elements * elem_size), - stream)); - ET_CUDA_CHECK_OR_RETURN_ERROR(cudaMallocAsync( - &temp_indices_buf, - static_cast(total_elements) * sizeof(int64_t), - stream)); + auto values = CudaAllocator::allocate_async( + static_cast(total_elements * elem_size), -1, stream); + ET_CHECK_OR_RETURN_ERROR( + values.ok(), + MemoryAllocationFailed, + "sort: allocating the transposed values failed"); + temp_values_buf = values.get(); + auto indices = CudaAllocator::allocate_async( + static_cast(total_elements) * sizeof(int64_t), -1, stream); + if (!indices.ok()) { + CudaAllocator::deallocate_async(temp_values_buf, -1, stream); + ET_LOG(Error, "sort: allocating the transposed indices failed"); + return Error::MemoryAllocationFailed; + } + temp_indices_buf = indices.get(); // Gather: [outer, sort, inner] → [outer, inner, sort] launch_permute( @@ -450,8 +459,8 @@ AOTITorchError aoti_torch_cuda_sort_stable( stream); ET_CUDA_KERNEL_LAUNCH_CHECK_OR_RETURN_ERROR(); - ET_CUDA_CHECK_OR_RETURN_ERROR(cudaFreeAsync(temp_values_buf, stream)); - ET_CUDA_CHECK_OR_RETURN_ERROR(cudaFreeAsync(temp_indices_buf, stream)); + CudaAllocator::deallocate_async(temp_values_buf, -1, stream); + CudaAllocator::deallocate_async(temp_indices_buf, -1, stream); } return Error::Ok; diff --git a/backends/cuda/runtime/test/test_cuda_allocator.cpp b/backends/cuda/runtime/test/test_cuda_allocator.cpp index 4ffa81a98e0..2214b375877 100644 --- a/backends/cuda/runtime/test/test_cuda_allocator.cpp +++ b/backends/cuda/runtime/test/test_cuda_allocator.cpp @@ -205,15 +205,86 @@ uint64_t reserved_bytes(cudaMemPool_t pool) { cudaSuccess); return reserved; } + +bool device_supports_memory_pools(int device) { + int value = 0; + return cudaDeviceGetAttribute( + &value, cudaDevAttrMemoryPoolsSupported, device) == cudaSuccess && + value != 0; +} } // namespace +// The pool tests below only mean something on a device with memory pools. +// Data-center GPUs in TCC mode on Windows have none, and there the allocator +// takes the synchronous path the two tests after this fixture check. +class CudaAllocatorPoolTest : public CudaAllocatorTest { + protected: + void SetUp() override { + CudaAllocatorTest::SetUp(); + if (!IsSkipped() && !device_supports_memory_pools(0)) { + GTEST_SKIP() << "device 0 does not support memory pools"; + } + } +}; + +// Works on every device: with memory pools through the pool, without them +// through cudaMalloc. This is what failed on a device without pools, where +// cudaMallocAsync returned cudaErrorNotSupported and nothing fell back. +TEST_F(CudaAllocatorTest, AllocateAsyncRoundtrip) { + cudaStream_t stream; + ASSERT_EQ(cudaStreamCreate(&stream), cudaSuccess); + + constexpr size_t kBytes = 1u << 20; + auto res = CudaAllocator::allocate_async(kBytes, 0, stream); + ASSERT_TRUE(res.ok()) << "allocate_async failed on device 0"; + std::vector h_src(kBytes, 9), h_dst(kBytes, 0); + ASSERT_EQ( + cudaMemcpyAsync( + res.get(), h_src.data(), kBytes, cudaMemcpyHostToDevice, stream), + cudaSuccess); + ASSERT_EQ( + cudaMemcpyAsync( + h_dst.data(), res.get(), kBytes, cudaMemcpyDeviceToHost, stream), + cudaSuccess); + ASSERT_EQ(cudaStreamSynchronize(stream), cudaSuccess); + EXPECT_EQ(h_src, h_dst); + + CudaAllocator::deallocate_async(res.get(), 0, stream); + ASSERT_EQ(cudaStreamSynchronize(stream), cudaSuccess); + ASSERT_EQ(cudaStreamDestroy(stream), cudaSuccess); +} + +// Without memory pools the allocator must not create one or leave an error +// behind, and releasing cached memory must have nothing to do. +TEST_F(CudaAllocatorTest, FallsBackWithoutMemoryPools) { + if (device_supports_memory_pools(0)) { + GTEST_SKIP() << "device 0 supports memory pools; covered by the pool tests"; + } + cudaStream_t stream; + ASSERT_EQ(cudaStreamCreate(&stream), cudaSuccess); + // Earlier tests switch to a missing device on purpose, which leaves + // cudaErrorInvalidDevice as this thread's last error; only the calls below + // are under test. + (void)cudaGetLastError(); + + auto res = CudaAllocator::allocate_async(4096, -1, stream); + ASSERT_TRUE(res.ok()); + EXPECT_EQ(CudaAllocator::pool_for_device(0), nullptr) + << "no pool can exist on a device without memory pools"; + CudaAllocator::deallocate_async(res.get(), -1, stream); + CudaAllocator::release_cached_memory(-1); + EXPECT_EQ(cudaGetLastError(), cudaSuccess); + + ASSERT_EQ(cudaStreamDestroy(stream), cudaSuccess); +} + // The delegate allocates from a pool it owns, so its retained memory must not // land in the device default pool that other users of the async allocator // share. // The retention threshold is the whole point of owning a pool: at the default // of zero the driver empties it on every synchronize. Nothing else in this // suite notices a smaller value, so it is asserted directly. -TEST_F(CudaAllocatorTest, PoolRetainsMemoryWithoutLimit) { +TEST_F(CudaAllocatorPoolTest, PoolRetainsMemoryWithoutLimit) { cudaStream_t stream; ASSERT_EQ(cudaStreamCreate(&stream), cudaSuccess); @@ -236,7 +307,7 @@ TEST_F(CudaAllocatorTest, PoolRetainsMemoryWithoutLimit) { ASSERT_EQ(cudaStreamDestroy(stream), cudaSuccess); } -TEST_F(CudaAllocatorTest, AllocatesFromItsOwnPool) { +TEST_F(CudaAllocatorPoolTest, AllocatesFromItsOwnPool) { cudaStream_t stream; ASSERT_EQ(cudaStreamCreate(&stream), cudaSuccess); @@ -262,7 +333,7 @@ TEST_F(CudaAllocatorTest, AllocatesFromItsOwnPool) { // Freed memory is kept so repeated allocation stays cheap, which means a plain // free no longer shrinks the pool. Without an explicit release a long lived // process would hold that memory after every program was gone. -TEST_F(CudaAllocatorTest, ReleaseCachedMemoryReturnsPoolMemory) { +TEST_F(CudaAllocatorPoolTest, ReleaseCachedMemoryReturnsPoolMemory) { cudaStream_t stream; ASSERT_EQ(cudaStreamCreate(&stream), cudaSuccess); @@ -287,7 +358,7 @@ TEST_F(CudaAllocatorTest, ReleaseCachedMemoryReturnsPoolMemory) { } // Releasing must not disturb allocations that are still in use. -TEST_F(CudaAllocatorTest, ReleaseCachedMemoryKeepsLiveAllocations) { +TEST_F(CudaAllocatorPoolTest, ReleaseCachedMemoryKeepsLiveAllocations) { cudaStream_t stream; ASSERT_EQ(cudaStreamCreate(&stream), cudaSuccess); @@ -326,7 +397,7 @@ TEST_F(CudaAllocatorTest, ReleaseCachedMemoryKeepsLiveAllocations) { // current one. This runner has a single GPU, so the two cannot be told apart // here; what it pins is that the sentinel is resolved rather than passed to the // driver. -TEST_F(CudaAllocatorTest, ReleaseCachedMemoryAcceptsTheAllDevicesSentinel) { +TEST_F(CudaAllocatorPoolTest, ReleaseCachedMemoryAcceptsTheAllDevicesSentinel) { cudaStream_t stream; ASSERT_EQ(cudaStreamCreate(&stream), cudaSuccess); @@ -351,7 +422,7 @@ TEST_F(CudaAllocatorTest, ReleaseCachedMemoryAcceptsTheAllDevicesSentinel) { // Memory allocated during a graph capture goes to the device graph pool, which // the pool trim cannot reach, so releasing has to trim that too. Without the // graph trim this is the only new test that fails. -TEST_F(CudaAllocatorTest, ReleaseCachedMemoryReturnsGraphMemory) { +TEST_F(CudaAllocatorPoolTest, ReleaseCachedMemoryReturnsGraphMemory) { cudaStream_t stream; ASSERT_EQ(cudaStreamCreate(&stream), cudaSuccess); diff --git a/tools/cmake/Test.cmake b/tools/cmake/Test.cmake index 652f7df5589..396da740980 100644 --- a/tools/cmake/Test.cmake +++ b/tools/cmake/Test.cmake @@ -52,4 +52,19 @@ function(et_cxx_test target_name) # add_test adds a test target to be used by ctest add_test(NAME ${target_name} COMMAND ${target_name}) + # Windows records no search path in an executable, so a test linking a shared + # library built in another directory cannot start. Copy the DLLs it links + # beside it. The command is dropped when the test links none, since + # copy_if_different given only a destination fails. + if(WIN32) + set(_dlls "$") + add_custom_command( + TARGET ${target_name} + POST_BUILD + COMMAND + "$<$:${CMAKE_COMMAND};-E;copy_if_different;${_dlls};$>" + COMMAND_EXPAND_LISTS + ) + endif() + endfunction()