cuda原子操作概念
CUDA的原子操作可以理解为对一个变量进行“读取-修改-写入”这三个操作的一个最小单位的执行过程,这个执行过程不能够再分解为更小的部分,在它执行过程中,不允许其他并行线程对该变量进行读取和写入的操作。基于这个机制,原子操作实现了对在多个线程间共享的变量的互斥保护,确保任何一次对变量的操作的结果的正确性。
原子操作确保了在多个并行线程间共享的内存的读写保护,每次只能有一个线程对该变量进行读写操作,一个线程对该变量操作的时候,其他线程如果也要操作该变量,只能等待前一线程执行完成。原子操作确保了安全,代价是牺牲了性能。
cuda原子操作常用api
int atomicAdd(int* address, int val);
读取位于全局或共享内存中地址 address 的 16 位、32 位或 64 位字 old,计算 (old + val),并将结果存储回同一地址的内存中。这三个操作在一个原子事务中执行。该函数返回old。
int atomicSub(int* address, int val);
读取位于全局或共享内存中地址address的 32 位字 old,计算 (old - val),并将结果存储回同一地址的内存中。 这三个操作在一个原子事务中执行。 该函数返回old。
int atomicExch(int* address, int val);
读取位于全局或共享内存中地址address的 32 位或 64 位字 old 并将 val 存储回同一地址的内存中。 这两个操作在一个原子事务中执行。 该函数返回old。
int atomicMin(int* address, int val);
读取位于全局或共享内存中地址address的 32 位或 64 位字 old,计算 old 和 val 的最小值,并将结果存储回同一地址的内存中。 这三个操作在一个原子事务中执行。 该函数返回old。
int atomicMax(int* address, int val);
读取位于全局或共享内存中地址address的 32 位或 64 位字 old,计算 old 和 val 的最大值,并将结果存储回同一地址的内存中。 这三个操作在一个原子事务中执行。 该函数返回old。
unsigned int atomicInc(unsigned int* address,unsigned int val);
读取位于全局或共享内存中地址address的 32 位字 old,计算 ((old >= val) ? 0 : (old+1)),并将结果存储回同一地址的内存中。 这三个操作在一个原子事务中执行。 该函数返回old。
unsigned int atomicDec(unsigned int* address, unsigned int val);
读取位于全局或共享内存中地址address的 32 位字 old,计算 (((old == 0) || (old > val)) ? val : (old-1) ),并将结果存储回同一个地址的内存。 这三个操作在一个原子事务中执行。 该函数返回old。
unsigned int atomicCAS(unsigned int* address, unsigned int compare, unsigned int val);
读取位于全局或共享内存中地址address的 16 位、32 位或 64 位字 old,计算 (old == compare ? val : old) ,并将结果存储回同一地址的内存中。 这三个操作在一个原子事务中执行。 该函数返回old(Compare And Swap)。
unsigned int atomicAnd(unsigned int* address, unsigned int val);
读取位于全局或共享内存中地址address的 32 位或 64 位字 old,计算 (old & val),并将结果存储回同一地址的内存中。 这三个操作在一个原子事务中执行。 该函数返回old。
cuda原子操作实验
完成下面的一个实例:
给定一个数组A,它包含1000000个int类型的元素,求他所有的元素之和:
输入:A[1000000]
输出:output(A中所有元素之和)
#include<stdio.h>
#include<stdint.h>
#include<time.h> //for time()
#include<stdlib.h> //for srand()/rand()
#include<sys/time.h> //for gettimeofday()/struct timeval
#include"error.cuh"
#define N 10000000
#define BLOCK_SIZE 256
#define BLOCKS 32
__managed__ int source[N]; //input data
__managed__ int final_result[1] = {0}; //scalar output
__global__ void _sum_gpu(int *input, int count, int *output)
{ const int bid = blockIdx.x;
const int tid = threadIdx.x;
__shared__ int sum_per_block[BLOCK_SIZE];
int temp = 0;
for (int idx = threadIdx.x + blockDim.x * blockIdx.x;
idx < count;
idx += gridDim.x * blockDim.x
)
{
temp += input[idx];
// printf("Hello World from block %d and thread %d! line==%d by henry\n", bid, tid, __LINE__);
}
//printf("Hello World from block %d and thread %d! line==%d by henry\n", bid, tid, __LINE__);
sum_per_block[threadIdx.x] = temp; //the per-thread partial sum is temp!
__syncthreads();
//**********shared memory summation stage***********
for (int length = BLOCK_SIZE / 2; length >= 1; length /= 2)
{
int double_kill = -1;
if (threadIdx.x < length)
{
double_kill = sum_per_block[threadIdx.x] + sum_per_block[threadIdx.x + length];
}
__syncthreads(); //why we need two __syncthreads() here, and,
if (threadIdx.x < length)
{
sum_per_block[threadIdx.x] = double_kill;
}
__syncthreads(); //....here ?
} //the per-block partial sum is sum_per_block[0]
if (blockDim.x * blockIdx.x < count) //in case that our users are naughty
{
//the final reduction performed by atomicAdd()
if (threadIdx.x == 0) atomicAdd(output, sum_per_block[0]);
}
}
int _sum_cpu(int *ptr, int count)
{
int sum = 0;
for (int i = 0; i < count; i++)
{
sum += ptr[i];
}
return sum;
}
void _init(int *ptr, int count)
{
uint32_t seed = (uint32_t)time(NULL); //make huan happy
srand(seed); //reseeding the random generator
//filling the buffer with random data
for (int i = 0; i < count; i++) ptr[i] = rand();
}
double get_time()
{
struct timeval tv;
gettimeofday(&tv, NULL);
return ((double)tv.tv_usec * 0.000001 + tv.tv_sec);
}
int main()
{
//**********************************
fprintf(stderr, "filling the buffer with %d elements...\n", N);
_init(source, N);
//**********************************
//Now we are going to kick start your kernel.
cudaDeviceSynchronize(); //steady! ready! go!
fprintf(stderr, "Running on GPU...\n");
double t0 = get_time();
_sum_gpu<<<BLOCKS, BLOCK_SIZE>>>(source, N, final_result);
CHECK(cudaGetLastError()); //checking for launch failures
CHECK(cudaDeviceSynchronize()); //checking for run-time failurs
double t1 = get_time();
int A = final_result[0];
fprintf(stderr, "GPU sum: %u\n", A);
//**********************************
//Now we are going to exercise your CPU...
fprintf(stderr, "Running on CPU...\n");
double t2 = get_time();
int B = _sum_cpu(source, N);
double t3 = get_time();
fprintf(stderr, "CPU sum: %u\n", B);
//******The last judgement**********
if (A == B)
{
fprintf(stderr, "Test Passed!\n");
}
else
{
fprintf(stderr, "Test failed!\n");
exit(-1);
}
//****and some timing details*******
fprintf(stderr, "GPU time %.3f ms\n", (t1 - t0) * 1000.0);
fprintf(stderr, "CPU time %.3f ms\n", (t3 - t2) * 1000.0);
return 0;
}
编译和运行效果
root@ubuntu-desktop:/home/henry# nvprof ./e.out
filling the buffer with 10000000 elements...
==17659== NVPROF is profiling process 17659, command: ./e.out
==17659== Warning: Unified Memory Profiling is not supported on the underlying platform. System requirements for unified memory can be found at: http://docs.nvs
Running on GPU...
GPU sum: 2581111373
Running on CPU...
CPU sum: 2581111373
Test Passed!
GPU time 9.753 ms
CPU time 105.469 ms
==17659== Profiling application: ./e.out
==17659== Profiling result:
Type Time(%) Time Calls Avg Min Max Name
GPU activities: 100.00% 7.8306ms 1 7.8306ms 7.8306ms 7.8306ms _sum_gpu(int*, int, int*)
API calls: 97.25% 370.36ms 1 370.36ms 370.36ms 370.36ms cuDevicePrimaryCtxRetain
2.32% 8.8354ms 2 4.4177ms 88.223us 8.7471ms cudaDeviceSynchronize
0.24% 897.05us 1 897.05us 897.05us 897.05us cudaLaunchKernel
0.10% 395.16us 2 197.58us 11.904us 383.26us cuModuleGetGlobal
0.06% 246.08us 97 2.5360us 992ns 42.400us cuDeviceGetAttribute
0.01% 36.799us 1 36.799us 36.799us 36.799us cuModuleGetFunction
0.00% 15.168us 1 15.168us 15.168us 15.168us cuDeviceTotalMem
0.00% 9.4400us 3 3.1460us 1.8240us 5.4400us cuDeviceGetCount
0.00% 8.1280us 1 8.1280us 8.1280us 8.1280us cuCtxSetCurrent
0.00% 3.7120us 2 1.8560us 1.5680us 2.1440us cuCtxGetCurrent
0.00% 3.3920us 2 1.6960us 1.2480us 2.1440us cuDeviceGet
0.00% 2.3040us 1 2.3040us 2.3040us 2.3040us cudaGetLastError
0.00% 2.0480us 1 2.0480us 2.0480us 2.0480us cuDeviceGetName
0.00% 1.6000us 1 1.6000us 1.6000us 1.6000us cuCtxGetDevice
0.00% 1.2800us 1 1.2800us 1.2800us 1.2800us cuDeviceGetUuid