关闭

Atomic Functions

标签: cudaatomicatomic functions
286人阅读 评论(0) 收藏 举报
分类:

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



如果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);
}





0
0

查看评论
* 以上用户言论只代表其个人观点,不代表CSDN网站的观点或立场
    个人资料
    • 访问:8444次
    • 积分:175
    • 等级:
    • 排名:千里之外
    • 原创:9篇
    • 转载:6篇
    • 译文:0篇
    • 评论:0条
    文章分类