做网站要具备些什么条件wordpress自定义文章页面模板

做网站要具备些什么条件,wordpress自定义文章页面模板,wordpress小程序推荐,网站的图片尺寸文为CUDA并行规约系列文章的下篇#xff0c;本文介绍了5种并行规约算法的实现#xff0c;并从硬件的角度对它们进行分析和优化#xff0c;最终给出一个开箱即用的模板代码及其使用示例。勘误首先是一个勘误#xff0c;在上篇中存在一个拼写错误#xff0c;线程束的正确单词…文为CUDA并行规约系列文章的下篇本文介绍了5种并行规约算法的实现并从硬件的角度对它们进行分析和优化最终给出一个开箱即用的模板代码及其使用示例。勘误首先是一个勘误在上篇中存在一个拼写错误线程束的正确单词是Warp而不是Wrap非常抱歉给读者朋友们造成了误解。写在前面这是本系列文章的下篇。上篇介绍了一些CUDA并行规约优化涉及到的GPU硬件知识并给出了两种并不完美的实现。本文将接着介绍剩下的五种实现并给出一个开箱即用的CUDA并行规约模板。算法具体实现下篇V2: Sequential Addressing先简单回顾一下在上篇的最后我们发现V1版本的实现存在Bank Conflict的问题具体表现为当时会访问会访问会访问.....造成一个Warp里所有线程都同时访问Bank 0导致这些访问被串行化严重影响性能。造成这一问题的根本原因是同一个Warp里的线程它们访问的地址存在可变的StrideBank机制设计之初的预期就是一个Warp里32个线程访问连续的32个地址而不是访问分散在各处的地址。明确了这一点之后优化方法就很明确了我们只需要让每个线程都负责其threadIdx对应的内存地址的规约即可比如就只管就只管这样就不会出现Bank Conflict了。由于要防止Warp Divergence所以第一轮循环只有前一半线程在工作以为例第一轮循环只有到在工作那么很显然的就需要规约和需要规约和以此类推。这个过程如下图所示在代码实现上也不难只需要做如下改变即可一个容易混淆的点可能会有读者感到疑惑总会有一次循环会访问和这怎么能算解决了Bank Conflict呢这里需要明确的一点是Bank Conflict是发生在线程间的而不是一个线程内的指令间的。从指令执行的角度来分析就以为例这条语句会分4步完成分别是读取读取计算写入。对而言读取和读取一定是先后发生的不会并行所以也就不存在Bank Conflict。这里我们再从线程间执行的角度来看看在读取时在读取这不会发生冲突。但是如果是V1的实现在读取时可能会同时在读取这就导致了Bank Conflict。优化效果根据NVIDIA的数据这一优化将性能提升了接近1倍也是相当可观的。V3: First Add During LoadV2版本已经把硬件上踩过的坑基本都填完了接下来就是一些细节上的优化。其中有一个思路就是在把数据加载到共享内存的阶段做一些预处理。比如我们可以先把和规约了然后存储到中把和规约了存储到中。这样就只需要启动一半的线程以几乎一半的时间执行完所有的任务。这一想法还可以进一步推广如果在加载阶段就提前规约4个元素那就可以把时间压缩到。这一思想就是所谓的算法级联即结合并行执行和串行执行的策略。这部分更详细的会在后面V6版本里分析。由于这里NVIDIA的实现和我们的问题定义有冲突这里就不详细展开了仅贴出NVIDIA的实现供参考。实测的数据也验证了这一理论的正确性时间确实缩短了接近一半。V4: Unroll the Last WarpV3版本计算出的有效带宽为17GB/s远没有达到硬件的上限所以有理由怀疑这里还存在指令执行上的瓶颈。观察之后发现这里的部分循环指令也许可以优化掉。在上篇中我们提到过在同一个Warp里指令执行可以认为是同步的所以我们可以在最后32个线程工作时去掉__syncthreads。此外既然都已经知道这里是最后32个线程在工作了那循环也不必需要了可以直接硬编码写......最终的修改如下图所示为什么不需要线程同步了因为这里可以保证进入warpReduce的线程是到即一个Warp内的线程。根据之前的内容在执行函数第二行时是一定也在执行第二行换言之是一定执行完成了第一行的这就保证了在执行取这个指令时存储的一定是规约了和之后的值。那么还有一个疑问这里写入了新的的值是这不就破坏了整个规约过程了吗事实上这里写入了什么并不重要这是基于两点事实首先写入和写入一定是在同一个时钟周期内发生的所以在一开始执行第二行时读取到的一定是没被破坏的这确保了执行的正确性其次在执行完第一行之后实际上只有到真正起作用了后面16个线程只要不乱写数据妨碍到前16个线程的工作无论做什么都对结果没影响。所以理论上讲这里加个判断只让的线程执行最终的结果都是正确的。这里没有这么做主要是为了防止Warp Divergence。优化效果实验数据表明仅仅是对最后一个Warp做了循环展开最终性能就优化了接近一倍。V5: Completely Unrolled既然只对最后一个Warp进行循环展开优化就这么明显那能不能再激进一些把所有循环都丢掉呢答案是Yes。因为一个线程块内的线程数量上限是1024并且我们要求线程数量是2的整数次幂所以我们可以枚举所有可能的线程数量然后仿照之前展开最后一个循环的方法针对每种情况直接写出对应的展开后的代码。如果只是在参数列表里加一个blockSize那么在运行时还是会经过多个无意义的if-else影响执行效率。这里可以借助C的模板机制进一步地提升执行性能。具体的修改如下图所示这里的blockSize是在调用时使用模板传递的这里面的if会在编译时就进行优化。编译器会使用Dead Code Elimination根据blockSize删掉不可达的代码所以最终编译出的二进制里面是不会有这些红色的if的只有一串顺序执行的指令。但是模板要求我们在编译阶段就确定blockSize在实际实现时是比较困难的这个该如何解决呢答案是在调用时使用switch来枚举简单粗暴如下图所示优化效果V6: Multiple Adds理论分析这里我们要从成本的角度来看待一下现有的方案。和我们租用算力服务器一样成本可以用使用时间 * GPU核心数量来衡量这里GPU核心数量可以认为是线程数量两者之间只有常数倍的差别。我们假设这样一个场景来计算成本只有1个Batch共有个元素启动个线程来处理在V5实现中这里使用Brent定理其中为算法的估计实际执行时间表示算法的总工作量即一共执行多少次运算表示线程数量表示算法在最理想的情况下算法执行的最短用时很显然这里由于最少也需要执行步规约所以一通计算可以得到这里具体计算过程见NVIDIA的PPT这里不再赘述。带入公式可得那么可以计算得到V5版本算法的成本为。但是如果用一个线程串行处理所有数据成本却只有也就是说我们的并行算法的成本比纯串行处理还要高。那么怎么把这个成本降下来呢这里成本变高的本质是线程数量过多使得每个线程的工作量太少导致了整体成本的增大。那么相对应的我们可以通过减少线程数量来实现降本增效。具体来说我们可以使用V3中提到的方法即在加载数据的阶段就提前做几次规约。那么具体应该做多少次呢这里NVIDIA的PPT里给出的数据是应该做次如此优化之后最终的成本能降到可以和串行执行持平。这里把串行和并行搭配使用的策略就是算法级联。实现方式这里NVIDIA的实现和我们的问题定义相冲突所以这里不再赘述了后面在开箱即用的部分会解释我们是如何处理这一冲突的。下图是V5到V6版本的变化以及V6的完整代码实现。优化效果这里贴出7个版本的优化效果数据表格还能再优化吗NVIDIA官方的PPT到这里就结束了。我们可以思考一下V6版本还能进一步优化吗也许还可以比如用循环展开的技巧和最开始的while循环展开一下这个就作为open issue供大家探讨了。开箱即用的实现数据要求输入矩阵按行优先展开成的一个向量输出维向量要求线程块内线程的数量block_size并且为2的整数次幂要求一定要能整除block_sizeNVIDIA官方的实现里似乎没有batch的概念感觉是想要把输入的所有数据都规约到一个值因此在V3和V6里面会有跨Block规约的情况。我们这里就不采用这个方案而是把分为了若干个fold最终的线程数量就是在一开始加载数据时就提前规约fold_num次。这个fold_num由调用方传递。对于一个batch数量大于1024的情况和数据数量不是2的整数次幂的情况则需要调用CUDA的上层框架做处理了这里暂时不考虑这些case。代码实现这一实现使用了C的模板特性支持调用者自行选择数据类型和规约函数。templatetypename Tusing ReduceFunc T(*)(T, T);templatetypename T, ReduceFuncT reduce_func, size_t block_size__device__ void reduceWarp(volatile T *shared_memory, size_t tid) {if (block_size 64) shared_memory[tid] reduce_func(shared_memory[tid], shared_memory[tid 32]);if (block_size 32) shared_memory[tid] reduce_func(shared_memory[tid], shared_memory[tid 16]);if (block_size 16) shared_memory[tid] reduce_func(shared_memory[tid], shared_memory[tid 8]);if (block_size 8) shared_memory[tid] reduce_func(shared_memory[tid], shared_memory[tid 4]);if (block_size 4) shared_memory[tid] reduce_func(shared_memory[tid], shared_memory[tid 2]);if (block_size 2) shared_memory[tid] reduce_func(shared_memory[tid], shared_memory[tid 1]);}