CUDA学习笔记08: 原子规约/向量求和
参考资料
CUDA编程模型系列一(核心函数)_哔哩哔哩_bilibili
代码
#include <iostream>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <stdio.h>
#include <math.h>#define N 10000000
#define BLOCK 256
#define GRID_SIZE 32__managed__ int source[N];
__managed__ int gpu_result[1] = { 0 };__global__ void sum_gpu(int* in, int count, int* out)
{__shared__ int ken[BLOCK];//grid_loopint shared_tmp = 0;for (int idx = blockDim.x * blockIdx.x + threadIdx.x; idx < count; idx += blockDim.x * gridDim.x){shared_tmp += in[idx];}ken[threadIdx.x] = shared_tmp;__syncthreads();int tmp = 0;for (int total_threads = BLOCK / 2; total_threads >= 1; total_threads /= 2){if (threadIdx.x < total_threads){tmp = ken[threadIdx.x] + ken[threadIdx.x + total_threads];}__syncthreads();if (threadIdx.x < total_threads){ken[threadIdx.x] = tmp;}}// block_sum -> share memory[0]if (blockIdx.x * blockDim.x < count){if (threadIdx.x == 0){atomicAdd(out, ken[0]);// memory space wmr}}}// 规约
void test01()
{int cpu_result = 0;/* 初始化 */for (int i = 0; i < N; i++) {source[i] = rand() % 10;}cudaEvent_t start, stop_cpu, stop_gpu;cudaEventCreate(&start);cudaEventCreate(&stop_cpu);cudaEventCreate(&stop_gpu);cudaEventRecord(start);cudaEventSynchronize(start);for (int i = 0; i < 20; i++) {gpu_result[0] = 0;sum_gpu<<<GRID_SIZE, BLOCK>>>(source, N, gpu_result);cudaDeviceSynchronize();}cudaEventRecord(stop_gpu);cudaEventSynchronize(stop_gpu);for (int i = 0; i < N; i++) {cpu_result += source[i];}cudaEventRecord(stop_cpu);cudaEventSynchronize(stop_cpu);float time_cpu = 0, time_gpu = 0;cudaEventElapsedTime(&time_cpu, stop_gpu, stop_cpu);cudaEventElapsedTime(&time_cpu, start, stop_gpu);printf("CPU time: %.2f\nGPU time: %.2f\n", time_cpu, time_gpu / 20);printf("Result: %s\nGPU_result: %d;\nCPU_result: %d;\n", (gpu_result[0] == cpu_result) ? "Pass" : "Error", gpu_result[0], cpu_result);
}代码在windows下可以运行.