wordpress更换域名后网站打不开,宜昌建设银行网站,电商网站商品页的优化目标是什么?,响应式网站底部怎么做本系列视频目的是帮助开发者们一步步地学会利用CUDA编程模型加速GPU应用, 我们的口号是: 让GPU飞起来
本期我介绍了cuda 当中规约算法的一种情况, 也是小何尚职业生涯中的第一道面试题, 计算数组中所有元素的和. CUDA编程模型系列八(原子操作 / 规约 / 向量元素求和) #include…本系列视频目的是帮助开发者们一步步地学会利用CUDA编程模型加速GPU应用, 我们的口号是: 让GPU飞起来
本期我介绍了cuda 当中规约算法的一种情况, 也是小何尚职业生涯中的第一道面试题, 计算数组中所有元素的和. CUDA编程模型系列八(原子操作 / 规约 / 向量元素求和) #include stdio.h
#include math.h#define N 100000000
#define BLOCK_SIZE 256
#define GRID_SIZE 32__managed__ int source[N];
__managed__ int gpu_result[1] {0};// source[N]: 1 2 3 4 ...............N
// cpu: for loop
// gpu: 1 2 3 4 ...............N 0 1 2 3 4 5 6 7 8 9 10 11 12 .... 31
// thread id step 0: tid0:source[0] source[4] - source[0]
// tid1:source[1] source[5] - source[1]
// tid2:source[2] source[6] - source[2]
// tid4:source[4] source[7] - source[3]
// step 1: tid0: source[0] source[2] - source[0]
// tid1: source[1] source[3] - source[1]
//
// step 2: tid0: source[0] source[1] - source[0]
// thread id: blockDim.x * blockIdx.x threadIdx.x step * blockDim.x * GridDim.x
// thread 0: source[0, 8, 16, 24] sum - shared memory
//
//__global__ void sum_gpu(int *in, int count, int *out)
{__shared__ int ken[BLOCK_SIZE];//grid_loopint shared_tmp0;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_SIZE/2; total_threads1; 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}}}int main()
{int cpu_result 0;printf(Init input source[N]\n);for(int i 0; iN; 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; i20; i){gpu_result[0] 0;sum_gpuGRID_SIZE, BLOCK_SIZE(source, N, gpu_result);cudaDeviceSynchronize();}cudaEventRecord(stop_gpu);cudaEventSynchronize(stop_gpu);for(int i 0; iN; i){cpu_result source[i];}cudaEventRecord(stop_cpu);cudaEventSynchronize(stop_cpu);float time_cpu, time_gpu;cudaEventElapsedTime(time_cpu, stop_gpu, stop_cpu);cudaEventElapsedTime(time_gpu, 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);return 0;
}
