【参加CUDA线上培训】4原子操作

CUDA的原子操作提供了一种在多线程环境中对共享变量进行安全读写的方法,确保互斥访问。这些操作包括加法、减法、交换、比较和交换等,且有相应的API如atomicAdd和atomicExch。在给定的示例中,使用CUDA原子操作计算一个大数组的元素之和,对比了GPU和CPU的执行时间,结果显示GPU在并行计算上具有显著优势。
摘要由CSDN通过智能技术生成

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
  • 0
    点赞
  • 0
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

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

余额充值