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

在linux2.6.18之后,删除了和头文件,编译器提供内建(built-in)原子操作函数。需要在gcc编译选项中指明CPU类型。如gcc -marth=i686 -o hello hello.c。...
  • vinowan
  • vinowan
  • 2011年08月06日 17:26
  • 518

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

利用atomic递增 Channels Unbuffered channels

Atomic RAR Password Recovery RAR解密用

  • 2013年05月22日 15:40
  • 2.98MB
  • 下载

CUDA atomic原子操作

和许多多线程并行问题一样,CUDA也存在互斥访问的问题,即当一个线程改变变量X,而另外一个线程在读取变量X的值,执行原子操作类似于有一个自旋锁,只有等X的变量在改变完成之后,才能执行读操作,这样可以保...

C++11多线程(十):atomic原子操作的高效率(实例)

参考链接:http://blog.csdn.net/yockie/article/details/8838686 目录 1.原子操作介绍 2.示例比较:不用锁及原子(结果错误) 3.示例比较:用锁(...

CE Moore Atomic Energy Levels Vol 1

  • 2010年05月07日 13:14
  • 3.48MB
  • 下载

Atomic Scala, 2nd Ed

  • 2015年12月20日 09:57
  • 3.54MB
  • 下载
内容举报
返回顶部
收藏助手
不良信息举报
您举报文章:Atomic Functions
举报原因:
原因补充:

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