Atomic Functions

原创 2015年07月09日 10:34:47

考虑两个线程均是要往同一个全局或者共享数据中。



如果x的初始值是10,那么x的最终结果是?答案是无法确定的,主要是因为有访问冲突。


Atomic Functions 就是读-修改-写操作时避免与其它线程冲突,计算时会将其地址锁定,直到结束计算。


atomic opeations:

intatomicAdd(int* address, intval); 

intatomicSub(int* address, intval);

intatomicExch(int* address, intval);

intatomicMin(int* address, intval);

intatomicMax(int* address, intval);

unsigned intatomicInc(unsigned int* address, unsigned intval);

unsigned intatomicDec(unsigned int* address, unsigned intval);

intatomicCAS(int* address, int compare, intval); //compare and swap

intatomicAnd(int* address, intval);

intatomicOr(int* address, intval);

intatomicXor(int* address, intval);  



测试例子:

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "device_functions.h"
#include <iostream>

using namespace std;

__device__ int gpu_hist[10];

__global__ void init()
{
	int tid = blockIdx.x * blockDim.x + threadIdx.x;
	gpu_hist[tid] = 0;
}

__global__ void gpu_histogram(int *a, int n)
{
	//int *ptr;
	int tid = blockIdx.x * blockDim.x + threadIdx.x;
	int numberThreads = blockDim.x * gridDim.x;
	while (tid < n)
	{
		//ptr = &gpu_hist[a[tid]];
		//atomicAdd(ptr, 1);
		gpu_hist[a[tid]]++; // have no atomic functions
		tid += numberThreads;
	}
}

int main()
{
	int N = 32;
	int *a, *dev_a;
	int hist[10];
	int size = N * sizeof(int);
	a = (int *)malloc(size);
	srand(1);
	for (int i = 0; i < N; ++i)
	{
		a[i] = rand() % 10;
		printf("%d ", a[i]);
	}
	printf("\n");
	cudaMalloc((void**)&dev_a, size);
	cudaMemcpy(dev_a, a, size, cudaMemcpyHostToDevice);
	init << <1, 10 >> >();
	gpu_histogram<<<1, 32>>>(dev_a, N);
	cudaThreadSynchronize();
	cudaMemcpyFromSymbol(&hist, gpu_hist, 10 * sizeof(int));
	printf("Histogram as computed on GPU\n");
	for (int i = 0; i < 10; ++i)
	{
		printf("Number of %d s = %d\n", i, hist[i]);
	}
	free(a);
	cudaFree(dev_a);
}





Built-in functions for atomic memory access

The following builtins are intended to be compatible with those described in the Intel Itanium Proce...
  • HedpatCzw
  • HedpatCzw
  • 2016年01月11日 12:18
  • 169

C++11中<atomic>的使用

C++11中的使用
  • fengbingchun
  • fengbingchun
  • 2017年06月18日 19:13
  • 1217

Java的Atomic类分析

Atomic包介绍 Java1.5的Atomic包名为java.util.concurrent.atomic。这个包提供了一系列原子类。这些类可以保证多线程环境下,当某个线程在执行atomic的方法...
  • goodlixueyong
  • goodlixueyong
  • 2016年05月08日 18:53
  • 3403

C++并发实战16: std::atomic原子操作

针对原子类型操作要不一步完成,要么不做,防止由于多线程串行化带来的错误。 1  std::atomic_flag是一个bool原子类型,其支持test_and_set和clear两个操作,atomic...
  • liuxuejiang158
  • liuxuejiang158
  • 2013年12月19日 15:50
  • 22905

关于atomic到底安不安全

当提到atomic的时候,大家第一反应就是与nonatomic对应。nonatomic是非原子性,不安全。atomic是原子性的,安全。 但是在网上又看到很多说atomic并不是安全的。是不是到这里...
  • H_Qiao
  • H_Qiao
  • 2018年01月12日 10:38
  • 77

Legacy __sync Built-in Functions for Atomic Memory Access

Legacy __sync Built-in Functions for Atomic Memory Access   copy from: http://gcc.gnu.org/online...
  • zlyong0018
  • zlyong0018
  • 2013年08月02日 18:12
  • 388

[docker]coreOS与atomic对比

摘自https://major.io/2014/05/13/coreos-vs-project-atomic-a-review/。 【部署】 coreOS: 通过云上的镜像或者PXE部署。在云上可以...
  • halcyonbaby
  • halcyonbaby
  • 2014年09月19日 21:31
  • 7817

《Go in Action 2015.11.pdf》之6.4.1 Atomic functions

利用atomic递增 Channels Unbuffered channels
  • KingEasternSun
  • KingEasternSun
  • 2017年03月29日 18:13
  • 274

linux下mutex与atomic性能比较

一种是用boost::atomic;一种直接加锁;代码很简单: #include #include #include #include #include #include #inclu...
  • boyhailong
  • boyhailong
  • 2015年01月28日 00:27
  • 2360

atomic 和noatomic

曾经面试时候被被问到最多的问题是 nonatomic是啥意思,
  • wandd
  • wandd
  • 2014年10月24日 17:17
  • 508
内容举报
返回顶部
收藏助手
不良信息举报
您举报文章:Atomic Functions
举报原因:
原因补充:

(最多只允许输入30个字)