如何将大于 VRAM 大小的数据传递到 GPU？
How to pass data bigger than the VRAM size into the GPU?
我试图将比 VRAM 更多的数据传递到我的 GPU，这导致了以下错误。 CudaAPIError: Call to cuMemAlloc results in CUDA_ERROR_OUT_OF_MEMORY
我创建了这段代码来重现问题：
from numba import cuda
import numpy as np
@cuda.jit()
def addingNumbers (big_array, big_array2, save_array):
i = cuda.grid(1)
if i < big_array.shape[0]:
for j in range (big_array.shape[1]):
save_array[i][j] = big_array[i][j] * big_array2[i][j]
big_array = np.random.random_sample((1000000, 500))
big_array2 = np.random.random_sample((1000000, 500))
save_array = np.zeros(shape=(1000000, 500))
arraysize = 1000000
threadsperblock = 64
blockspergrid = (arraysize + (threadsperblock - 1))
d_big_array = cuda.to_device(big_array)
d_big_array2 = cuda.to_device(big_array2)
d_save_array = cuda.to_device(save_array)
addingNumbers[blockspergrid, threadsperblock](d_big_array, d_big_array2, d_save_array)
save_array = d_save_array.copy_to_host()
有没有办法将数据动态传递到 GPU 以处理比 VRAM 能容纳的更多的数据？如果没有，将所有这些数据手动传递给 gpu 的推荐方法是什么。使用 dask_cuda 是一个选项，还是类似的东西？
一个 well-written 如何处理一个更大的问题（即数据集）并将其分解，并在 numba CUDA 中处理处理 piece-wise 的示例是 here。特别是，感兴趣的变体是 pricer_cuda_overlap.py。不幸的是，该示例使用了我认为在 accelerate.cuda.rand 中已弃用的 运行dom 数字生成功能，因此在今天的 numba 中它不能直接 运行nable（我认为）。
然而，就这里的问题而言，运行dom 编号生成过程是无关紧要的，因此我们可以简单地删除它而不影响重要的观察结果。在该示例中，接下来是一个由多个文件中的多个部分组合而成的单个文件：
$ cat t45.py
#! /usr/bin/env python
"""
This version demonstrates copy-compute overlapping through multiple streams.
"""
from __future__ import print_function
import math
import sys
import numpy as np
from numba import cuda, jit
from math import sqrt, exp
from timeit import default_timer as timer
from collections import deque
StockPrice = 20.83
StrikePrice = 21.50
Volatility = 0.021 # per year
InterestRate = 0.20
Maturity = 5. / 12.
NumPath = 500000
NumStep = 200
def driver(pricer, pinned=False):
paths = np.zeros((NumPath, NumStep + 1), order='F')
paths[:, 0] = StockPrice
DT = Maturity / NumStep
if pinned:
from numba import cuda
with cuda.pinned(paths):
ts = timer()
pricer(paths, DT, InterestRate, Volatility)
te = timer()
else:
ts = timer()
pricer(paths, DT, InterestRate, Volatility)
te = timer()
ST = paths[:, -1]
PaidOff = np.maximum(paths[:, -1] - StrikePrice, 0)
print('Result')
fmt = '%20s: %s'
print(fmt % ('stock price', np.mean(ST)))
print(fmt % ('standard error', np.std(ST) / sqrt(NumPath)))
print(fmt % ('paid off', np.mean(PaidOff)))
optionprice = np.mean(PaidOff) * exp(-InterestRate * Maturity)
print(fmt % ('option price', optionprice))
print('Performance')
NumCompute = NumPath * NumStep
print(fmt % ('Mstep/second', '%.2f' % (NumCompute / (te - ts) / 1e6)))
print(fmt % ('time elapsed', '%.3fs' % (te - ts)))
class MM(object):
"""Memory Manager
Maintain a freelist of device memory for reuse.
"""
def __init__(self, shape, dtype, prealloc):
self.device = cuda.get_current_device()
self.freelist = deque()
self.events = {}
for i in range(prealloc):
gpumem = cuda.device_array(shape=shape, dtype=dtype)
self.freelist.append(gpumem)
self.events[gpumem] = cuda.event(timing=False)
def get(self, stream=0):
assert self.freelist
gpumem = self.freelist.popleft()
evnt = self.events[gpumem]
if not evnt.query(): # not ready?
# querying is faster then waiting
evnt.wait(stream=stream) # future works must wait
return gpumem
def free(self, gpumem, stream=0):
evnt = self.events[gpumem]
evnt.record(stream=stream)
self.freelist.append(gpumem)
if sys.version_info[0] == 2:
range = xrange
@jit('void(double[:], double[:], double, double, double, double[:])',
target='cuda')
def cu_step(last, paths, dt, c0, c1, normdist):
i = cuda.grid(1)
if i >= paths.shape[0]:
return
noise = normdist[i]
paths[i] = last[i] * math.exp(c0 * dt + c1 * noise)
def monte_carlo_pricer(paths, dt, interest, volatility):
n = paths.shape[0]
num_streams = 2
part_width = int(math.ceil(float(n) / num_streams))
partitions = [(0, part_width)]
for i in range(1, num_streams):
begin, end = partitions[i - 1]
begin, end = end, min(end + (end - begin), n)
partitions.append((begin, end))
partlens = [end - begin for begin, end in partitions]
mm = MM(shape=part_width, dtype=np.double, prealloc=10 * num_streams)
device = cuda.get_current_device()
blksz = device.MAX_THREADS_PER_BLOCK
gridszlist = [int(math.ceil(float(partlen) / blksz))
for partlen in partlens]
strmlist = [cuda.stream() for _ in range(num_streams)]
# Allocate device side array - in original example this would be initialized with random numbers
d_normlist = [cuda.device_array(partlen, dtype=np.double, stream=strm)
for partlen, strm in zip(partlens, strmlist)]
c0 = interest - 0.5 * volatility ** 2
c1 = volatility * math.sqrt(dt)
# Configure the kernel
# Similar to CUDA-C: cu_monte_carlo_pricer<<<gridsz, blksz, 0, stream>>>
steplist = [cu_step[gridsz, blksz, strm]
for gridsz, strm in zip(gridszlist, strmlist)]
d_lastlist = [cuda.to_device(paths[s:e, 0], to=mm.get(stream=strm))
for (s, e), strm in zip(partitions, strmlist)]
for j in range(1, paths.shape[1]):
d_pathslist = [cuda.to_device(paths[s:e, j], stream=strm,
to=mm.get(stream=strm))
for (s, e), strm in zip(partitions, strmlist)]
for step, args in zip(steplist, zip(d_lastlist, d_pathslist, d_normlist)):
d_last, d_paths, d_norm = args
step(d_last, d_paths, dt, c0, c1, d_norm)
for d_paths, strm, (s, e) in zip(d_pathslist, strmlist, partitions):
d_paths.copy_to_host(paths[s:e, j], stream=strm)
mm.free(d_paths, stream=strm)
d_lastlist = d_pathslist
for strm in strmlist:
strm.synchronize()
if __name__ == '__main__':
driver(monte_carlo_pricer, pinned=True)
$ python t45.py
Result
stock price: 22.6720614385
standard error: 0.0
paid off: 1.17206143849
option price: 1.07834858009
Performance
Mstep/second: 336.40
time elapsed: 0.297s
$
这个例子中有很多内容，关于如何在 CUDA 中编写 pipelined/overlapped 代码的一般主题本身就是一个完整的答案，所以我只介绍要点。 this blog post 很好地涵盖了一般主题，尽管考虑的是 CUDA C++，而不是 numba CUDA (python)。然而，numba CUDA 中大多数感兴趣的项目与其在 CUDA C++ 中的对应项之间存在 1:1 对应关系。因此，我将假定已理解 CUDA 流等基本概念，以及它们如何用于 ar运行ge 异步并发 activity。
那么这个例子是做什么的呢？我将主要关注 CUDA 方面。
- 考虑到复制和计算操作的重叠，输入数据 (
paths) 被转换为主机上的 CUDA 固定内存
- 为了以块的形式处理工作，定义了内存管理器 (
MM)，这将允许在处理过程中重复使用设备内存的块分配。
- python 列表被创建来表示块处理的顺序。有一个列表定义每个块或分区的开始和结束。有一个列表定义了要使用的 cuda 流的顺序。有一个 CUDA 内核将使用的数据数组分区列表。
- 然后，有了这些列表，"depth-first-order" 就有了作品的发布。对于每个流，该流所需的数据（块）在 t运行sferred 到设备（排队等待 t运行sfer），将处理该数据的内核启动（排队），并且t运行sfer 会将结果从该块发送回主机内存，它已排队。此过程在
monte_carlo_pricer 中的 for j 循环中重复，步骤数为 (paths.shape[1])。
当我 运行 使用分析器执行上述代码时，我们可以看到如下所示的时间线：
在这种特殊情况下，我 运行 在 Quadro K2000 上进行此操作，这是一种旧的小型 GPU，只有一个复制引擎。因此我们在配置文件中看到最多有 1 个复制操作与 CUDA 内核 activity 重叠，并且没有复制操作与其他复制操作重叠。但是，如果我 运行 在具有 2 个复制引擎的设备上执行此操作，我希望 tighter/denser 时间线是可能的，同时重叠 2 个复制操作和计算操作，以获得最大吞吐量。为此，使用中的流 (num_streams) 也必须至少增加到 3.
这里的代码不保证运行没有缺陷。提供它用于演示目的。使用它需要您自担风险。
我试图将比 VRAM 更多的数据传递到我的 GPU，这导致了以下错误。 CudaAPIError: Call to cuMemAlloc results in CUDA_ERROR_OUT_OF_MEMORY
我创建了这段代码来重现问题：
from numba import cuda
import numpy as np
@cuda.jit()
def addingNumbers (big_array, big_array2, save_array):
i = cuda.grid(1)
if i < big_array.shape[0]:
for j in range (big_array.shape[1]):
save_array[i][j] = big_array[i][j] * big_array2[i][j]
big_array = np.random.random_sample((1000000, 500))
big_array2 = np.random.random_sample((1000000, 500))
save_array = np.zeros(shape=(1000000, 500))
arraysize = 1000000
threadsperblock = 64
blockspergrid = (arraysize + (threadsperblock - 1))
d_big_array = cuda.to_device(big_array)
d_big_array2 = cuda.to_device(big_array2)
d_save_array = cuda.to_device(save_array)
addingNumbers[blockspergrid, threadsperblock](d_big_array, d_big_array2, d_save_array)
save_array = d_save_array.copy_to_host()
有没有办法将数据动态传递到 GPU 以处理比 VRAM 能容纳的更多的数据？如果没有，将所有这些数据手动传递给 gpu 的推荐方法是什么。使用 dask_cuda 是一个选项，还是类似的东西？
一个 well-written 如何处理一个更大的问题（即数据集）并将其分解，并在 numba CUDA 中处理处理 piece-wise 的示例是 here。特别是，感兴趣的变体是 pricer_cuda_overlap.py。不幸的是，该示例使用了我认为在 accelerate.cuda.rand 中已弃用的 运行dom 数字生成功能，因此在今天的 numba 中它不能直接 运行nable（我认为）。
然而，就这里的问题而言，运行dom 编号生成过程是无关紧要的，因此我们可以简单地删除它而不影响重要的观察结果。在该示例中，接下来是一个由多个文件中的多个部分组合而成的单个文件：
$ cat t45.py
#! /usr/bin/env python
"""
This version demonstrates copy-compute overlapping through multiple streams.
"""
from __future__ import print_function
import math
import sys
import numpy as np
from numba import cuda, jit
from math import sqrt, exp
from timeit import default_timer as timer
from collections import deque
StockPrice = 20.83
StrikePrice = 21.50
Volatility = 0.021 # per year
InterestRate = 0.20
Maturity = 5. / 12.
NumPath = 500000
NumStep = 200
def driver(pricer, pinned=False):
paths = np.zeros((NumPath, NumStep + 1), order='F')
paths[:, 0] = StockPrice
DT = Maturity / NumStep
if pinned:
from numba import cuda
with cuda.pinned(paths):
ts = timer()
pricer(paths, DT, InterestRate, Volatility)
te = timer()
else:
ts = timer()
pricer(paths, DT, InterestRate, Volatility)
te = timer()
ST = paths[:, -1]
PaidOff = np.maximum(paths[:, -1] - StrikePrice, 0)
print('Result')
fmt = '%20s: %s'
print(fmt % ('stock price', np.mean(ST)))
print(fmt % ('standard error', np.std(ST) / sqrt(NumPath)))
print(fmt % ('paid off', np.mean(PaidOff)))
optionprice = np.mean(PaidOff) * exp(-InterestRate * Maturity)
print(fmt % ('option price', optionprice))
print('Performance')
NumCompute = NumPath * NumStep
print(fmt % ('Mstep/second', '%.2f' % (NumCompute / (te - ts) / 1e6)))
print(fmt % ('time elapsed', '%.3fs' % (te - ts)))
class MM(object):
"""Memory Manager
Maintain a freelist of device memory for reuse.
"""
def __init__(self, shape, dtype, prealloc):
self.device = cuda.get_current_device()
self.freelist = deque()
self.events = {}
for i in range(prealloc):
gpumem = cuda.device_array(shape=shape, dtype=dtype)
self.freelist.append(gpumem)
self.events[gpumem] = cuda.event(timing=False)
def get(self, stream=0):
assert self.freelist
gpumem = self.freelist.popleft()
evnt = self.events[gpumem]
if not evnt.query(): # not ready?
# querying is faster then waiting
evnt.wait(stream=stream) # future works must wait
return gpumem
def free(self, gpumem, stream=0):
evnt = self.events[gpumem]
evnt.record(stream=stream)
self.freelist.append(gpumem)
if sys.version_info[0] == 2:
range = xrange
@jit('void(double[:], double[:], double, double, double, double[:])',
target='cuda')
def cu_step(last, paths, dt, c0, c1, normdist):
i = cuda.grid(1)
if i >= paths.shape[0]:
return
noise = normdist[i]
paths[i] = last[i] * math.exp(c0 * dt + c1 * noise)
def monte_carlo_pricer(paths, dt, interest, volatility):
n = paths.shape[0]
num_streams = 2
part_width = int(math.ceil(float(n) / num_streams))
partitions = [(0, part_width)]
for i in range(1, num_streams):
begin, end = partitions[i - 1]
begin, end = end, min(end + (end - begin), n)
partitions.append((begin, end))
partlens = [end - begin for begin, end in partitions]
mm = MM(shape=part_width, dtype=np.double, prealloc=10 * num_streams)
device = cuda.get_current_device()
blksz = device.MAX_THREADS_PER_BLOCK
gridszlist = [int(math.ceil(float(partlen) / blksz))
for partlen in partlens]
strmlist = [cuda.stream() for _ in range(num_streams)]
# Allocate device side array - in original example this would be initialized with random numbers
d_normlist = [cuda.device_array(partlen, dtype=np.double, stream=strm)
for partlen, strm in zip(partlens, strmlist)]
c0 = interest - 0.5 * volatility ** 2
c1 = volatility * math.sqrt(dt)
# Configure the kernel
# Similar to CUDA-C: cu_monte_carlo_pricer<<<gridsz, blksz, 0, stream>>>
steplist = [cu_step[gridsz, blksz, strm]
for gridsz, strm in zip(gridszlist, strmlist)]
d_lastlist = [cuda.to_device(paths[s:e, 0], to=mm.get(stream=strm))
for (s, e), strm in zip(partitions, strmlist)]
for j in range(1, paths.shape[1]):
d_pathslist = [cuda.to_device(paths[s:e, j], stream=strm,
to=mm.get(stream=strm))
for (s, e), strm in zip(partitions, strmlist)]
for step, args in zip(steplist, zip(d_lastlist, d_pathslist, d_normlist)):
d_last, d_paths, d_norm = args
step(d_last, d_paths, dt, c0, c1, d_norm)
for d_paths, strm, (s, e) in zip(d_pathslist, strmlist, partitions):
d_paths.copy_to_host(paths[s:e, j], stream=strm)
mm.free(d_paths, stream=strm)
d_lastlist = d_pathslist
for strm in strmlist:
strm.synchronize()
if __name__ == '__main__':
driver(monte_carlo_pricer, pinned=True)
$ python t45.py
Result
stock price: 22.6720614385
standard error: 0.0
paid off: 1.17206143849
option price: 1.07834858009
Performance
Mstep/second: 336.40
time elapsed: 0.297s
$
这个例子中有很多内容，关于如何在 CUDA 中编写 pipelined/overlapped 代码的一般主题本身就是一个完整的答案，所以我只介绍要点。 this blog post 很好地涵盖了一般主题，尽管考虑的是 CUDA C++，而不是 numba CUDA (python)。然而，numba CUDA 中大多数感兴趣的项目与其在 CUDA C++ 中的对应项之间存在 1:1 对应关系。因此，我将假定已理解 CUDA 流等基本概念，以及它们如何用于 ar运行ge 异步并发 activity。
那么这个例子是做什么的呢？我将主要关注 CUDA 方面。
- 考虑到复制和计算操作的重叠，输入数据 (
paths) 被转换为主机上的 CUDA 固定内存 - 为了以块的形式处理工作，定义了内存管理器 (
MM)，这将允许在处理过程中重复使用设备内存的块分配。 - python 列表被创建来表示块处理的顺序。有一个列表定义每个块或分区的开始和结束。有一个列表定义了要使用的 cuda 流的顺序。有一个 CUDA 内核将使用的数据数组分区列表。
- 然后，有了这些列表，"depth-first-order" 就有了作品的发布。对于每个流，该流所需的数据（块）在 t运行sferred 到设备（排队等待 t运行sfer），将处理该数据的内核启动（排队），并且t运行sfer 会将结果从该块发送回主机内存，它已排队。此过程在
monte_carlo_pricer中的for j循环中重复，步骤数为 (paths.shape[1])。
当我 运行 使用分析器执行上述代码时，我们可以看到如下所示的时间线：
在这种特殊情况下，我 运行 在 Quadro K2000 上进行此操作，这是一种旧的小型 GPU，只有一个复制引擎。因此我们在配置文件中看到最多有 1 个复制操作与 CUDA 内核 activity 重叠，并且没有复制操作与其他复制操作重叠。但是，如果我 运行 在具有 2 个复制引擎的设备上执行此操作，我希望 tighter/denser 时间线是可能的，同时重叠 2 个复制操作和计算操作，以获得最大吞吐量。为此，使用中的流 (num_streams) 也必须至少增加到 3.
这里的代码不保证运行没有缺陷。提供它用于演示目的。使用它需要您自担风险。