【CUDA学习笔记】第四篇:线程以及线程同步(附案例代码下载方式)(二)

简介: 【CUDA学习笔记】第四篇:线程以及线程同步(附案例代码下载方式)(二)

3、线程同步


3.1、共享内存

   共享内存位于芯片内部,因此它比全局内存快得多。(CUDA里面存储器的快慢有两方面,一个是延迟低,一个是带宽大。这里特指延迟低),相比没有经过缓存的全局内存访问,共享内存大约在延迟上低100倍。同一个块中的线程可以访问相同的一段共享内存(注意:不同块中的线程所见到的共享内存中的内容是不相同的),这在许多线程需要与其他线程共享它们的结果的应用程序中非常有用。但是如果不同步,也可能会造成混乱或错误的结果。如果某线程的计算结果在写入到共享内存完成之前被其他线程读取,那么将会导致错误。因此,应该正确地控制或管理内存访问。这是由__syncthreads()指令完成的,该指令确保在继续执行程序之前完成对内存的所有写入操作。这也被称为barrierbarrier的含义是块中的所有线程都将到达该代码行,然后在此等待其他线程完成。当所有线程都到达了这里之后,它们可以一起继续往下执行。

#include <stdio.h>
__global__ void gpu_shared_memory(float *d_a)
{
// Defining local variables which are private to each thread
int i, index = threadIdx.x;
float average, sum = 0.0f;
//Define shared memory
__shared__ float sh_arr[10];
sh_arr[index] = d_a[index];
__syncthreads();    // This ensures all the writes to shared memory have completed
for (i = 0; i<= index; i++) 
{ 
sum += sh_arr[i]; 
}
average = sum / (index + 1.0f);
d_a[index] = average; 
sh_arr[index] = average;
}
int main(int argc, char **argv)
{
//Define Host Array
float h_a[10];   
//Define Device Pointer
float *d_a;       
for (int i = 0; i < 10; i++) {
h_a[i] = i;
}
// allocate global memory on the device
cudaMalloc((void **)&d_a, sizeof(float) * 10);
// now copy data from host memory  to device memory 
cudaMemcpy((void *)d_a, (void *)h_a, sizeof(float) * 10, cudaMemcpyHostToDevice);
gpu_shared_memory << <1, 10 >> >(d_a);
// copy the modified array back to the host memory
cudaMemcpy((void *)h_a, (void *)d_a, sizeof(float) * 10, cudaMemcpyDeviceToHost);
printf("Use of Shared Memory on GPU:  \n");
//Printing result on console
for (int i = 0; i < 10; i++) {
printf("The running average after %d element is %f \n", i, h_a[i]);
}
return 0;
}

   在main函数中,当分配好主机和设备上的数组后,用0.0到9.0填充主机上的数组,然后将这个数组复制到显存。内核将对显存中的数据进行读取,计算并保存结果。最后结果从显存中传输到内存,然后在控制台上输出。控制台上的输出结果如图所示:

   这个程序还含有额外的一个CUDA函数调用:cudaDeviceSynchronize()。为何要加这句?这是因为启动内核是一个异步操作,只要发布了内核启动命令,不等内核执行完成,控制权就会立刻返回给调用内核的CPU线程。在上述的代码中,CPU线程返回,继续执行的下一句是printf()。而再之后,在内核完成之前,进程就会结束,终止控制台窗口。所以,如果不加上这句同步函数,你就看不到任何的内核执行结果输出。在程序退出后内核生成的输出结果,将没有地方可去,你没法看到它们,因此,如果我们不包含这个指令,你将不会看到任何内核执行的printf语句的输出结果。要能看到内核生成的输出结果,我们必须包含这句同步函数。这样,内核的结果将通过可用的标准输出显示,而应用程序则会在内核执行完成之后才退出。


3.2、原子操作

   考虑当大量的线程需要试图修改一段较小的内存区域的情形,这是(在日常的算法实现中)常发生的现象。当我们试图进行“读取-修改-写入”操作序列的时候,这种情形经常会带来很多麻烦。

   

   一个例子是代码d_out[i]++,这代码首先将d_out[i]的原值从存储器中读取出来,然后执行了+1操作,再将结果回写到存储器。然而,如果多个线程试图在同一个内存区域中进行这个操作,则可能会得到错误的结果。


   假设某内存区域中有初始值6,两个线程p和q分别试图将这段区域中的内容+1,则最终的结果应当是8。但是在实际执行的时候,可能p和q两个线程同时读取了这个初始值,两者都得到了6,执行+1操作都得到了7,然后它们将7写回这个内存区域。这样,和正确的结果8不同,我们得到的最终结果是7,这是错误的。这种错误是如何的危险,我们通过ATM取现操作来演示。假设你的账户余额为5000卢比,你的账户下面开了两张银行卡,你和你的朋友同时去2个不同的ATM上取现4000卢比,你俩在同一瞬间刷卡取现。所以,当两个ATM检查余额的时候,都将显示5000卢比的余额。当你俩同时取现4000卢比的时候,两个ATM机都只根据初始值5000卢比判断,要取的现金4000卢比小于当前余额。所以两个机器将会给你们每人4000卢比。即使你之前只有5000卢比的余额,你们也能得到8000卢比,这很危险。为了示范一下这种情形,做了一个很多线程试图同时访问一个小数组的例子:

#include <stdio.h>
#define NUM_THREADS 10000
#define SIZE  10
#define BLOCK_WIDTH 100
__global__ void gpu_increment_without_atomic(int *d_a)
{
// Calculate thread id for current thread
int tid = blockIdx.x * blockDim.x + threadIdx.x;
// each thread increments elements wrapping at SIZE variable
tid = tid % SIZE;
d_a[tid] += 1;
}
int main(int argc, char **argv)
{
printf("%d total threads in %d blocks writing into %d array elements\n",
NUM_THREADS, NUM_THREADS / BLOCK_WIDTH, SIZE);
// declare and allocate host memory
int h_a[SIZE];
const int ARRAY_BYTES = SIZE * sizeof(int);
// declare and allocate GPU memory
int * d_a;
cudaMalloc((void **)&d_a, ARRAY_BYTES);
//Initialize GPU memory to zero
cudaMemset((void *)d_a, 0, ARRAY_BYTES);
gpu_increment_without_atomic << <NUM_THREADS / BLOCK_WIDTH, BLOCK_WIDTH >> >(d_a);
// copy back the array to host memory
cudaMemcpy(h_a, d_a, ARRAY_BYTES, cudaMemcpyDeviceToHost);
printf("Number of times a particular Array index has been incremented without atomic add is: \n");
for (int i = 0; i < SIZE; i++)
{
printf("index: %d --> %d times\n ", i, h_a[i]);
}
cudaFree(d_a);
return 0;
}

image.png

   可能如同已经猜到的那样,每次运行你的程序,每个内存区域中的元素值都可能会不同。这是设备上不定顺序的多线程执行导致的。

   为了解决这个问题,CUDA提供了atomicAdd这种原子操作函数。该函数会从逻辑上保证,每个调用它的线程对相同的内存区域上的“读取旧值-累加-回写新值”操作是不可被其他线程扰乱的原子性的整体完成的。使用atomicAdd进行原子累加的内核函数代码如下:

#include <stdio.h>
#define NUM_THREADS 10000
#define SIZE  10
#define BLOCK_WIDTH 100
__global__ void gpu_increment_atomic(int *d_a)
{
// Calculate thread id for current thread
int tid = blockIdx.x * blockDim.x + threadIdx.x;
// each thread increments elements wrapping at SIZE variable
tid = tid % SIZE;
atomicAdd(&d_a[tid], 1);
}
int main(int argc, char **argv)
{
printf("%d total threads in %d blocks writing into %d array elements\n",
NUM_THREADS, NUM_THREADS / BLOCK_WIDTH, SIZE);
// declare and allocate host memory
int h_a[SIZE];
const int ARRAY_BYTES = SIZE * sizeof(int);
// declare and allocate GPU memory
int * d_a;
cudaMalloc((void **)&d_a, ARRAY_BYTES);
//Initialize GPU memory to zero
cudaMemset((void *)d_a, 0, ARRAY_BYTES);
gpu_increment_atomic << <NUM_THREADS / BLOCK_WIDTH, BLOCK_WIDTH >> >(d_a);
// copy back the array to host memory
cudaMemcpy(h_a, d_a, ARRAY_BYTES, cudaMemcpyDeviceToHost);
printf("Number of times a particular Array index has been incremented is: \n");
for (int i = 0; i < SIZE; i++) 
{ 
printf("index: %d --> %d times\n ", i, h_a[i]); 
}
cudaFree(d_a);
return 0;
}

   在main函数中,具有10个元素的数组被初始化成0值,然后传递给了内核,但现在,内核中的代码将执行原子累加操作。所以,这个程序输出的结果将是对的,数组中的每个元素将被累加1000。运行结果显示如图:

   如果你测量一下这个程序的运行时间,相比之前的那个简单地在全局内存上直接进行加法操作的程序它用的时间更长。这是因为使用原子操作后程序具有更大的执行代价。可以通过使用共享内存来加速这些原子累加操作。如果线程规模不变,但原子操作的元素数量扩大,则这些同样次数的原子操作会更快地完成。这是因为更广泛的分布范围上的原子操作有利于利用多个能执行原子操作的单元,以及每个原子操作单元上面的竞争性的原子事务也相应减少了。

相关实践学习
部署Stable Diffusion玩转AI绘画(GPU云服务器)
本实验通过在ECS上从零开始部署Stable Diffusion来进行AI绘画创作,开启AIGC盲盒。
相关文章
|
1月前
|
编解码 数据安全/隐私保护 计算机视觉
Opencv学习笔记(十):同步和异步(多线程)操作打开海康摄像头
如何使用OpenCV进行同步和异步操作来打开海康摄像头,并提供了相关的代码示例。
79 1
Opencv学习笔记(十):同步和异步(多线程)操作打开海康摄像头
|
3月前
|
数据采集 XML JavaScript
C# 中 ScrapySharp 的多线程下载策略
C# 中 ScrapySharp 的多线程下载策略
|
24天前
|
安全 Java
Java多线程通信新解:本文通过生产者-消费者模型案例,深入解析wait()、notify()、notifyAll()方法的实用技巧
【10月更文挑战第20天】Java多线程通信新解:本文通过生产者-消费者模型案例,深入解析wait()、notify()、notifyAll()方法的实用技巧,包括避免在循环外调用wait()、优先使用notifyAll()、确保线程安全及处理InterruptedException等,帮助读者更好地掌握这些方法的应用。
16 1
|
1月前
|
安全 Java 开发者
在多线程编程中,确保数据一致性与防止竞态条件至关重要。Java提供了多种线程同步机制
【10月更文挑战第3天】在多线程编程中,确保数据一致性与防止竞态条件至关重要。Java提供了多种线程同步机制,如`synchronized`关键字、`Lock`接口及其实现类(如`ReentrantLock`),还有原子变量(如`AtomicInteger`)。这些工具可以帮助开发者避免数据不一致、死锁和活锁等问题。通过合理选择和使用这些机制,可以有效管理并发,确保程序稳定运行。例如,`synchronized`可确保同一时间只有一个线程访问共享资源;`Lock`提供更灵活的锁定方式;原子变量则利用硬件指令实现无锁操作。
20 2
|
1月前
FFmpeg学习笔记(二):多线程rtsp推流和ffplay拉流操作,并储存为多路avi格式的视频
这篇博客主要介绍了如何使用FFmpeg进行多线程RTSP推流和ffplay拉流操作,以及如何将视频流保存为多路AVI格式的视频文件。
170 0
|
3月前
|
调度
多线程学习笔记
这篇文章是多线程学习笔记,涵盖了线程与进程的概念、多线程实现方式、线程状态、线程同步与不安全示例、死锁问题以及生产者与消费者问题等多线程编程的关键知识点。
多线程学习笔记
|
2月前
|
安全 Java 调度
python3多线程实战(python3经典编程案例)
该文章提供了Python3中多线程的应用实例,展示了如何利用Python的threading模块来创建和管理线程,以实现并发执行任务。
40 0
|
3月前
|
Java Windows
【Azure Developer】Windows中通过pslist命令查看到Java进程和线程信息,但为什么和代码中打印出来的进程号不一致呢?
【Azure Developer】Windows中通过pslist命令查看到Java进程和线程信息,但为什么和代码中打印出来的进程号不一致呢?
|
3月前
|
Java
多线程线程同步
多线程的锁有几种方式
|
3月前
|
消息中间件 安全 Kafka
"深入实践Kafka多线程Consumer:案例分析、实现方式、优缺点及高效数据处理策略"
【8月更文挑战第10天】Apache Kafka是一款高性能的分布式流处理平台,以高吞吐量和可扩展性著称。为提升数据处理效率,常采用多线程消费Kafka数据。本文通过电商订单系统的案例,探讨了多线程Consumer的实现方法及其利弊,并提供示例代码。案例展示了如何通过并行处理加快订单数据的处理速度,确保数据正确性和顺序性的同时最大化资源利用。多线程Consumer有两种主要模式:每线程一个实例和单实例多worker线程。前者简单易行但资源消耗较大;后者虽能解耦消息获取与处理,却增加了系统复杂度。通过合理设计,多线程Consumer能够有效支持高并发数据处理需求。
168 4