CUD Stream

• https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#c-language-extensions
  中指出在kenel的調(diào)用函數(shù)中最后一個(gè)可選參數(shù)表示該核函數(shù)處在哪個(gè)流之中。
  

- 參數(shù)Dg用于定義整個(gè)grid的維度和尺寸，即一個(gè)grid有多少個(gè)block。為dim3類型。Dim3 Dg(Dg.x, Dg.y, 1)表示grid中每行有Dg.x個(gè)block，每列有Dg.y個(gè)block，第三維恒為1(目前一個(gè)核函數(shù)只有一個(gè)grid)。整個(gè)grid中共有Dg.x*Dg.y個(gè)block，其中Dg.x和Dg.y最大值為65535。
- 參數(shù)Db用于定義一個(gè)block的維度和尺寸，即一個(gè)block有多少個(gè)thread。為dim3類型。Dim3 Db(Db.x, Db.y, Db.z)表示整個(gè)block中每行有Db.x個(gè)thread，每列有Db.y個(gè)thread，高度為Db.z。Db.x和Db.y最大值為512，Db.z最大值為62。 一個(gè)block中共有Db.x*Db.y*Db.z個(gè)thread。計(jì)算能力為1.0,1.1的硬件該乘積的最大值為768，計(jì)算能力為1.2,1.3的硬件支持的最大值為1024。
- Ns 的類型為 size_t，用于設(shè)置每個(gè)block除了靜態(tài)分配的shared Memory以外，最多能動(dòng)態(tài)分配的shared memory大小，單位為byte。不需要?jiǎng)討B(tài)分配時(shí)該值為0或省略不寫。如[__shared__](https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#shared)中所述，此動(dòng)態(tài)分配的內(nèi)存由聲明為外部數(shù)組的任何變量使用; 
- 參數(shù)S是一個(gè)cudaStream_t類型的可選參數(shù)，初始值為零，表示該核函數(shù)處在哪個(gè)流之中。

• CUDA編程中，默認(rèn)使用默認(rèn)流非并行執(zhí)行kernel，每個(gè)kernel由許多thread并行的執(zhí)行在GPU上。Stream的概念是相對(duì)Grid level來說的，使得kernel在一個(gè)device上同時(shí)執(zhí)行。
  

• 官方提供的用例

// https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#streams
cudaStream_t stream[2];
for (int i = 0; i < 2; ++i)cudaStreamCreate(&stream[i]);
float* hostPtr;
cudaMallocHost(&hostPtr, 2 * size);
// 以下代碼示例將其中每個(gè)流定義為從主機(jī)到設(shè)備的一個(gè)內(nèi)存副本、一個(gè)內(nèi)核啟動(dòng)和一個(gè)從設(shè)備到主機(jī)的內(nèi)存副本的序列：
for (int i = 0; i < 2; ++i) {cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size,size, cudaMemcpyHostToDevice, stream[i]);MyKernel <<<100, 512, 0, stream[i]>>>(outputDevPtr + i * size, inputDevPtr + i * size, size);cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size,size, cudaMemcpyDeviceToHost, stream[i]);
}
// 通過調(diào)用 釋放流
for (int i = 0; i < 2; ++i)cudaStreamDestroy(stream[i]);

PyTorch Stream

• 在PyTorch中，默認(rèn)情況下，GPU上的操作是在默認(rèn)流（default stream）中執(zhí)行的。默認(rèn)流是一個(gè)序列化的流，其中的操作按照它們出現(xiàn)的順序逐個(gè)執(zhí)行。這意味著在沒有顯式指定其他流的情況下，所有的操作都會(huì)在默認(rèn)流中執(zhí)行。

• 然而，PyTorch還提供了功能可以將操作提交到其他流中執(zhí)行，以充分利用GPU的并行性。這對(duì)于并行處理多個(gè)任務(wù)或同時(shí)執(zhí)行多個(gè)獨(dú)立操作非常有用。

• 您可以使用torch.cuda.Stream()來創(chuàng)建其他流，并使用torch.cuda.current_stream()來獲取當(dāng)前流。然后，您可以將操作提交到指定的流中執(zhí)行，例如：

import torchdevice = torch.device('cuda')# 創(chuàng)建一個(gè)默認(rèn)流
default_stream = torch.cuda.current_stream()# 創(chuàng)建一個(gè)自定義流
custom_stream = torch.cuda.Stream()# 在默認(rèn)流中執(zhí)行操作
with torch.cuda.stream(default_stream):# 執(zhí)行操作...# 在自定義流中執(zhí)行操作
with torch.cuda.stream(custom_stream):# 執(zhí)行操作...

例子

import torch
s1 = torch.cuda.Stream()
s2 = torch.cuda.Stream()
# Initialise cuda tensors here. E.g.:
A = torch.rand(1000, 1000, device = 'cuda')
B = torch.rand(1000, 1000, device = 'cuda')
# Wait for the above tensors to initialise.
torch.cuda.synchronize()
with torch.cuda.stream(s1):C = torch.mm(A, A)
with torch.cuda.stream(s2):D = torch.mm(B, B)
# Wait for C and D to be computed.
torch.cuda.synchronize()
# Do stuff with C and D.
print(C)
print(D)
// https://stackoverflow.com/questions/70128833/why-and-when-to-use-torch-cuda-stream

這樣可以利用多個(gè)流來并行執(zhí)行計(jì)算，并在計(jì)算和數(shù)據(jù)傳輸之間實(shí)現(xiàn)重疊。這對(duì)于提高GPU利用率和加速訓(xùn)練或推理過程非常有幫助。

錯(cuò)誤示例

• 沒有使用 synchronize() 或者 wait_stream()進(jìn)行同步，可能導(dǎo)致再未完成歸一化前執(zhí)行求和

// https://pytorch.org/docs/stable/notes/cuda.html
cuda = torch.device('cuda')
s = torch.cuda.Stream()  # Create a new stream.
A = torch.empty((100, 100), device=cuda).normal_(0.0, 1.0)
with torch.cuda.stream(s):# sum() may start execution before normal_() finishes!B = torch.sum(A)

CG

• https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#streams

• https://pytorch.org/docs/stable/notes/cuda.html#multistream-capture

• https://pytorch.org/cppdocs/notes/tensor_cuda_stream.html

• https://pypi.org/project/pytorch-stream/

• CUDA 的 Stream and Event https://zhuanlan.zhihu.com/p/369367933

• GITHUBGIST Gist就是小型代碼片段的分享https://www.cnblogs.com/leader755/p/14284716.html

• [JIT] 在 TorchScript 中支持 CUDA 流 https://github.com/pytorch/pytorch/issues/41355

• https://pytorch.org/docs/stable/notes/cuda.html#cuda-semantics

• https://github.com/pytorch/pytorch/issues/41355

多設(shè)備

// https://pytorch.org/docs/stable/notes/cuda.html#cuda-semantics
cuda = torch.device('cuda')     # Default CUDA device
cuda0 = torch.device('cuda:0')
cuda2 = torch.device('cuda:2')  # GPU 2 (these are 0-indexed)x = torch.tensor([1., 2.], device=cuda0)
# x.device is device(type='cuda', index=0)
y = torch.tensor([1., 2.]).cuda()
# y.device is device(type='cuda', index=0)with torch.cuda.device(1):# allocates a tensor on GPU 1a = torch.tensor([1., 2.], device=cuda)# transfers a tensor from CPU to GPU 1b = torch.tensor([1., 2.]).cuda()# a.device and b.device are device(type='cuda', index=1)# You can also use ``Tensor.to`` to transfer a tensor:b2 = torch.tensor([1., 2.]).to(device=cuda)# b.device and b2.device are device(type='cuda', index=1)c = a + b# c.device is device(type='cuda', index=1)z = x + y# z.device is device(type='cuda', index=0)# even within a context, you can specify the device# (or give a GPU index to the .cuda call)d = torch.randn(2, device=cuda2)e = torch.randn(2).to(cuda2)f = torch.randn(2).cuda(cuda2)# d.device, e.device, and f.device are all device(type='cuda', index=2)