[6] CUDA之线程同步

CUDA之线程同步

  • 共享内存:线程时间需要互相交换数据才能完成任务的情况并不少见,因此,必须存在某种能让线程彼此交流的机制
  • 当很多线程并行工作并且访问相同的数据或者存储器位置的时候,线程间必须正确的同步
  • 线程之间交换数据并不一定要需要使用共享内存,只是共享内存较快而已

1.共享内存

  • 共享内存位于芯片内部,因此它比全局内存要快得多,相比没有经过缓存的全局内存访问,共享内存大约在延迟上第100倍
  • 同一个块中的线程可以访问相同的一段共享内存,不同块中的线程所见到的共享内存中的内容是不相同的
  • 如果某线程的计算结果在写入到共享内存完成之前被其他线程读取,那么将会导致错误。因此应该正确的控制和管理内存访问,这是由 __syncthreads() 指令完成的,该指令确保在继续执行程序之前完成对内存的所有写入操作,即同步,也被称为 barrierbarrier的含义是块中的所有线程都将到达该代码行,然后在此等待其他线程完成,当所有线程都到达了这里之后,他们可以一起继续往下执行
  • 举个例子:

#include
### CUDA 线程同步 示例代码 为了确保线程间的正确协作,在CUDA编程中常常需要使用不同的同步原语。下面展示了一个利用`__syncthreads()`函数进行线程同步的例子,该函数用于阻塞直到同一线程块内的所有线程都到达这一点[^1]。 ```cpp #include <cuda_runtime.h> #include <stdio.h> // Kernel function to add elements of two arrays __global__ void add(int n, float *x, float *y) { int index = blockIdx.x * blockDim.x + threadIdx.x; int stride = blockDim.x * gridDim.x; for (int i = index; i < n; i += stride){ y[i] = x[i] + y[i]; // Synchronize all threads within a block before proceeding. __syncthreads(); } } int main(void) { int N = 1 << 20; int size = N * sizeof(float); // Allocate host memory float *h_x = (float *)malloc(size); float *h_y = (float *)malloc(size); // Initialize input vectors for (int i = 0; i < N; ++i) { h_x[i] = 1.0f; h_y[i] = 2.0f; } // Declare device pointers and allocate GPU memory float *d_x = NULL; float *d_y = NULL; cudaMalloc((void **)&d_x, size); cudaMalloc((void **)&d_y, size); // Copy data from CPU to GPU cudaMemcpy(d_x, h_x, size, cudaMemcpyHostToDevice); cudaMemcpy(d_y, h_y, size, cudaMemcpyHostToDevice); // Define number of blocks and threads per block int threads_per_block = 256; int blocks_per_grid =(N + threads_per_block - 1) / threads_per_block; // Launch kernel with synchronization between threads inside each block add<<<blocks_per_grid, threads_per_block>>>(N, d_x, d_y); // Wait until the GPU has completed its tasks cudaDeviceSynchronize(); // Copy result back to host cudaMemcpy(h_y, d_y, size, cudaMemcpyDeviceToHost); // Free GPU memory cudaFree(d_x); cudaFree(d_y); // Check results on some values... bool success = true; for (int i = 0; i < N && success; ++i) { if ((h_x[i] + h_y[i]) != (1.0f + 2.0f)) { printf("Mismatch at element %d: was %f but should be %f.\n", i, h_y[i], 3.0f); success = false; } } free(h_x); free(h_y); return !success ? EXIT_FAILURE : EXIT_SUCCESS; } ``` 上述代码展示了如何通过调用`__syncthreads()`来保证同一个线程块内部的所有线程在继续之前都已经完成了当前迭代的操作[^5]。
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包

打赏作者

明月醉窗台

你的鼓励将是我创作的最大动力

¥1 ¥2 ¥4 ¥6 ¥10 ¥20
扫码支付:¥1
获取中
扫码支付

您的余额不足,请更换扫码支付或充值

打赏作者

实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值