[CUDA编程原理] CUDA Stream - 流同步和流管理

关于CUDA流[1]

CUDA流表示一个GPU操作队列,该队列中的操作将以添加到流中的先后顺序而依次执行。可以将一个流看做是GPU上的一个任务,不同任务可以并行执行。

CUDA流执行

CUDA流同步[2]

  • 默认流是一个隐式流(无需显式创建,CUDA中默认存在),它与同一CUcontext[3]中的所有其他流同步,非阻塞流(non-blocking streams)除外。

例如,在如下代码中,在流 s 中启动核函数 k_A,然后在默认流中启动 k_B,在流 s 中启动 k_C

k_A<<<1, 1, 0, s>>>();
k_B<<<1, 1>>>();
k_C<<<1, 1, 0, s>>>();

其运行结果k_A阻塞k_B的运行,k_B阻塞k_C的运行。

  • 如果操作[4]A和操作B位于同一个CUDA流中,操作A和操作B按照先后顺序执行;

例如,在如下代码中,在流 s 中先后启动核函数 k_Ak_B

k_A<<<1, 1, 0, s>>>();
k_B<<<1, 1, 0, s>>>();

其运行结果k_A阻塞k_B的运行。

  • 如果两个操作位于不同的CUDA流中,则操作A和操作B并没有确定的执行顺序;

例如,在如下代码中,先在流 s1 中先后启动核函数 k_A1k_B1,然后在流 s2 中先后启动核函数 k_A2k_B2

k_A1<<<1, 1, 0, s1>>>();
k_B1<<<1, 1, 0, s1>>>();
k_A2<<<1, 1, 0, s2>>>();
k_B2<<<1, 1, 0, s2>>>();

在实际运行结果中,流s1的操作并不会影响流s2的操作,k_A1会阻塞k_B1的运行,而不会阻塞k_B2的运行。

CUDA流管理[5]

创建流

在CUDA中,流的创建有三种方式:cudaStreamCreatecudaStreamCreateWithFlagscudaStreamCreateWithPriority

创建CUDA流

__host__ cudaError_t cudaStreamCreate ( cudaStream_t* pStream )

创建不同标记的CUDA流

__host__ cudaError_t cudaStreamCreateWithFlags ( cudaStream_t* pStream , unsigned int flags )

  • cudaStreamDefault: CUDA流的默认创建标记;
  • cudaStreamNonBlocking: 创建非阻塞流。在通过该标记创建的CUDA流中运行的工作可以与流0(NULL 流)中的工作同时运行,并且通过该标记创建的流不应与流0执行隐式同步(implicit synchronization)。

创建不同优先级的CUDA流

__host__ cudaError_t cudaStreamCreateWithPriority ( cudaStream_t* pStream , unsigned int flags, int priority )

通过priority参数,指定CUDA 流的优先级,数值越小其优先级越高(“0”表示默认优先级)。高优先级流中的计算核函数(compute kernels)可以抢占已经在低优先级流中执行的核函数,而主机和设备之间的内存操作不受优先级的影响。

销毁流

__host____device__ cudaError_t cudaStreamDestroy ( cudaStream_t stream )

销毁并清理异步流(asynchronous stream)。

流同步操作

__host__ cudaError_t cudaStreamSynchronize ( cudaStream_t stream )

调用该操作时会发生阻塞,等待指定流中操作完成。


  1. CUDA CONCURRENCY Bob Crovella, 7/21/2020

  2. CUDA Runtime API >> Stream synchronization behavior

  3. CUDA Drive API >> CUDA context >> CUcontext

  4. 这里的操作通常指异步内存拷贝(cudaMemcpyAsync)和核函数调用。

  5. CUDA Runtime API >> Stream Management

最后编辑于
©著作权归作者所有,转载或内容合作请联系作者
  • 序言:七十年代末,一起剥皮案震惊了整个滨河市,随后出现的几起案子,更是在滨河造成了极大的恐慌,老刑警刘岩,带你破解...
    沈念sama阅读 157,801评论 4 361
  • 序言:滨河连续发生了三起死亡事件,死亡现场离奇诡异,居然都是意外死亡,警方通过查阅死者的电脑和手机,发现死者居然都...
    沈念sama阅读 66,870评论 1 290
  • 文/潘晓璐 我一进店门,熙熙楼的掌柜王于贵愁眉苦脸地迎上来,“玉大人,你说我怎么就摊上这事。” “怎么了?”我有些...
    开封第一讲书人阅读 107,593评论 0 239
  • 文/不坏的土叔 我叫张陵,是天一观的道长。 经常有香客问我,道长,这世上最难降的妖魔是什么? 我笑而不...
    开封第一讲书人阅读 43,771评论 0 203
  • 正文 为了忘掉前任,我火速办了婚礼,结果婚礼上,老公的妹妹穿的比我还像新娘。我一直安慰自己,他们只是感情好,可当我...
    茶点故事阅读 52,142评论 3 286
  • 文/花漫 我一把揭开白布。 她就那样静静地躺着,像睡着了一般。 火红的嫁衣衬着肌肤如雪。 梳的纹丝不乱的头发上,一...
    开封第一讲书人阅读 40,486评论 1 215
  • 那天,我揣着相机与录音,去河边找鬼。 笑死,一个胖子当着我的面吹牛,可吹牛的内容都是我干的。 我是一名探鬼主播,决...
    沈念sama阅读 31,777评论 2 311
  • 文/苍兰香墨 我猛地睁开眼,长吁一口气:“原来是场噩梦啊……” “哼!你这毒妇竟也来了?” 一声冷哼从身侧响起,我...
    开封第一讲书人阅读 30,461评论 0 196
  • 序言:老挝万荣一对情侣失踪,失踪者是张志新(化名)和其女友刘颖,没想到半个月后,有当地人在树林里发现了一具尸体,经...
    沈念sama阅读 34,181评论 1 240
  • 正文 独居荒郊野岭守林人离奇死亡,尸身上长有42处带血的脓包…… 初始之章·张勋 以下内容为张勋视角 年9月15日...
    茶点故事阅读 30,450评论 2 244
  • 正文 我和宋清朗相恋三年,在试婚纱的时候发现自己被绿了。 大学时的朋友给我发了我未婚夫和他白月光在一起吃饭的照片。...
    茶点故事阅读 31,971评论 1 258
  • 序言:一个原本活蹦乱跳的男人离奇死亡,死状恐怖,灵堂内的尸体忽然破棺而出,到底是诈尸还是另有隐情,我是刑警宁泽,带...
    沈念sama阅读 28,305评论 2 252
  • 正文 年R本政府宣布,位于F岛的核电站,受9级特大地震影响,放射性物质发生泄漏。R本人自食恶果不足惜,却给世界环境...
    茶点故事阅读 32,952评论 3 235
  • 文/蒙蒙 一、第九天 我趴在偏房一处隐蔽的房顶上张望。 院中可真热闹,春花似锦、人声如沸。这庄子的主人今日做“春日...
    开封第一讲书人阅读 26,039评论 0 8
  • 文/苍兰香墨 我抬头看了看天上的太阳。三九已至,却和暖如春,着一层夹袄步出监牢的瞬间,已是汗流浃背。 一阵脚步声响...
    开封第一讲书人阅读 26,789评论 0 193
  • 我被黑心中介骗来泰国打工, 没想到刚下飞机就差点儿被人妖公主榨干…… 1. 我叫王不留,地道东北人。 一个月前我还...
    沈念sama阅读 35,519评论 2 270
  • 正文 我出身青楼,却偏偏与公主长得像,于是被迫代替她去往敌国和亲。 传闻我的和亲对象是个残疾皇子,可洞房花烛夜当晚...
    茶点故事阅读 35,410评论 2 264

推荐阅读更多精彩内容