CUDA 7 流并发性优化

异构计算是指高效地使用系统中的所有处理器，包括 CPU 和 GPU 。为此，应用程序必须在多个处理器上并发执行函数。 CUDA 应用程序通过在 streams 中执行异步命令来管理并发性，这些命令是按顺序执行的。不同的流可以并发地执行它们的命令，也可以彼此无序地执行它们的命令。

在不指定流的情况下执行异步 CUDA 命令时，runtime使用默认流。在 CUDA 7 之前，默认流是一个特殊流，它隐式地与设备上的所有其它流同步。

CUDA 7 引入了大量强大的新功能，包括一个新的选项，可以为每个主机线程使用独立的默认流，这避免了传统默认流的序列化。本文将展示如何在 CUDA 程序中简化实现内核和数据副本之间的并发。

CUDA 中的异步命令

如 CUDA C 编程指南所述，异步命令在设备完成请求的任务之前将控制权返回给调用主机线程（非阻塞的）。这些命令是：

·
内核启动；

·
存储器在两个地址之间复制到同一设备存储器；

·
从主机到设备的 64kb 或更少内存块的内存拷贝；

·
由后缀为 Async 的函数执行的内存复制；

·
内存设置函数调用。

为内核启动或主机设备内存复制指定流是可选的；可以调用 CUDA 命令而不指定流（或通过将 stream 参数设置为零）。下面两行代码都在默认流上启动内核。

kernel<<< blocks, threads, bytes >>>(); // default stream

kernel<<< blocks, threads, bytes, 0 >>>(); // stream 0

默认流

在并发性对性能不重要的情况下，默认流很有用。在 CUDA 7 之前，每个设备都有一个用于所有主机线程的默认流，这会导致隐式同步。正如 CUDA C 编程指南中的“隐式同步”一节所述，如果主机线程向它们之间的默认流发出任何 CUDA 命令，来自不同流的两个命令就不能并发运行。

CUDA 7 引入了一个新选项，每线程默认流，它有两个效果。首先，它为每个主机线程提供自己的默认流。这意味着不同主机线程向默认流发出的命令可以并发运行。其次，这些默认流是常规流。这意味着默认流中的命令可以与非默认流中的命令同时运行。

要在 nvcc 7 及更高版本中启用每线程默认流，可以在包含 CUDA 头（ cuda.h 或 cuda_runtime.h ）之前，使用 nvcc 命令行选项 CUDA 或 #define 编译 CUDA_API_PER_THREAD_DEFAULT_STREAM 预处理器宏。需要注意的是：当代码由 nvcc 编译时，不能使用 #define
CUDA_API_PER_THREAD_DEFAULT_STREAM 在. cu 文件中启用此行为，因为 nvcc 在翻译单元的顶部隐式包含了 cuda_runtime.h 。

多流示例

看一个小例子。下面的代码简单地在八个流上启动一个简单内核的八个副本。只为每个网格启动一个线程块，就有足够的资源同时运行多个线程块。默认流如何导致序列化的示例，在默认流上添加了不起作用的虚拟内核启动。这是密码。

const int N = 1 << 20;

global void kernel(float *x, int n)

{

int tid = threadIdx.x + blockIdx.x * blockDim.x;

for (int i = tid; i < n; i += blockDim.x * gridDim.x) {

x[i] = sqrt(pow(3.14159,i));

}

int main()

{

const int num_streams = 8;

cudaStream_t streams[num_streams];

float *data[num_streams];

for (int i = 0; i < num_streams; i++) {

cudaStreamCreate(&streams[i]);

cudaMalloc(&data[i], N * sizeof(float));

// launch one worker kernel per stream

kernel<<<1, 64, 0, streams[i]>>>(data[i], N);

// launch a dummy kernel on the default stream

kernel<<<1, 1>>>(0, 0);

}

cudaDeviceReset();

return 0;

}

首先检查遗留行为，通过不带选项的编译。

nvcc ./stream_test.cu -o stream_legacy

可以在 NVIDIA visualprofiler （ nvvp ）中运行该程序，以获得显示所有流和内核启动的时间轴。图 1 显示了 Macbook Pro 上生成的内核时间线，该 Macbook Pro 带有 NVIDIA GeForce GT 750M （一台开普勒 GPU ）。可以看到默认流上虚拟内核的非常小，以及它们如何导致所有其他流序列化。

一个简单的多流示例在将任何交错内核发送到默认流时不会实现并发。

现在尝试新的单线程默认流。

nvcc --default-stream per-thread ./stream_test.cu -o stream_per-thread

图 2 显示了来自 nvvp 的结果。可以看到九个流之间的完全并发：默认流（在本例中映射到流 14 ）和创建的其它八个流。虚拟内核运行得如此之快，以至于很难看到在这个图像中默认流上有八个调用。

图 2 ：使用新的每线程默认流选项的多流示例，它支持完全并发执行。

多线程示例

来看另一个例子，该示例旨在演示新的默认流行为如何使多线程应用程序，更容易实现执行并发。下面的例子创建了八个 POSIX 线程，每个线程在默认流上调用内核，然后同步默认流。（需要在本例中进行同步，以确保探查器在程序退出之前获得内核开始和结束时间戳。）

#include <pthread.h>

#include <stdio.h>

const int N = 1 << 20;

global void kernel(float *x, int n)

{

int tid = threadIdx.x + blockIdx.x * blockDim.x;

for (int i = tid; i < n; i += blockDim.x * gridDim.x) {

x[i] = sqrt(pow(3.14159,i));

}

void *launch_kernel(void *dummy)

{

float *data;

cudaMalloc(&data, N * sizeof(float));

kernel<<<1, 64>>>(data, N);

cudaStreamSynchronize(0);

return NULL;

}

int main()

{

const int num_threads = 8;

pthread_t threads[num_threads];

for (int i = 0; i < num_threads; i++) {

if (pthread_create(&threads[i], NULL, launch_kernel, 0)) {

        fprintf(stderr, "Error

creating threadn");

        return 1;

}

for (int i = 0; i < num_threads; i++) {

if(pthread_join(threads[i], NULL)) {

        fprintf(stderr, "Error joining

threadn");

        return 2;

}

cudaDeviceReset();

return 0;

}

首先，编译时不使用任何选项来测试遗留的默认流行为。

nvcc ./pthread_test.cu -o pthreads_legacy

在 nvvp 中运行它时，看到一个流，默认流，所有内核启动都序列化，如图 3 所示。

图 3 ：一个具有默认流行为的多线程示例：所有八个线程都被序列化。

用新的 per-thread default stream 选项编译它。

nvcc --default-stream per-thread ./pthread_test.cu -o pthreads_per_thread

图 4 显示，对于每个线程的默认流，每个线程都会自动创建一个新的流，它们不会同步，因此所有八个线程的内核都会并发运行。

图 4 ：单线程默认流的多线程示例：所有八个线程的内核同时运行。

注意

在为并发进行编程时，还需要记住以下几点。

·
记住：对于每线程的默认流，每个线程中的默认流的行为与常规流相同，只要同步和并发就可以了。对于传统的默认流，这是不正确的。

·
–default-stream 选项是按编译单元应用的，确保将其应用于所有需要它的 nvcc 命令行。

·
cudaDeviceSynchronize() 继续同步设备上的所有内容，甚至使用新的每线程默认流选项。如果只想同步单个流，请使用 cudaStreamSynchronize(cudaStream_t stream) ，如第二个示例所示。

· 从 CUDA 7 开始，还可以使用句柄 cudaStreamPerThread 显式地访问每线程的默认流，也可以使用句柄 cudaStreamLegacy 访问旧的默认流。请注意， cudaStreamLegacy 仍然隐式地与每个线程的默认流同步，如果碰巧在一个程序中混合使用它们。

· 可以通过将 cudaStreamCreate() 标志传递给 cudaStreamCreate() 来创建不与传统默认流同步的非阻塞流。

CUDA 7 流并发性优化相关推荐

CUDA Stream流并发性
目录 1.CUDA 中的异步命令 2.默认流 3.Multistream多流示例异构计算是指高效地使用系统中的所有处理器,包括 CPU 和 GPU .为此,应用程序必须在多个处理器上并发执行函数. ...
CUDA 7 Stream流简化并发性
CUDA 7 Stream流简化并发性异构计算是指高效地使用系统中的所有处理器,包括 CPU 和 GPU .为此,应用程序必须在多个处理器上并发执行函数. CUDA 应用程序通过在 streams ...
CUDA 7流简化并发
CUDA 7流简化并发异构计算是指有效使用系统中的所有处理器,包括CPU和GPU.为此,应用程序必须在多个处理器上同时执行功能.CUDA应用程序通过在流(按顺序执行的命令序列)中,执行异步命令来管理 ...
一个iOS流畅性优化工具
????????关注后回复 "进群" ,拉你进程序员交流群???????? 转自:掘金 BangRaJun https://juejin.cn/post/693472015254 ...
Cuda Stream流分析
Cuda Stream流分析 Stream 一般来说,cuda c并行性表现在下面两个层面上: • Kernel level • Grid level Stream和event简介 Cuda stre ...
万字干货 | Python后台开发的高并发场景优化解决方案
嘉宾 | 黄思涵来源 | AI科技大本营在线公开课互联网发展到今天,规模变得越来越大,也对所有的后端服务提出了更高的要求.在平时的工作中,我们或多或少都遇到过服务器压力过大问题.针对该问题,本次公 ...
探索 ConcurrentHashMap 高并发性的实现机制
简介 ConcurrentHashMap 是 util.concurrent 包的重要成员.本文将结合 Java 内存模型,分析 JDK 源代码,探索 ConcurrentHashMap 高并发的具体 ...
DB2 9 根蒂根基底细（730 考试）认证指南，第 6 局部: 数据并发性(5)
锁和机能锁兼容性假设数据本钱上的一种锁形态允许在同一本钱上布置另一个锁,就感受这两种锁(或两种形态)是兼容的.每当一个事项持有数据本钱上的锁,而第二个事项哀告同一本钱上的锁时,DB2 数据库管理按 ...
提高代码性能及并发性的方法浅谈
最近在做系统调优,总结了下cache相关知识,以及如何提高性能和并发性能的方法. 一CACHE相关 1. cache概述 cache,中译名高速缓冲存储器,其作用是为了更好的利用局部性原理,减少CPU ...

CUDA 7 流并发性优化

CUDA 7 流并发性优化相关推荐

最新文章

热门文章