From 2ab27cd80e093153b103296588ef6d3cc9b705dc Mon Sep 17 00:00:00 2001 From: Songhao Jia Date: Mon, 5 Oct 2026 00:42:35 -0700 Subject: [PATCH] Fall back to cudaMalloc on CUDA devices without memory pools The CUDA backend's stream-ordered allocations fail on any device that does not support memory pools, which includes data-center GPUs in TCC mode on Windows (A10G, A100, H100 and the like). There cudaMemPoolCreate and cudaMallocAsync return cudaErrorNotSupported, and nothing falls back: E cuda_allocator.cpp:99] cudaMemPoolCreate failed for device 0: operation not supported. Using the default pool. E cuda_allocator.cpp:448] cudaMallocAsync failed: operation not supported (requested 512 bytes on device 0) F storage.h:126] In function allocate(), assert failed (result.ok()): CudaAllocator::allocate_async failed for 512 bytes on device 0 That is from a CUDA program running on the Windows GPU CI runner (an A10G) in #23123, the first CI job to run the delegate on Windows with a driver installed. A GeForce card in WDDM mode supports memory pools, which is why a local run passes. ## Change - CudaAllocator asks each device once whether it supports memory pools (cudaDevAttrMemoryPoolsSupported). On a device that does not, allocate_async allocates with cudaMalloc and deallocate_async frees with cudaFree, deciding by the device the pointer lives on. cudaFree waits for the device's outstanding work, so work still using the block on the stream finishes first. A failed attribute query keeps the stream-ordered path, so nothing changes where the query cannot answer. - release_cached_memory skips cudaDeviceGraphMemTrim on such a device; graph memory is served by memory pools, so there is none to release. - sort.cu and rand.cu called cudaMallocAsync/cudaFreeAsync directly; they now go through CudaAllocator, so they take the same fallback. Devices with memory pools, which is every Linux runner and every WDDM GPU, take exactly the path they took before. ## Tests and CI - test_cuda_allocator: AllocateAsyncRoundtrip runs on every device and fails without the fallback on a device without pools. FallsBackWithoutMemoryPools checks that no pool is created and no error is left behind there. The pool tests move to a fixture that skips on a device without pools, where they cannot mean anything. - cuda-windows.yml gains unittest-cuda-runtime-windows, the Windows counterpart of cuda.yml's unittest-cuda-runtime: it builds llm-release-cuda with tests and runs the same five runtime tests. The Windows runner's A10G has no memory pools, so this covers the fallback; the Linux job covers the pool path. - The Windows GPU runners come without an NVIDIA driver (nvcuda.dll is missing and cudaGetDeviceCount returns cudaErrorInsufficientDriver), so the job first installs the data-center driver that pytorch/pytorch's Windows CUDA smoke tests install (driver_update.bat, 580.88, from the same bucket); install_nvidia_driver_windows.ps1. - et_cxx_test copies a test's runtime DLLs ($) beside it on Windows, which records no search path: the CUDA tests link extension_cuda.dll from another directory and could not start otherwise. ## Test plan - Windows 11, RTX 5080 (WDDM, memory pools supported), CUDA 13.0, llm-release-cuda with tests: all five runtime tests pass (test_cuda_allocator 16 passed, the no-pool test skipped as expected; mutable_state 15, weight_cache 11, guard 7, stream_guard 21). - The no-pool path runs in the new Windows CI job on the A10G. - Linux: the pool path is unchanged and covered by the existing unittest-cuda-runtime job. Found while adding the Windows CUDA wheel (#23123), which needs this to run on the Windows GPU runner; this PR does not depend on it. --- .ci/scripts/install_nvidia_driver_windows.ps1 | 28 +++++++ .ci/scripts/test_cuda_runtime_windows.ps1 | 44 ++++++++++ .github/workflows/cuda-windows.yml | 34 ++++++++ backends/cuda/runtime/cuda_allocator.cpp | 62 +++++++++++++- backends/cuda/runtime/shims/rand.cu | 10 ++- backends/cuda/runtime/shims/sort.cu | 41 +++++---- .../cuda/runtime/test/test_cuda_allocator.cpp | 83 +++++++++++++++++-- tools/cmake/Test.cmake | 15 ++++ 8 files changed, 293 insertions(+), 24 deletions(-) create mode 100644 .ci/scripts/install_nvidia_driver_windows.ps1 create mode 100644 .ci/scripts/test_cuda_runtime_windows.ps1 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()