WSL2 から cudaMallocManaged を呼ぶと確保できるメモリ量が少なくなる

質問内容

cudaMallocManaged について、Windows側で呼んだときとWSL2上で呼んだときとで確保できるメモリ量が異なっていました。
WSL2 (Ubuntu20.04) 上で CUDA を使っているときに、メモリ確保に失敗して気がつきました。
GPUのメモリ 4.0GiB に対して、Windows側の呼び出しでは 3,729,723,392 バイト (約 3557 MiB) まで確保できましたが、WSL2側の呼び出しでは 1,231,290,368 バイト (1174.25 MiB) しか確保できませんでした。
これはなぜでしょうか？

実現したいこと

理由が知りたいです。
また、可能であればWSL2上でもWindows側と同じだけメモリを確保できるようにしたいです。

発生している問題・エラーメッセージ

下のコードで確保できたメモリ量です。

# GeForce GTX 1650 (4GB VRAM) のとき
## Windows側での実行結果
succeeded size = 0x0000_0000_de4f_1000 (3729723392)
## WSL2側での実行結果
succeeded size = 0x0000_0000_4964_0000 (1231290368)

# GeForce RTX 2080 SUPER (8GB VRAM) のとき
## Windows側での実行結果
succeeded size = 0x0000_0001_d35d_1000 (7841058816)
## WSL2側での実行結果
succeeded size = 0x0000_0000_49c3_0000 (1237516288)

確保量が1/3くらいになっています。
【9/25編集　GPUを変更して試した結果を追加】
VRAMが増えてもWSL2側で確保できるメモリ量は変わっていませんでした。
cudaMalloc ではVRAMの限界まで確保できています。

ソースコード

c++
1// memory.cu
2
3#include <cstdio>
4#include <cstdlib>
5#include <cstdint>
6
7// メモリ確保用のFunctor。
8// operator()(size_t) でメモリの確保を試みて成功したかどうかを返す。
9// それぞれメモリ確保の方法が異なる。
10//   alloc_m: malloc
11//   alloc_cuda: cudaMalloc
12//   alloc_cuda_managed_gloabl: cudaMallocManaged(ptr, size, cudaMemAttachGlobal)
13//   alloc_cuda_managed_host: cudaMallocManaged(ptr, size, cudaMemAttachHost)
14
15struct alloc_m {
16    __device__
17    bool operator()(size_t size) {
18        auto p = std::malloc(size);
19        free(p);
20        return p;
21    }
22};
23
24struct alloc_cuda {
25    __host__
26    bool operator()(size_t size) {
27        void *p;
28        auto e = cudaMalloc(&p, size);
29        if (e == cudaSuccess) {
30            cudaFree(p);
31            return true;
32        } else {
33            return false;
34        }
35    }
36};
37
38struct alloc_cuda_managed_gloabl {
39    __host__
40    bool operator()(size_t size) {
41        void *p;
42        auto e = cudaMallocManaged(&p, size, cudaMemAttachGlobal);
43        if (e == cudaSuccess) {
44            cudaFree(p);
45            return true;
46        } else {
47            return false;
48        }
49    }
50};
51
52struct alloc_cuda_managed_host {
53    __host__
54    bool operator()(size_t size) {
55        void *p;
56        auto e = cudaMallocManaged(&p, size, cudaMemAttachHost);
57        if (e == cudaSuccess) {
58            cudaFree(p);
59            return true;
60        } else {
61            return false;
62        }
63    }
64};
65
66// 確保できるメモリの上限を二分探索で探す。
67// ホストとカーネルで同じ内容。
68
69template <typename Allocator>
70__host__ size_t search_host(size_t min, size_t max) {
71    size_t pivot = (max + min) / 2;
72    if (pivot == min || pivot == max) return pivot;
73    
74    if (Allocator()(pivot)) {
75        return search_host<Allocator>(pivot, max);
76    } else {
77        return search_host<Allocator>(min, pivot);
78    }
79}
80
81template <typename Allocator>
82__device__ size_t search_kernel(size_t min, size_t max) {
83    size_t pivot = (max + min) / 2;
84    if (pivot == min || pivot == max) return pivot;
85    
86    if (Allocator()(pivot)) {
87        return search_kernel<Allocator>(pivot, max);
88    } else {
89        return search_kernel<Allocator>(min, pivot);
90    }
91}
92
93// 確保できるメモリの上限を探す。
94// まず確保に失敗する量を探してそこを上限とし、search_* で二分探索を行う。
95// ホストとカーネルでほぼ同じ内容。nvcc の仕様により printf のフォーマット文字列だけ変えている。
96
97template <typename Allocator>
98__host__ void do_check_host() {
99    // 失敗するまで2倍ずつメモリ確保量を増やしていく
100    size_t size = 1 << 10; // 1 KiB
101    while (true) {
102        if (!Allocator()(size)) {
103            size_t x1 = size & 0xffff;
104            size_t x2 = (size >> 16) & 0xffff;
105            size_t x3 = (size >> 32) & 0xffff;
106            size_t x4 = (size >> 48) & 0xffff;
107            printf("failed size = 0x%04zx_%04zx_%04zx_%04zx (%zd)\n", x4, x3, x2, x1, size);
108            break;
109        }
110        size <<= 1;
111    }
112
113    // now `size` indicate upper limit of memory allocation
114    // `size` を上限として、二分探索で確保可能なメモリ量を探す
115    size_t avail = search_host<Allocator>(size / 2, size);
116    {
117        size_t x1 = avail & 0xffff;
118        size_t x2 = (avail >> 16) & 0xffff;
119        size_t x3 = (avail >> 32) & 0xffff;
120        size_t x4 = (avail >> 48) & 0xffff;
121        printf("succeeded size = 0x%04zx_%04zx_%04zx_%04zx (%zd)\n", x4, x3, x2, x1, avail);
122    }
123}
124
125template <typename Allocator>
126__global__ void do_check_kernel() {
127    size_t size = 1 << 10; // 1 KiB
128    while (true) {
129        if (!Allocator()(size)) {
130            uint32_t x1 = size & 0xffff;
131            uint32_t x2 = (size >> 16) & 0xffff;
132            uint32_t x3 = (size >> 32) & 0xffff;
133            uint32_t x4 = (size >> 48) & 0xffff;
134            printf("failed size = 0x%04x_%04x_%04x_%04x (%llu)\n", x4, x3, x2, x1, (uint64_t)size);
135            break;
136        }
137        size <<= 1;
138    }
139
140    // now `size` indicate upper limit of memory allocation
141    size_t avail = search_kernel<Allocator>(size / 2, size);
142    {
143        uint32_t x1 = avail & 0xffff;
144        uint32_t x2 = (avail >> 16) & 0xffff;
145        uint32_t x3 = (avail >> 32) & 0xffff;
146        uint32_t x4 = (avail >> 48) & 0xffff;
147        printf("succeeded size = %04x_%04x_%04x_%04x (%llu)\n", x4, x3, x2, x1, (uint64_t)avail);
148    }
149}
150
151int main() {
152    printf("=== HOST =============================================\n");
153    printf("* cudaMalloc\n");
154    do_check_host<alloc_cuda>();
155    printf("* cudaMallocManaged(global)\n");
156    do_check_host<alloc_cuda_managed_gloabl>();
157    printf("* cudaMallocManaged(host)\n");
158    do_check_host<alloc_cuda_managed_host>();
159    
160    printf("=== KERNEL ===========================================\n");
161    printf("* malloc\n");
162    do_check_kernel<alloc_m><<<1,1>>>();
163    cudaDeviceSynchronize();
164    
165    return 0;
166}

コンパイルと実行

$ nvcc memory.cu -o memory && ./memory

補足情報（FW/ツールのバージョンなど）

Windows側の情報は以下の通りです。

OS: Windows 10 Home 21H2 (build 19044.2006) [x86_64]
CPU: Intel Core i5 4670
メモリ: 16GB
GPU: GeForce GTX 1650
VRAM: 4GB
$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2021 NVIDIA Corporation
Built on Fri_Dec_17_18:28:54_Pacific_Standard_Time_2021
Cuda compilation tools, release 11.6, V11.6.55
Build cuda_11.6.r11.6/compiler.30794723_0

WSL2側の情報は以下の通りです。

$ uname -r
5.10.102.1-microsoft-standard-WSL2
$ cat /etc/lsb-release
DISTRIB_ID=Ubuntu
DISTRIB_RELEASE=20.04
DISTRIB_CODENAME=focal
DISTRIB_DESCRIPTION="Ubuntu 20.04.5 LTS"
$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2022 NVIDIA Corporation
Built on Tue_Mar__8_18:18:20_PST_2022
Cuda compilation tools, release 11.6, V11.6.124
Build cuda_11.6.r11.6/compiler.31057947_0

Windows側で nvidia-smi を実行した結果は以下の通りでした。
WSL2側でもドライバのバージョン、メモリ等は同じだったので、そちらは省略します。

$ nvidia-smi
+-----------------------------------------------------------------------------+
| NVIDIA-SMI 512.15       Driver Version: 512.15       CUDA Version: 11.6     |
|-------------------------------+----------------------+----------------------+
| GPU  Name            TCC/WDDM | Bus-Id        Disp.A | Volatile Uncorr. ECC |
| Fan  Temp  Perf  Pwr:Usage/Cap|         Memory-Usage | GPU-Util  Compute M. |
|                               |                      |               MIG M. |
|===============================+======================+======================|
|   0  NVIDIA GeForce ... WDDM  | 00000000:01:00.0  On |                  N/A |
|  0%   54C    P8    N/A /  75W |    451MiB /  4096MiB |      4%      Default |
|                               |                      |                  N/A |
+-------------------------------+----------------------+----------------------+
(...snip...)