CUDA 原子操作#

1
#include <stdio.h>
2
#include <stdlib.h>
3
#include <math.h>
4
#include <cuda_runtime.h>
5
#define N 12345
6
#define BLOCK 256
7
__global__ void cuda_sum_comp(float *d_data, float *d_data_sum, int n)
8
{
9
  int i1 = threadIdx.x + blockIdx.x * blockDim.x;
10
  if(i1<n) d_data_sum[0] += d_data[i1];
11
}
12

13
__global__ void cuda_sum_atom(float *d_data, float *d_data_sum, int n)
14
{
15
  int i1 = threadIdx.x + blockIdx.x * blockDim.x;
16
  if(i1<n) atomicAdd(&d_data_sum[0], d_data[i1]);
17
}
18

19
int main(void)
20
{
21
  float *data, *data_sum;
22
  float *d_data, *d_data_sum;
23
  data = (float*)malloc(N*sizeof(float)); memset(data,0,N*sizeof(float));
24
  cudaMalloc(&d_data, N*sizeof(float)); cudaMemset(d_data, 0, N*sizeof(float));
25
  data_sum = (float*)malloc(sizeof(float)); memset(data_sum,0,sizeof(float));
26
  cudaMalloc(&d_data_sum, sizeof(float)); cudaMemset(d_data_sum, 0, sizeof(float));
27
  for(int i=0; i<N; i++) data[i] = i-5432;
28
  cudaMemcpy(d_data, data, N*sizeof(float), cudaMemcpyHostToDevice);
29
  cuda_sum_comp<<<(N+BLOCK-1)/BLOCK, BLOCK>>>(d_data, d_data_sum, N);
30
  cudaMemcpy(data_sum, d_data_sum, sizeof(float), cudaMemcpyDeviceToHost);
31
  printf("CUDA_compute:sum = %f\n", data_sum[0]);
32
  memset(data_sum,0,sizeof(float));cudaMemset(d_data_sum, 0, sizeof(float));
33
  cuda_sum_atom<<<(N+BLOCK-1)/BLOCK, BLOCK>>>(d_data, d_data_sum, N);
34
  cudaMemcpy(data_sum, d_data_sum, sizeof(float), cudaMemcpyDeviceToHost);
35
  printf("CUDA_atom:sum = %f\n", data_sum[0]);
36
  free(data); free(data_sum);
37
  cudaFree(d_data); cudaFree(d_data_sum);
38
}