学校网站开发实际意义网站建设计划方案模板下载

学校网站开发实际意义,网站建设计划方案模板下载,做百度网站排,东莞网络推广平从简单CUDA开始 我们将从一个简单的C程序开始#xff0c;这个程序将两个包含一百万个元素的数组相加。下面是完整的代码。 #include iostream #include math.h// 函数#xff1a;将两个数组的元素相加 void add(int n, float *x, float *y) {for (int i 0…从简单CUDA开始我们将从一个简单的C程序开始这个程序将两个包含一百万个元素的数组相加。下面是完整的代码。#includeiostream#includemath.h// 函数将两个数组的元素相加voidadd(intn,float*x,float*y){for(inti0;in;i)y[i]x[i]y[i];}intmain(void){intN120;// 1M elementsfloat*xnewfloat[N];float*ynewfloat[N];// 在CPU上初始化数组for(inti0;iN;i){x[i]1.0f;y[i]2.0f;}// 在CPU上运行add函数add(N,x,y);// 检查结果中是否有错误floatmaxError0.0f;for(inti0;iN;i)maxErrorfmax(maxError,fabs(y[i]-3.0f));std::coutMax error: maxErrorstd::endl;// 释放内存delete[]x;delete[]y;return0;}这个程序在CPU上初始化两个一百万元素的数组x和y然后调用add函数将它们相加并将结果存回数组y。最后它检查结果是否正确。这是一个典型的内存带宽受限的计算因为每个加法操作都需要两次内存读取和一次内存写入。转换为CUDA内核为了将这个函数转换为可以在GPU上运行的CUDA内核我们需要做的主要改动是添加__global__修饰符。这个修饰符告诉CUDA编译器这是一个我们将在GPU上调用的函数。__global__voidadd(intn,float*x,float*y){for(inti0;in;i)y[i]x[i]y[i];}在CUDA中在CPU上运行的代码称为主机代码(host code)而在GPU上运行的代码称为设备代码(device code)。__global__表示一个函数可以在设备上执行并且可以从主机全局调用。我们通过一种特殊的语法...来从主机调用这个内核。add1,1(N,x,y);现在add函数将在GPU上执行。但我们还需要处理内存。CUDA中的内存分配为了让GPU能够访问数据我们需要在GPU内存中分配数据或者使用一种特殊的CUDA特性称为统一内存(Unified Memory)。统一内存创建了一个托管的内存池CPU和GPU都可以访问。要使用统一内存我们用cudaMallocManaged()替换new用cudaFree()替换delete。#includeiostream#includemath.h// CUDA内核在GPU上执行__global__voidadd(intn,float*x,float*y){for(inti0;in;i)y[i]x[i]y[i];}intmain(void){intN120;float*x,*y;// 使用统一内存分配x和ycudaMallocManaged(x,N*sizeof(float));cudaMallocManaged(y,N*sizeof(float));// 在CPU上初始化数组for(inti0;iN;i){x[i]1.0f;y[i]2.0f;}// 在GPU上运行内核add1,1(N,x,y);// 等待GPU完成计算cudaDeviceSynchronize();// 检查结果floatmaxError0.0f;for(inti0;iN;i)maxErrorfmax(maxError,fabs(y[i]-3.0f));std::coutMax error: maxErrorstd::endl;// 释放内存cudaFree(x);cudaFree(y);return0;}我们还需要在调用内核后添加cudaDeviceSynchronize()。这是因为CUDA内核启动是异步的。主机代码在启动内核后会立即继续执行而不会等待内核完成。cudaDeviceSynchronize()会阻塞主机线程直到所有先前发出的CUDA任务包括内核完成。这确保了我们在CPU上访问结果之前GPU已经完成了计算。现在我们可以编译并运行这个CUDA程序了。# 编译CUDA代码nvcc add.cu -o add_cuda# 运行程序./add_cuda# 输出: Max error: 0.0程序成功运行但速度非常慢。这是因为内核只在单个GPU线程上运行。性能分析为了衡量性能我们可以使用NVIDIA Nsight Systems命令行工具nsys。我们可以创建一个简单的bash脚本nsys_easy来简化这个过程。#!/bin/bashnsys profile -t cuda --statstrue$现在我们可以用这个脚本来分析我们的程序。# 分析单线程CUDA版本./nsys_easy ./add_cuda在NVIDIA T4 GPU上单线程内核的执行时间大约是91.8毫秒。这比CPU版本慢得多。为了进行比较我们可以测量CPU版本的性能。在我的系统上CPU版本大约需要2.5毫秒。GPU版本慢了36倍这是因为我们的CUDA内核只使用了一个GPU线程。引入线程为了利用GPU的并行能力我们需要使用多个线程。我们可以修改内核启动配置使用256个线程。// 启动256个线程add1,256(N,x,y);现在我们需要修改内核代码让每个线程处理一部分数据。我们可以使用内置变量threadIdx.x它提供了当前线程在块内的索引。__global__voidadd(intn,float*x,float*y){intindexthreadIdx.x;intstrideblockDim.x;for(intiindex;in;istride)y[i]x[i]y[i];}在这个修改后的内核中每个线程从threadIdx.x开始以blockDim.x块中的线程总数这里是256为步长处理数组中的元素。这种循环方式被称为grid-stride loop它有几个好处可扩展性无论我们用多少线程启动内核它都能正确工作。高效性它能很好地合并内存访问。再次运行性能分析我们看到执行时间降到了2.05毫秒比单线程版本快了45倍引入块我们已经使用了256个线程但现代GPU可以同时运行成千上万个线程。为了进一步扩展我们可以使用多个线程块(thread blocks)。我们可以修改内核启动配置使用多个块。intblockSize256;intnumBlocks(NblockSize-1)/blockSize;addnumBlocks,blockSize(N,x,y);这里我们计算了需要的块数以确保每个元素都至少被一个线程访问。现在我们需要修改内核使用blockIdx.x当前块在grid中的索引和gridDim.xgrid中的块总数来计算全局索引。__global__voidadd(intn,float*x,float*y){intindexblockIdx.x*blockDim.xthreadIdx.x;intstrideblockDim.x*gridDim.x;for(intiindex;in;istride)y[i]x[i]y[i];}再次运行性能分析执行时间降到了47.5微秒比单块版本快了43倍比单线程版本快了1932倍这个版本的内核实现了265 GB/s的内存带宽达到了NVIDIA T4 GPU峰值带宽(320 GB/s)的80%以上。统一内存预取虽然我们的内核现在非常快但Nsight Systems的分析显示仍然存在一些性能瓶颈。这是因为统一内存的按需页面迁移(on-demand page migration)。当GPU内核首次访问尚未迁移到GPU内存的数据时会发生页面错误(page fault)导致数据从CPU内存迁移到GPU内存。这些迁移会增加内核的执行时间。为了解决这个问题我们可以使用cudaMemPrefetchAsync()在内核启动前将数据**预取(prefetch)**到GPU。// 在内核启动前预取数据到GPUcudaMemPrefetchAsync(x,N*sizeof(float),0,0);cudaMemPrefetchAsync(y,N*sizeof(float),0,0);addnumBlocks,blockSize(N,x,y);cudaDeviceSynchronize();预取操作是异步的并且可以与数据传输和内核执行重叠。通过预取我们告诉CUDA运行时我们即将在GPU上使用这些数据从而避免了运行时的页面错误。总结通过这个简单的向量加法示例我们学习了如何将一个CPU程序转换为一个高度并行的CUDA程序。我们使用了统一内存来简化内存管理并使用grid-stride loop来实现可扩展的并行内核。最后我们使用预取来优化数据传输。以下是不同版本的性能对比版本时间 (ns)相对单线程加速带宽单线程91,811,2061x137 MB/s单块(256线程)2,049,03445x6 GB/s多块47,5201932x265 GB/s下一步这只是CUDA编程的开始。要继续学习我们推荐以下资源CUDA Toolkit文档: Quick Start Guide, Programming Guide, Best Practices Guide后续教程系列: 探索更高级的CUDA主题如共享内存、原子操作和多维索引。NVIDIA DLI课程: Getting Started with Accelerated Computing in Modern CUDA C