A Simple CUDA Renderer
前置条件：
如果是在自己的电脑上做，你可能需要安装一些额外的东西
尽量确保自己有支持64位的gcc
如果没有表现为编译时报-m64参数不可用，但似乎粗暴的改为-m32就可正常运行
windows系统下：
放弃吧，做不了的，换平台才是王道
配置好之后
如果你可以查看版本号，但是在第一题运行make文件时出问题了，请再回到此处
先按这个解决nvcc fatal : Cannot find compiler ‘cl.exe‘ in PATH解决方法_w36680130的博客-CSDN博客
如果和我一样根本没有那个文件的话
前往Thank You for Downloading Visual Studio Community Edition (microsoft.com)安装
只需选择c/c++桌面版开发即可，可以依据自己的需求再选一点
总体目标：
在 CUDA 中编写一个绘制彩色圆圈的并行渲染器。虽然此渲染器非常简单，但并行化渲染器需要我们设计和实现可以有效地并行构造和操作的数据结构。
Part1 CUDA的简单应用
原题翻译：
Homework 1 您的预热任务是在 CUDA 中重新实现业 1 中的 SAXPY 函数。作为编写 CUDA 程序的练习，您需要在这部分作业中了解cuda工作的基本原理。这部分作业的起始代码位于作业存储库的 /saxpy 目录中。您可以通过在 /saxpy 目录中调用 make 和 ./cudaSaxpy 来构建和运行 saxpy CUDA 程序。
1.1 请在 saxpy.cu 年的函数 saxpyCuda 中完成 SAXPY 的实现。
在执行计算之前，您需要分配设备全局内存阵列此处应当理解为GPU的内存并将主机输入数组 X、Y 和结果的内容复制到 CUDA 设备内存中。CUDA 计算完成后，必须将结果复制回主机内存。
1.2 在 saxpyCuda 中的 CUDA 内核调用周围添加计时器。添加后，程序应可测量传输数据所用的时间。
提供的初始代码包含一组计时器，用于测量将数据复制到 GPU、运行内核以及将数据复制回 CPU 的整个过程。
插入另一组仅测量运行内核所花费的时间的计时器。它们不应包括 CPU 到 GPU 数据传输或将结果从 GPU 传输回 CPU 的时间。
Hint 1.2 在测量GPU内核运行时间时，需要注意：CUDA 内核在 GPU 上的执行与在 CPU 上运行的主应用程序线程是异步的。
因此，我们需要在内核调用之后调用 cudaDeviceSynchronize() 以等待 GPU 上的所有 CUDA 工作完成。cudaDeviceSync()将在 GPU 上所有之前的 CUDA 工作都完成后返回。
请注意，在 cudaMemcpy（） 之后不需要 cudaDeviceSynchronize（） 来确保内存传输到 GPU 的完成，因为 cudaMempy（） 在我们使用它的条件下是同步的。
<span style="background-color:#f8f8f8"><span style="color:#333333"> <span style="color:#008855">double</span> <span style="color:#000000">startTime</span> <span style="color:#981a1a">=</span> <span style="color:#000000">CycleTimer</span>::<span style="color:#000000">currentSeconds</span>();
<span style="color:#000000">saxpy_kernel</span><span style="color:#981a1a"><<<</span><span style="color:#000000">blocks</span>, <span style="color:#000000">threadsPerBlock</span><span style="color:#981a1a">>>></span>(<span style="color:#000000">N</span>, <span style="color:#000000">alpha</span>, <span style="color:#000000">device_x</span>, <span style="color:#000000">device_y</span>, <span style="color:#000000">device_result</span>);
<span style="color:#008855">double</span> <span style="color:#000000">endTime</span> <span style="color:#981a1a">=</span> <span style="color:#000000">CycleTimer</span>::<span style="color:#000000">currentSeconds</span>();
<span style="color:#aa5500">// 此时我们将得到一个非常短的运行时间</span>
<span style="color:#aa5500">// 原因在于 endTime的计算和GPU的执行是异步进行的</span></span></span><span style="background-color:#f8f8f8"><span style="color:#333333"> <span style="color:#008855">double</span> <span style="color:#000000">startTime</span> <span style="color:#981a1a">=</span> <span style="color:#000000">CycleTimer</span>::<span style="color:#000000">currentSeconds</span>();
<span style="color:#000000">saxpy_kernel</span><span style="color:#981a1a"><<<</span><span style="color:#000000">blocks</span>, <span style="color:#000000">threadsPerBlock</span><span style="color:#981a1a">>>></span>(<span style="color:#000000">N</span>, <span style="color:#000000">alpha</span>, <span style="color:#000000">device_x</span>, <span style="color:#000000">device_y</span>, <span style="color:#000000">device_result</span>);
<span style="color:#000000">cudaDeviceSynchronize</span>();
<span style="color:#008855">double</span> <span style="color:#000000">endTime</span> <span style="color:#981a1a">=</span> <span style="color:#000000">CycleTimer</span>::<span style="color:#000000">currentSeconds</span>();
<span style="color:#aa5500">// 此时我们将正确测量cuda内核所用的执行时间</span></span></span>问题1.与基于 CPU 的顺序 SAXPY 实施相比，您观察到的性能如何（回想一下作业 1 中程序 5 中 saxpy 的结果）？
问题2.比较并解释两组计时器提供的结果之间的差异（仅对内核执行进行计时，而不是对将数据移动到 GPU 并移回内核执行的整个过程进行计时）。观察到的带宽值是否与机器不同组件报告的可用带宽大致一致？
Hint Question 2 有几个因素会影响峰值带宽，包括 CPU 主板芯片组性能以及用作传输源的主机 CPU 内存是否“固定”DMA技术——允许 GPU 直接访问内存，而无需经过虚拟内存地址转换。这使得期望值可能更高或更低
关于DMA技术，说一点题外话，我了解比较多的是DMA用于游戏外挂研发（多有趣啊，直接读取内存，可以绕过许多基于进程、地址保护的反作弊手段QAQ
预备知识：
CUDA运行于GPU端，我们一般称呼CPU端为HOST端，GPU端为Device端，简单考虑CUDA工作情况时，我们认为CUDA的任务由CPU分配，其指定需要在GPU上启动的“线程块”数和线程结构，并指定每个线程的工作，此时GPU与CPU是异步工作的。
线程块： GPU中的线程并不是和CPU中一样独立启动的，其启动的最小单元是线程块，一个线程块中往往含有多个线程，即使本次任务无需这么多线程，也最少会分配一个线程块供你使用。且线程块内的线程组成可以是一维、二维、三维的，同一个线程块内的线程可以通讯（基于共享内存
Grid：多个线程块组成Grid，块的组织方式可以是一维、二维、三维的。
Warp：线程束，其是逻辑上的最小执行单元，往往有32个并行线程组成，这些并行线程使用不同的数据资源
SM：一个SM由多个CUDAcore组成，每个SM根据GPU架构不同有不同数量的CUDA core，Pascal架构中一个SM有128个CUDA core。 SM还包括特殊运算单元(SFU)，共享内存(shared memory)，寄存器文件(Register File)和调度器(Warp Scheduler)等。register和shared memory是稀缺资源，这些有限的资源就使每个SM中active warps有非常严格的限制，也就限制了并行能力。
当一个kernel被执行时，grid中的线程块被分配到SM上，一个线程块的thread只能在一个SM上调度，SM一般可以调度多个线程块，大量的thread可能被分到不同的SM上。每个thread拥有它自己的程序计数器和状态寄存器，并且用该线程自己的数据执行指令，这就是所谓的Single Instruction Multiple Thread(SIMT)。
一个CUDA core可以执行一个thread，一个SM的CUDA core会分成几个warp（即CUDA core在SM中分组)，由warp scheduler负责调度。尽管warp中的线程从同一程序地址，但可能具有不同的行为，比如分支结构，因为GPU规定warp中所有线程在同一周期执行相同的指令，warp发散会导致性能下降。一个SM同时并发的warp是有限的，因为资源限制，SM要为每个线程块分配共享内存，而也要为每个线程束中的线程分配独立的寄存器，所以SM的配置会影响其所支持的线程块和warp并发数量。
CUDA中常用的功能：
详见ppt，内存复制、内存申请、并行控制、计时的一组函数
实现过程：
我们需要实现的部分只有在GPU上申请空间、复制到GPU、复制回GPU、计时
主要工作都在saxpy.cu中
<span style="background-color:#f8f8f8"><span style="color:#333333"> <span style="color:#555555">#include <stdio.h></span>
<span style="color:#555555">#include <cuda.h></span>
<span style="color:#555555">#include <cuda_runtime.h></span>
<span style="color:#555555">#include <driver_functions.h></span>
<span style="color:#555555">#include "CycleTimer.h"</span>
<span style="color:#aa5500">// return GB/sec</span>
<span style="color:#aa5500">// 计算理论速率</span>
<span style="color:#008855">float</span> <span style="color:#0000ff">GBPerSec</span>(<span style="color:#008855">int</span> <span style="color:#000000">bytes</span>, <span style="color:#008855">float</span> <span style="color:#000000">sec</span>) {
<span style="color:#770088">return</span> <span style="color:#000000">static_cast</span><span style="color:#981a1a"><</span><span style="color:#008855">float</span><span style="color:#981a1a">></span>(<span style="color:#000000">bytes</span>) <span style="color:#981a1a">/</span> (<span style="color:#116644">1024.</span> <span style="color:#981a1a">*</span> <span style="color:#116644">1024.</span> <span style="color:#981a1a">*</span> <span style="color:#116644">1024.</span>) <span style="color:#981a1a">/</span> <span style="color:#000000">sec</span>;
}
<span style="color:#aa5500">// This is the CUDA "kernel" function that is run on the GPU. You</span>
<span style="color:#aa5500">// know this because it is marked as a __global__ function.</span>
<span style="color:#aa5500">// 运行在GPU上的函数，可以通过__global__来确认</span>
<span style="color:#000000">__global__</span> <span style="color:#008855">void</span>
<span style="color:#0000ff">saxpy_kernel</span>(<span style="color:#008855">int</span> <span style="color:#000000">N</span>, <span style="color:#008855">float</span> <span style="color:#000000">alpha</span>, <span style="color:#008855">float*</span> <span style="color:#000000">x</span>, <span style="color:#008855">float*</span> <span style="color:#000000">y</span>, <span style="color:#008855">float*</span> <span style="color:#000000">result</span>) {
<span style="color:#aa5500">// compute overall thread index from position of thread in current</span>
<span style="color:#aa5500">// block, and given the block we are in (in this example only a 1D</span>
<span style="color:#aa5500">// calculation is needed so the code only looks at the .x terms of</span>
<span style="color:#aa5500">// blockDim and threadIdx.</span>
<span style="color:#aa5500">//根据当前线程的位置计算整体线程索引块，并计算我们所在的块内位置（在本例中为一维</span>
<span style="color:#aa5500">//因此代码仅查看 .x 项blockDim 和 threadIdx。</span>
<span style="color:#008855">int</span> <span style="color:#000000">index</span> <span style="color:#981a1a">=</span> <span style="color:#000000">blockIdx</span>.<span style="color:#000000">x</span> <span style="color:#981a1a">*</span> <span style="color:#000000">blockDim</span>.<span style="color:#000000">x</span> <span style="color:#981a1a">+</span> <span style="color:#000000">threadIdx</span>.<span style="color:#000000">x</span>;
<span style="color:#aa5500">// this check is necessary to make the code work for values of N</span>
<span style="color:#aa5500">// that are not a multiple of the thread block size (blockDim.x)</span>
<span style="color:#aa5500">// 检查，使得代码对任意N均生效，即使得多分配的线程不进行任何工作</span>
<span style="color:#770088">if</span> (<span style="color:#000000">index</span> <span style="color:#981a1a"><</span> <span style="color:#000000">N</span>)
<span style="color:#000000">result</span>[<span style="color:#000000">index</span>] <span style="color:#981a1a">=</span> <span style="color:#000000">alpha</span> <span style="color:#981a1a">*</span> <span style="color:#000000">x</span>[<span style="color:#000000">index</span>] <span style="color:#981a1a">+</span> <span style="color:#000000">y</span>[<span style="color:#000000">index</span>];
}
<span style="color:#aa5500">// saxpyCuda --</span>
<span style="color:#aa5500">//</span>
<span style="color:#aa5500">// This function is regular C code running on the CPU. It allocates</span>
<span style="color:#aa5500">// memory on the GPU using CUDA API functions, uses CUDA API functions</span>
<span style="color:#aa5500">// to transfer data from the CPU's memory address space to GPU memory</span>
<span style="color:#aa5500">// address space, and launches the CUDA kernel function on the GPU.</span>
<span style="color:#aa5500">// 本函数要求实现一个简单的cuda程序用于加速一组运算</span>
<span style="color:#aa5500">// 输入： N 数组长度 、 alpha 系数、 xarray 运算数1 、 yarray 运算数2</span>
<span style="color:#aa5500">// 输出 resultarray 结果数组</span>
<span style="color:#aa5500">// 思路：</span>
<span style="color:#aa5500">// 1. 使用cudamalloc申请合适大小的空间</span>
<span style="color:#aa5500">// 2. 使用cudaMemcpy将CPU中的数据移到GPU</span>
<span style="color:#aa5500">// 3. 在GPU上进行计算</span>
<span style="color:#aa5500">// 4. 使用cudaMemcpy将GPU中的数据移到CPU</span>
<span style="color:#008855">void</span> <span style="color:#0000ff">saxpyCuda</span>(<span style="color:#008855">int</span> <span style="color:#000000">N</span>, <span style="color:#008855">float</span> <span style="color:#000000">alpha</span>, <span style="color:#008855">float*</span> <span style="color:#000000">xarray</span>, <span style="color:#008855">float*</span> <span style="color:#000000">yarray</span>, <span style="color:#008855">float*</span> <span style="color:#000000">resultarray</span>) {
<span style="color:#aa5500">// must read both input arrays (xarray and yarray) and write to</span>
<span style="color:#aa5500">// output array (resultarray)</span>
<span style="color:#008855">int</span> <span style="color:#000000">totalBytes</span> <span style="color:#981a1a">=</span> <span style="color:#770088">sizeof</span>(<span style="color:#008855">float</span>) <span style="color:#981a1a">*</span> <span style="color:#116644">3</span> <span style="color:#981a1a">*</span> <span style="color:#000000">N</span>;
<span style="color:#aa5500">// compute number of blocks and threads per block. In this</span>
<span style="color:#aa5500">// application we've hardcoded thread blocks to contain 512 CUDA</span>
<span style="color:#aa5500">// threads.</span>
<span style="color:#aa5500">// 计算每个块的线程数，此处已经通过硬件信息得知为512</span>
<span style="color:#aa5500">// 在我们的机器上可能不是512，需要修改</span>
<span style="color:#770088">const</span> <span style="color:#008855">int</span> <span style="color:#000000">threadsPerBlock</span> <span style="color:#981a1a">=</span> <span style="color:#116644">512</span>;
<span style="color:#aa5500">// Notice the round up here. The code needs to compute the number</span>
<span style="color:#aa5500">// of threads blocks needed such that there is one thread per</span>
<span style="color:#aa5500">// element of the arrays. This code is written to work for values</span>
<span style="color:#aa5500">// of N that are not multiples of threadPerBlock.</span>
<span style="color:#aa5500">//代码计算运行所需的线程块数，以便每个数字有一个线程进行计算。 </span>
<span style="color:#aa5500">//编写此代码是为了处理 N 不是 threadPerBlock 的倍数。</span>
<span style="color:#770088">const</span> <span style="color:#008855">int</span> <span style="color:#000000">blocks</span> <span style="color:#981a1a">=</span> (<span style="color:#000000">N</span> <span style="color:#981a1a">+</span> <span style="color:#000000">threadsPerBlock</span> <span style="color:#981a1a">-</span> <span style="color:#116644">1</span>) <span style="color:#981a1a">/</span> <span style="color:#000000">threadsPerBlock</span>;
<span style="color:#aa5500">// These are pointers that will be pointers to memory allocated</span>
<span style="color:#aa5500">// *one the GPU*. You should allocate these pointers via</span>
<span style="color:#aa5500">// cudaMalloc. You can access the resulting buffers from CUDA</span>
<span style="color:#aa5500">// device kernel code (see the kernel function saxpy_kernel()</span>
<span style="color:#aa5500">// above) but you cannot access the contents these buffers from</span>
<span style="color:#aa5500">// this thread. CPU threads cannot issue loads and stores from GPU</span>
<span style="color:#aa5500">// memory!</span>
<span style="color:#aa5500">//这些指针将是指向分配的内存的指针</span>
<span style="color:#aa5500">//需要注意，CPU和GPU之间的指针不互通</span>
<span style="color:#008855">float*</span> <span style="color:#000000">device_x</span> <span style="color:#981a1a">=</span> <span style="color:#000000">nullptr</span>;
<span style="color:#008855">float*</span> <span style="color:#000000">device_y</span> <span style="color:#981a1a">=</span> <span style="color:#000000">nullptr</span>;
<span style="color:#008855">float*</span> <span style="color:#000000">device_result</span> <span style="color:#981a1a">=</span> <span style="color:#000000">nullptr</span>;
<span style="color:#aa5500">//</span>
<span style="color:#aa5500">// CS149 TODO: allocate device memory buffers on the GPU using cudaMalloc.</span>
<span style="color:#aa5500">//</span>
<span style="color:#aa5500">// We highly recommend taking a look at NVIDIA's</span>
<span style="color:#aa5500">// tutorial, which clearly walks you through the few lines of code</span>
<span style="color:#aa5500">// you need to write for this part of the assignment:</span>
<span style="color:#aa5500">//</span>
<span style="color:#aa5500">// https://devblogs.nvidia.com/easy-introduction-cuda-c-and-c/</span>
<span style="color:#aa5500">//</span>
<span style="color:#000000">cudaMalloc</span>(<span style="color:#981a1a">&</span><span style="color:#000000">device_x</span>, <span style="color:#000000">N</span><span style="color:#981a1a">*</span><span style="color:#770088">sizeof</span>(<span style="color:#008855">float</span>));
<span style="color:#000000">cudaMalloc</span>(<span style="color:#981a1a">&</span><span style="color:#000000">device_y</span>, <span style="color:#000000">N</span><span style="color:#981a1a">*</span><span style="color:#770088">sizeof</span>(<span style="color:#008855">float</span>));
<span style="color:#000000">cudaMalloc</span>(<span style="color:#981a1a">&</span><span style="color:#000000">device_result</span>, <span style="color:#000000">N</span><span style="color:#981a1a">*</span><span style="color:#770088">sizeof</span>(<span style="color:#008855">float</span>));
<span style="color:#aa5500">// start timing after allocation of device memory</span>
<span style="color:#008855">double</span> <span style="color:#000000">startTime</span> <span style="color:#981a1a">=</span> <span style="color:#000000">CycleTimer</span>::<span style="color:#000000">currentSeconds</span>();
<span style="color:#aa5500">//</span>
<span style="color:#aa5500">// CS149 TODO: copy input arrays to the GPU using cudaMemcpy</span>
<span style="color:#aa5500">//</span>
<span style="color:#000000">cudaMemcpy</span>(<span style="color:#000000">device_x</span>, <span style="color:#000000">xarray</span>, <span style="color:#000000">N</span><span style="color:#981a1a">*</span><span style="color:#770088">sizeof</span>(<span style="color:#008855">float</span>), <span style="color:#000000">cudaMemcpyHostToDevice</span>);
<span style="color:#000000">cudaMemcpy</span>(<span style="color:#000000">device_y</span>, <span style="color:#000000">yarray</span>, <span style="color:#000000">N</span><span style="color:#981a1a">*</span><span style="color:#770088">sizeof</span>(<span style="color:#008855">float</span>), <span style="color:#000000">cudaMemcpyHostToDevice</span>);
<span style="color:#aa5500">// run CUDA kernel. (notice the <<< >>> brackets indicating a CUDA</span>
<span style="color:#aa5500">// kernel launch) Execution on the GPU occurs here.</span>
<span style="color:#008855">double</span> <span style="color:#000000">startTime_kernel</span> <span style="color:#981a1a">=</span> <span style="color:#000000">CycleTimer</span>::<span style="color:#000000">currentSeconds</span>();
<span style="color:#000000">saxpy_kernel</span><span style="color:#981a1a"><<<</span><span style="color:#000000">blocks</span>, <span style="color:#000000">threadsPerBlock</span><span style="color:#981a1a">>>></span>(<span style="color:#000000">N</span>, <span style="color:#000000">alpha</span>, <span style="color:#000000">device_x</span>, <span style="color:#000000">device_y</span>, <span style="color:#000000">device_result</span>);
<span style="color:#000000">cudaDeviceSynchronize</span>();
<span style="color:#008855">double</span> <span style="color:#000000">endTime_kernel</span> <span style="color:#981a1a">=</span> <span style="color:#000000">CycleTimer</span>::<span style="color:#000000">currentSeconds</span>();
<span style="color:#aa5500">//</span>
<span style="color:#aa5500">// CS149 TODO: copy result from GPU back to CPU using cudaMemcpy</span>
<span style="color:#aa5500">//</span>
<span style="color:#000000">cudaMemcpy</span>(<span style="color:#000000">resultarray</span>, <span style="color:#000000">device_result</span>, <span style="color:#000000">N</span><span style="color:#981a1a">*</span><span style="color:#770088">sizeof</span>(<span style="color:#008855">float</span>), <span style="color:#000000">cudaMemcpyDeviceToHost</span>);
<span style="color:#aa5500">// end timing after result has been copied back into host memory</span>
<span style="color:#008855">double</span> <span style="color:#000000">endTime</span> <span style="color:#981a1a">=</span> <span style="color:#000000">CycleTimer</span>::<span style="color:#000000">currentSeconds</span>();
<span style="color:#000000">cudaError_t</span> <span style="color:#000000">errCode</span> <span style="color:#981a1a">=</span> <span style="color:#000000">cudaPeekAtLastError</span>();
<span style="color:#770088">if</span> (<span style="color:#000000">errCode</span> <span style="color:#981a1a">!=</span> <span style="color:#000000">cudaSuccess</span>) {
<span style="color:#000000">fprintf</span>(<span style="color:#000000">stderr</span>, <span style="color:#aa1111">"WARNING: A CUDA error occured: code=%d, %s\n"</span>,
<span style="color:#000000">errCode</span>, <span style="color:#000000">cudaGetErrorString</span>(<span style="color:#000000">errCode</span>));
}
<span style="color:#008855">double</span> <span style="color:#000000">overallDuration</span> <span style="color:#981a1a">=</span> <span style="color:#000000">endTime</span> <span style="color:#981a1a">-</span> <span style="color:#000000">startTime</span>;
<span style="color:#008855">double</span> <span style="color:#000000">kernelDuration</span> <span style="color:#981a1a">=</span> <span style="color:#000000">endTime_kernel</span> <span style="color:#981a1a">-</span> <span style="color:#000000">startTime_kernel</span>;
<span style="color:#000000">printf</span>(<span style="color:#aa1111">"Effective BW by CUDA saxpy: %.3f ms\t\t[%.3f GB/s]\n"</span>, <span style="color:#116644">1000.f</span> <span style="color:#981a1a">*</span> <span style="color:#000000">overallDuration</span>, <span style="color:#000000">GBPerSec</span>(<span style="color:#000000">totalBytes</span>, <span style="color:#000000">overallDuration</span>));
<span style="color:#000000">printf</span>(<span style="color:#aa1111">"\t Kernel time: %.3f ms\n"</span>, <span style="color:#116644">1000.f</span> <span style="color:#981a1a">*</span> <span style="color:#000000">kernelDuration</span>);
<span style="color:#aa5500">//</span>
<span style="color:#aa5500">// CS149 TODO: free memory buffers on the GPU using cudaFree</span>
<span style="color:#aa5500">//</span>
<span style="color:#000000">cudaFree</span>(<span style="color:#000000">device_x</span>);
<span style="color:#000000">cudaFree</span>(<span style="color:#000000">device_y</span>);
<span style="color:#000000">cudaFree</span>(<span style="color:#000000">device_result</span>);
}
<span style="color:#008855">void</span> <span style="color:#0000ff">printCudaInfo</span>() {
<span style="color:#aa5500">// print out stats about the GPU in the machine. Useful if</span>
<span style="color:#aa5500">// students want to know what GPU they are running on.</span>
<span style="color:#008855">int</span> <span style="color:#000000">deviceCount</span> <span style="color:#981a1a">=</span> <span style="color:#116644">0</span>;
<span style="color:#000000">cudaError_t</span> <span style="color:#000000">err</span> <span style="color:#981a1a">=</span> <span style="color:#000000">cudaGetDeviceCount</span>(<span style="color:#981a1a">&</span><span style="color:#000000">deviceCount</span>);
<span style="color:#000000">printf</span>(<span style="color:#aa1111">"---------------------------------------------------------\n"</span>);
<span style="color:#000000">printf</span>(<span style="color:#aa1111">"Found %d CUDA devices\n"</span>, <span style="color:#000000">deviceCount</span>);
<span style="color:#770088">for</span> (<span style="color:#008855">int</span> <span style="color:#000000">i</span><span style="color:#981a1a">=</span><span style="color:#116644">0</span>; <span style="color:#000000">i</span><span style="color:#981a1a"><</span><span style="color:#000000">deviceCount</span>; <span style="color:#000000">i</span><span style="color:#981a1a">++</span>) {
<span style="color:#000000">cudaDeviceProp</span> <span style="color:#000000">deviceProps</span>;
<span style="color:#000000">cudaGetDeviceProperties</span>(<span style="color:#981a1a">&</span><span style="color:#000000">deviceProps</span>, <span style="color:#000000">i</span>);
<span style="color:#000000">printf</span>(<span style="color:#aa1111">"Device %d: %s\n"</span>, <span style="color:#000000">i</span>, <span style="color:#000000">deviceProps</span>.<span style="color:#000000">name</span>);
<span style="color:#000000">printf</span>(<span style="color:#aa1111">" SMs: %d\n"</span>, <span style="color:#000000">deviceProps</span>.<span style="color:#000000">multiProcessorCount</span>);
<span style="color:#000000">printf</span>(<span style="color:#aa1111">" Global mem: %.0f MB\n"</span>,
<span style="color:#000000">static_cast</span><span style="color:#981a1a"><</span><span style="color:#008855">float</span><span style="color:#981a1a">></span>(<span style="color:#000000">deviceProps</span>.<span style="color:#000000">totalGlobalMem</span>) <span style="color:#981a1a">/</span> (<span style="color:#116644">1024</span> <span style="color:#981a1a">*</span> <span style="color:#116644">1024</span>));
<span style="color:#000000">printf</span>(<span style="color:#aa1111">" CUDA Cap: %d.%d\n"</span>, <span style="color:#000000">deviceProps</span>.<span style="color:#000000">major</span>, <span style="color:#000000">deviceProps</span>.<span style="color:#000000">minor</span>);
}
<span style="color:#000000">printf</span>(<span style="color:#aa1111">"---------------------------------------------------------\n"</span>);
}
</span></span>然后执行makefile
运行时如果报
这意味着我们的gpu那个参数的值不对
可以试试改为compute_75 compute_86 compute_60
Part2 CUDA并行前缀和
原题翻译：
现在您已经熟悉了 CUDA 程序的基本结构和布局，作为第二个练习，您需要完成一个函数的并行实现。其作用是在给定整数 A 的列表的情况下，返回所有满足A[i] == A[i+1]的索引 i 的列表，并统计个数。
例如，给定数组 {1，2，2，1，1，1，3，5，3，3}，程序应输出数组 {1，3，4，8}，并返回4。
我们希望您通过首先实现并行独占前缀和操作来实现上面的函数
从简单化的角度考虑这个问题，直接见数组作为参数传递，每个线程负责比较自己对应的与下一个之间的大小关系
问题1： 对GPU上同一内存的同时读取
问题2： 输出的顺序不一定是按从前到后的
对于问题1可以参考之前学过的奇偶排序的思路
分两轮进行
第一轮用A[2k]和A[2k+1]比较
第二轮用A[2k+1]和A[2k+2]比较
对于问题2只需将结果数组进行排序即可，下面给出了该思路的具体实现
前缀和采用数组 A 并生成一个新的数组输出，该输出在每个索引 i 处具有所有元素的总和，但不包括 A[i]。例如，给定数组 A={1，4，6，8，2}，则独占前缀的输出总和输出={0，1，5，11，19}。
但是这一计算过程存在潜在的并行循环问题。
每次计算都需要用到之前的元素的前缀和，这里题目中直接给出了一个可执行的伪代码，不是很好懂，我来补充一下推导过程
前缀和本质是一个递推公式
a[n] = a[n-1] + an
但是也可以拆成
a[n] = a[n/2] + a[n/2 + 1] + '' +an
所以一个思路（归并思想）是
用这个树从底向上搜索，遇到向右，就加子元素，遇到向左，就什么都不做
但这个算法是nlogn的，问题就出在没有利用好已经计算出的前缀和
（递归加分治）
对于数组A，先计算A'，是数组A每两个相邻元素的和，如果我们能计算出A'的前缀和，那么A奇数位的前缀和即为A'的前缀和，偶数维可由奇数位求出，计算A'和已知A'的前缀和求A的前缀和都是可以并行的，问题就递归为求A'的前缀和
此时问题规模递减为n/2，以此类推即可
具体的实现十分优雅，没有开辟额外的数组空间（太牛了
最终实现为O(logn)的时间开销和O(n)的空间开销
这个图就是工作流程，感觉这个图比较好懂
<span style="background-color:#f8f8f8"> <span style="color:#aa5500">// 每次计算strat到end的值</span> <span style="color:#000000">oid</span> <span style="color:#0000ff">exclusive_scan_iterative</span>(<span style="color:#008855">int*</span> <span style="color:#000000">start</span>, <span style="color:#008855">int*</span> <span style="color:#000000">end</span>, <span style="color:#008855">int*</span> <span style="color:#000000">output</span>) { <span style="color:#008855">int</span> <span style="color:#000000">N</span> <span style="color:#981a1a">=</span> <span style="color:#000000">end</span> <span style="color:#981a1a">-</span> <span style="color:#000000">start</span>; <span style="color:#aa5500">// 初始化output = strat</span> <span style="color:#000000">memmove</span>(<span style="color:#000000">output</span>, <span style="color:#000000">start</span>, <span style="color:#000000">N</span><span style="color:#981a1a">*</span><span style="color:#770088">sizeof</span>(<span style="color:#008855">int</span>)); <span style="color:#aa5500">// upsweep phase</span> <span style="color:#aa5500">// 前向求和，先求每两个相邻元素的和，保存到后一个上</span> <span style="color:#aa5500">// output[i+2-1] += output[i+1-1]</span> <span style="color:#aa5500">// 1 2 3 4 5 6 7 8</span> <span style="color:#aa5500">// 1 3 3 7 5 11 7 15</span> <span style="color:#aa5500">// 然后实际上是对2的倍数位组成的新数组重复上次操作</span> <span style="color:#aa5500">// 3 7 11 15</span> <span style="color:#aa5500">// 10 26</span> <span style="color:#aa5500">// 1 3 3 10 5 11 7 26</span> <span style="color:#770088">for</span> (<span style="color:#008855">int</span> <span style="color:#000000">two_d</span> <span style="color:#981a1a">=</span> <span style="color:#116644">1</span>; <span style="color:#000000">two_d</span> <span style="color:#981a1a"><=</span> <span style="color:#000000">N</span><span style="color:#981a1a">/</span><span style="color:#116644">2</span>; <span style="color:#000000">two_d</span><span style="color:#981a1a">*=</span><span style="color:#116644">2</span>) { <span style="color:#008855">int</span> <span style="color:#000000">two_dplus1</span> <span style="color:#981a1a">=</span> <span style="color:#116644">2</span><span style="color:#981a1a">*</span><span style="color:#000000">two_d</span>; <span style="color:#000000">parallel_for</span> (<span style="color:#008855">int</span> <span style="color:#000000">i</span> <span style="color:#981a1a">=</span> <span style="color:#116644">0</span>; <span style="color:#000000">i</span> <span style="color:#981a1a"><</span> <span style="color:#000000">N</span>; <span style="color:#000000">i</span> <span style="color:#981a1a">+=</span> <span style="color:#000000">two_dplus1</span>) { <span style="color:#000000">output</span>[<span style="color:#000000">i</span><span style="color:#981a1a">+</span><span style="color:#000000">two_dplus1</span><span style="color:#981a1a">-</span><span style="color:#116644">1</span>] <span style="color:#981a1a">+=</span> <span style="color:#000000">output</span>[<span style="color:#000000">i</span><span style="color:#981a1a">+</span><span style="color:#000000">two_d</span><span style="color:#981a1a">-</span><span style="color:#116644">1</span>]; } } <span style="color:#000000">output</span>[<span style="color:#000000">N</span><span style="color:#981a1a">-</span><span style="color:#116644">1</span>] <span style="color:#981a1a">=</span> <span style="color:#116644">0</span>; <span style="color:#aa5500">// 接下来就该往回加了</span> <span style="color:#770088">for</span> (<span style="color:#008855">int</span> <span style="color:#000000">two_d</span> <span style="color:#981a1a">=</span> <span style="color:#000000">N</span><span style="color:#981a1a">/</span><span style="color:#116644">2</span>; <span style="color:#000000">two_d</span> <span style="color:#981a1a">>=</span> <span style="color:#116644">1</span>; <span style="color:#000000">two_d</span> <span style="color:#981a1a">/=</span> <span style="color:#116644">2</span>) { <span style="color:#008855">int</span> <span style="color:#000000">two_dplus1</span> <span style="color:#981a1a">=</span> <span style="color:#116644">2</span><span style="color:#981a1a">*</span><span style="color:#000000">two_d</span>; <span style="color:#000000">parallel_for</span> (<span style="color:#008855">int</span> <span style="color:#000000">i</span> <span style="color:#981a1a">=</span> <span style="color:#116644">0</span>; <span style="color:#000000">i</span> <span style="color:#981a1a"><</span> <span style="color:#000000">N</span>; <span style="color:#000000">i</span> <span style="color:#981a1a">+=</span> <span style="color:#000000">two_dplus1</span>) { <span style="color:#008855">int</span> <span style="color:#000000">t</span> <span style="color:#981a1a">=</span> <span style="color:#000000">output</span>[<span style="color:#000000">i</span><span style="color:#981a1a">+</span><span style="color:#000000">two_d</span><span style="color:#981a1a">-</span><span style="color:#116644">1</span>]; <span style="color:#000000">output</span>[<span style="color:#000000">i</span><span style="color:#981a1a">+</span><span style="color:#000000">two_d</span><span style="color:#981a1a">-</span><span style="color:#116644">1</span>] <span style="color:#981a1a">=</span> <span style="color:#000000">output</span>[<span style="color:#000000">i</span><span style="color:#981a1a">+</span><span style="color:#000000">two_dplus1</span><span style="color:#981a1a">-</span><span style="color:#116644">1</span>]; <span style="color:#000000">output</span>[<span style="color:#000000">i</span><span style="color:#981a1a">+</span><span style="color:#000000">two_dplus1</span><span style="color:#981a1a">-</span><span style="color:#116644">1</span>] <span style="color:#981a1a">+=</span> <span style="color:#000000">t</span>; } } }</span>下面就回到正规，继续讲题目要求
2.1 请在scan/scan.cu中实现exclusive_scan功能
我们希望您使用此算法在 CUDA 中实现并行前缀总和的版本。您必须在scan/scan.cu中实现exclusive_scan功能。您的实现将包含主机和设备代码。该实现将需要多次 CUDA 内核启动（上面伪代码中的每个 parallel_for 循环一个）。
实际上是要求我们完成循环的并行实现0.0
注意：在刚才给出的参考代码中，我们假定输入数组的长度 （N） 是 2 的幂。因此在 cudaScan 函数中，我们在 GPU 上分配相应的缓冲区时，通过将输入数组长度舍入到下一个 2 的幂。同时，在最后，代码应当仅将 N 个元素从 GPU 缓冲区复制回 CPU 缓冲区。
2.1 在scan/scan.cu中实现函数find_repeats
使用前缀总和实现“查找重复项” 编写exclusive_scan后，在scan/scan.cu中实现函数find_repeats。在实现的过程中，除了对 exclusive_scan（） 的一个或多个调用之外，还需要编写额外的设备代码。最终的实现效果应将重复元素的列表写入提供的输出指针（在设备内存中），然后返回输出列表的大小。
同时我们假定传递给exclusive_scan的数组位于设备内存中。
（此处还有评分标准、参考指标什么的，跳过
Hint 2.1 实现的过程应当不需要任何实现细节上的调优，仅需要按照题目实现即可。
<span style="background-color:#f8f8f8"><span style="color:#333333"> <span style="color:#555555">#include <stdio.h></span>
<span style="color:#555555">#include <cuda.h></span>
<span style="color:#555555">#include <cuda_runtime.h></span>
<span style="color:#555555">#include <driver_functions.h></span>
<span style="color:#555555">#include <thrust/scan.h></span>
<span style="color:#555555">#include <thrust/device_ptr.h></span>
<span style="color:#555555">#include <thrust/device_malloc.h></span>
<span style="color:#555555">#include <thrust/device_free.h></span>
<span style="color:#555555">#include "CycleTimer.h"</span>
<span style="color:#555555">#define THREADS_PER_BLOCK 256</span>
<span style="color:#aa5500">// helper function to round an integer up to the next power of 2</span>
<span style="color:#aa5500">// 快速翻倍</span>
<span style="color:#770088">static</span> <span style="color:#000000">inline</span> <span style="color:#008855">int</span> <span style="color:#0000ff">nextPow2</span>(<span style="color:#008855">int</span> <span style="color:#000000">n</span>) {
<span style="color:#000000">n</span><span style="color:#981a1a">--</span>;
<span style="color:#000000">n</span> <span style="color:#981a1a">|=</span> <span style="color:#000000">n</span> <span style="color:#981a1a">>></span> <span style="color:#116644">1</span>;
<span style="color:#000000">n</span> <span style="color:#981a1a">|=</span> <span style="color:#000000">n</span> <span style="color:#981a1a">>></span> <span style="color:#116644">2</span>;
<span style="color:#000000">n</span> <span style="color:#981a1a">|=</span> <span style="color:#000000">n</span> <span style="color:#981a1a">>></span> <span style="color:#116644">4</span>;
<span style="color:#000000">n</span> <span style="color:#981a1a">|=</span> <span style="color:#000000">n</span> <span style="color:#981a1a">>></span> <span style="color:#116644">8</span>;
<span style="color:#000000">n</span> <span style="color:#981a1a">|=</span> <span style="color:#000000">n</span> <span style="color:#981a1a">>></span> <span style="color:#116644">16</span>;
<span style="color:#000000">n</span><span style="color:#981a1a">++</span>;
<span style="color:#770088">return</span> <span style="color:#000000">n</span>;
}
<span style="color:#aa5500">// exclusive_scan --</span>
<span style="color:#aa5500">//</span>
<span style="color:#aa5500">// Implementation of an exclusive scan on global memory array `input`,</span>
<span style="color:#aa5500">// with results placed in global memory `result`.</span>
<span style="color:#aa5500">//</span>
<span style="color:#aa5500">// N is the logical size of the input and output arrays, however</span>
<span style="color:#aa5500">// students can assume that both the start and result arrays we</span>
<span style="color:#aa5500">// allocated with next power-of-two sizes as described by the comments</span>
<span style="color:#aa5500">// in cudaScan(). This is helpful, since your parallel scan</span>
<span style="color:#aa5500">// will likely write to memory locations beyond N, but of course not</span>
<span style="color:#aa5500">// greater than N rounded up to the next power of 2.</span>
<span style="color:#aa5500">//</span>
<span style="color:#aa5500">// Also, as per the comments in cudaScan(), you can implement an</span>
<span style="color:#aa5500">// "in-place" scan, since the timing harness makes a copy of input and</span>
<span style="color:#aa5500">// places it in result</span>
<span style="color:#aa5500">// 默认数据已经存放在GPU上</span>
<span style="color:#aa5500">// 结果应当放置在GPU上的output中</span>
<span style="color:#aa5500">// N 是输入和输出数组的逻辑大小，但是</span>
<span style="color:#aa5500">// 我们应当将N扩充为2的幂次</span>
<span style="color:#aa5500">// 另外，根据 cudaScan（） 中的注释，您可以实现一个</span>
<span style="color:#aa5500">// “就地”扫描，利用循环输出将被放置于output中的特性，减少内存复制的次数</span>
<span style="color:#aa5500">// 前向循环，利用了output内放置了输入的特性</span>
<span style="color:#aa5500">// 实现较为简单，只要理解了前缀和的计算算法即可</span>
<span style="color:#000000">__global__</span> <span style="color:#008855">void</span> <span style="color:#0000ff">upsweep</span>(<span style="color:#008855">int</span> <span style="color:#000000">N</span>, <span style="color:#008855">int</span> <span style="color:#000000">two_d</span>, <span style="color:#008855">int*</span> <span style="color:#000000">output</span>) {
<span style="color:#008855">int</span> <span style="color:#000000">two_dplus1</span> <span style="color:#981a1a">=</span> <span style="color:#116644">2</span><span style="color:#981a1a">*</span><span style="color:#000000">two_d</span>;
<span style="color:#008855">int</span> <span style="color:#000000">j</span> <span style="color:#981a1a">=</span> <span style="color:#000000">blockIdx</span>.<span style="color:#000000">x</span> <span style="color:#981a1a">*</span> <span style="color:#000000">blockDim</span>.<span style="color:#000000">x</span> <span style="color:#981a1a">+</span> <span style="color:#000000">threadIdx</span>.<span style="color:#000000">x</span>;
<span style="color:#008855">int</span> <span style="color:#000000">i</span> <span style="color:#981a1a">=</span> <span style="color:#000000">j</span> <span style="color:#981a1a">*</span> <span style="color:#000000">two_dplus1</span>;
<span style="color:#000000">output</span>[<span style="color:#000000">i</span><span style="color:#981a1a">+</span><span style="color:#000000">two_dplus1</span><span style="color:#981a1a">-</span><span style="color:#116644">1</span>] <span style="color:#981a1a">+=</span> <span style="color:#000000">output</span>[<span style="color:#000000">i</span><span style="color:#981a1a">+</span><span style="color:#000000">two_d</span><span style="color:#981a1a">-</span><span style="color:#116644">1</span>];
}
<span style="color:#aa5500">// 反向循环，利用了output内放置了输入的特性</span>
<span style="color:#000000">__global__</span> <span style="color:#008855">void</span> <span style="color:#0000ff">downsweep</span>(<span style="color:#008855">int</span> <span style="color:#000000">N</span>, <span style="color:#008855">int</span> <span style="color:#000000">two_d</span>, <span style="color:#008855">int*</span> <span style="color:#000000">output</span>) {
<span style="color:#aa5500">// Handle first iteration</span>
<span style="color:#770088">if</span> (<span style="color:#000000">two_d</span> <span style="color:#981a1a">==</span> <span style="color:#000000">N</span><span style="color:#981a1a">/</span><span style="color:#116644">2</span>)
<span style="color:#000000">output</span>[<span style="color:#000000">N</span><span style="color:#981a1a">-</span><span style="color:#116644">1</span>] <span style="color:#981a1a">=</span> <span style="color:#116644">0</span>;
<span style="color:#008855">int</span> <span style="color:#000000">two_dplus1</span> <span style="color:#981a1a">=</span> <span style="color:#116644">2</span><span style="color:#981a1a">*</span><span style="color:#000000">two_d</span>;
<span style="color:#008855">int</span> <span style="color:#000000">j</span> <span style="color:#981a1a">=</span> <span style="color:#000000">blockIdx</span>.<span style="color:#000000">x</span> <span style="color:#981a1a">*</span> <span style="color:#000000">blockDim</span>.<span style="color:#000000">x</span> <span style="color:#981a1a">+</span> <span style="color:#000000">threadIdx</span>.<span style="color:#000000">x</span>;
<span style="color:#008855">int</span> <span style="color:#000000">i</span> <span style="color:#981a1a">=</span> <span style="color:#000000">j</span> <span style="color:#981a1a">*</span> <span style="color:#000000">two_dplus1</span>;
<span style="color:#008855">int</span> <span style="color:#000000">t</span> <span style="color:#981a1a">=</span> <span style="color:#000000">output</span>[<span style="color:#000000">i</span><span style="color:#981a1a">+</span><span style="color:#000000">two_d</span><span style="color:#981a1a">-</span><span style="color:#116644">1</span>];
<span style="color:#000000">output</span>[<span style="color:#000000">i</span><span style="color:#981a1a">+</span><span style="color:#000000">two_d</span><span style="color:#981a1a">-</span><span style="color:#116644">1</span>] <span style="color:#981a1a">=</span> <span style="color:#000000">output</span>[<span style="color:#000000">i</span><span style="color:#981a1a">+</span><span style="color:#000000">two_dplus1</span><span style="color:#981a1a">-</span><span style="color:#116644">1</span>];
<span style="color:#000000">output</span>[<span style="color:#000000">i</span><span style="color:#981a1a">+</span><span style="color:#000000">two_dplus1</span><span style="color:#981a1a">-</span><span style="color:#116644">1</span>] <span style="color:#981a1a">+=</span> <span style="color:#000000">t</span>;
}
<span style="color:#008855">void</span> <span style="color:#0000ff">exclusive_scan</span>(<span style="color:#008855">int*</span> <span style="color:#000000">input</span>, <span style="color:#008855">int</span> <span style="color:#000000">N</span>, <span style="color:#008855">int*</span> <span style="color:#000000">result</span>)
{
<span style="color:#aa5500">// CS149 TODO:</span>
<span style="color:#aa5500">//</span>
<span style="color:#aa5500">// Implement your exclusive scan implementation here. Keep input</span>
<span style="color:#aa5500">// mind that although the arguments to this function are device</span>
<span style="color:#aa5500">// allocated arrays, this is a function that is running in a thread</span>
<span style="color:#aa5500">// on the CPU. Your implementation will need to make multiple calls</span>
<span style="color:#aa5500">// to CUDA kernel functions (that you must write) to implement the</span>
<span style="color:#aa5500">// scan.</span>
<span style="color:#aa5500">//在此处实现您的前缀和算法，并保留输入</span>
<span style="color:#aa5500">// 请注意，尽管此函数的参数是设备分配的数组</span>
<span style="color:#aa5500">// 但这是一个运行在CPU上的函数，我们需要多次调用CUDA内核来完成计算任务</span>
<span style="color:#008855">int</span> <span style="color:#000000">N2</span> <span style="color:#981a1a">=</span> <span style="color:#000000">nextPow2</span>(<span style="color:#000000">N</span>);
<span style="color:#770088">const</span> <span style="color:#008855">int</span> <span style="color:#000000">threadsPerBlock</span> <span style="color:#981a1a">=</span> <span style="color:#116644">512</span>;
<span style="color:#aa5500">// upsweep phase</span>
<span style="color:#770088">for</span> (<span style="color:#008855">int</span> <span style="color:#000000">two_d</span> <span style="color:#981a1a">=</span> <span style="color:#116644">1</span>; <span style="color:#000000">two_d</span> <span style="color:#981a1a"><</span> <span style="color:#000000">N2</span><span style="color:#981a1a">/</span><span style="color:#116644">2</span>; <span style="color:#000000">two_d</span><span style="color:#981a1a">*=</span><span style="color:#116644">2</span>) {
<span style="color:#008855">int</span> <span style="color:#000000">two_dplus1</span> <span style="color:#981a1a">=</span> <span style="color:#116644">2</span><span style="color:#981a1a">*</span><span style="color:#000000">two_d</span>;
<span style="color:#008855">int</span> <span style="color:#000000">numThreads</span> <span style="color:#981a1a">=</span> <span style="color:#000000">N2</span> <span style="color:#981a1a">/</span> <span style="color:#000000">two_dplus1</span>;
<span style="color:#008855">int</span> <span style="color:#000000">blocks</span> <span style="color:#981a1a">=</span> (<span style="color:#000000">numThreads</span> <span style="color:#981a1a">+</span> <span style="color:#000000">threadsPerBlock</span> <span style="color:#981a1a">-</span> <span style="color:#116644">1</span>) <span style="color:#981a1a">/</span> <span style="color:#000000">threadsPerBlock</span>;
<span style="color:#aa5500">// 相当于之前伪代码中的for循环</span>
<span style="color:#000000">upsweep</span><span style="color:#981a1a"><<<</span><span style="color:#000000">blocks</span>,<span style="color:#000000">threadsPerBlock</span><span style="color:#981a1a">>>></span>(<span style="color:#000000">N2</span>, <span style="color:#000000">two_d</span>, <span style="color:#000000">result</span>);
<span style="color:#000000">cudaDeviceSynchronize</span>();
}
<span style="color:#aa5500">// downsweep phase</span>
<span style="color:#770088">for</span> (<span style="color:#008855">int</span> <span style="color:#000000">two_d</span> <span style="color:#981a1a">=</span> <span style="color:#000000">N2</span><span style="color:#981a1a">/</span><span style="color:#116644">2</span>; <span style="color:#000000">two_d</span> <span style="color:#981a1a">>=</span> <span style="color:#116644">1</span>; <span style="color:#000000">two_d</span> <span style="color:#981a1a">/=</span> <span style="color:#116644">2</span>) {
<span style="color:#008855">int</span> <span style="color:#000000">two_dplus1</span> <span style="color:#981a1a">=</span> <span style="color:#116644">2</span><span style="color:#981a1a">*</span><span style="color:#000000">two_d</span>;
<span style="color:#008855">int</span> <span style="color:#000000">numThreads</span> <span style="color:#981a1a">=</span> <span style="color:#000000">N2</span> <span style="color:#981a1a">/</span> <span style="color:#000000">two_dplus1</span>;
<span style="color:#008855">int</span> <span style="color:#000000">blocks</span> <span style="color:#981a1a">=</span> (<span style="color:#000000">numThreads</span> <span style="color:#981a1a">+</span> <span style="color:#000000">threadsPerBlock</span> <span style="color:#981a1a">-</span> <span style="color:#116644">1</span>) <span style="color:#981a1a">/</span> <span style="color:#000000">threadsPerBlock</span>;
<span style="color:#aa5500">// 相当于之前伪代码中的for循环</span>
<span style="color:#000000">downsweep</span><span style="color:#981a1a"><<<</span><span style="color:#000000">blocks</span>,<span style="color:#000000">threadsPerBlock</span><span style="color:#981a1a">>>></span>(<span style="color:#000000">N2</span>, <span style="color:#000000">two_d</span>, <span style="color:#000000">result</span>);
<span style="color:#000000">cudaDeviceSynchronize</span>();
}
}
<span style="color:#aa5500">//</span>
<span style="color:#aa5500">// cudaScan --</span>
<span style="color:#aa5500">//</span>
<span style="color:#aa5500">// This function is a timing wrapper around the student's</span>
<span style="color:#aa5500">// implementation of scan - it copies the input to the GPU</span>
<span style="color:#aa5500">// and times the invocation of the exclusive_scan() function</span>
<span style="color:#aa5500">// above. Students should not modify it.</span>
<span style="color:#aa5500">// 你不应该修改它！</span>
<span style="color:#aa5500">// 但是可以读读，这个是需要实现的第二个函数的上层接口</span>
<span style="color:#aa5500">// 在其中完成了GPU端的内存申请、排序数组初始化</span>
<span style="color:#aa5500">// 这和我们实现的第一个函数的前提条件也能对得上</span>
<span style="color:#008855">double</span> <span style="color:#0000ff">cudaScan</span>(<span style="color:#008855">int*</span> <span style="color:#000000">inarray</span>, <span style="color:#008855">int*</span> <span style="color:#000000">end</span>, <span style="color:#008855">int*</span> <span style="color:#000000">resultarray</span>)
{
<span style="color:#008855">int*</span> <span style="color:#000000">device_result</span>;
<span style="color:#008855">int*</span> <span style="color:#000000">device_input</span>;
<span style="color:#aa5500">// 指针相减得长度</span>
<span style="color:#008855">int</span> <span style="color:#000000">N</span> <span style="color:#981a1a">=</span> <span style="color:#000000">end</span> <span style="color:#981a1a">-</span> <span style="color:#000000">inarray</span>;
<span style="color:#aa5500">// This code rounds the arrays provided to exclusive_scan up</span>
<span style="color:#aa5500">// to a power of 2, but elements after the end of the original</span>
<span style="color:#aa5500">// input are left uninitialized and not checked for correctness.</span>
<span style="color:#aa5500">//</span>
<span style="color:#aa5500">// Student implementations of exclusive_scan may assume an array's</span>
<span style="color:#aa5500">// allocated length is a power of 2 for simplicity. This will</span>
<span style="color:#aa5500">// result in extra work on non-power-of-2 inputs, but it's worth</span>
<span style="color:#aa5500">// the simplicity of a power of two only solution.</span>
<span style="color:#aa5500">// 将N扩充到2的幂次</span>
<span style="color:#aa5500">// 虽然我们需要做额外的工作，但这是解决2的幂次问题最简单的方法</span>
<span style="color:#008855">int</span> <span style="color:#000000">rounded_length</span> <span style="color:#981a1a">=</span> <span style="color:#000000">nextPow2</span>(<span style="color:#000000">end</span> <span style="color:#981a1a">-</span> <span style="color:#000000">inarray</span>);
<span style="color:#000000">cudaMalloc</span>((<span style="color:#008855">void</span> <span style="color:#008855">**</span>)<span style="color:#981a1a">&</span><span style="color:#000000">device_result</span>, <span style="color:#770088">sizeof</span>(<span style="color:#008855">int</span>) <span style="color:#981a1a">*</span> <span style="color:#000000">rounded_length</span>);
<span style="color:#000000">cudaMalloc</span>((<span style="color:#008855">void</span> <span style="color:#008855">**</span>)<span style="color:#981a1a">&</span><span style="color:#000000">device_input</span>, <span style="color:#770088">sizeof</span>(<span style="color:#008855">int</span>) <span style="color:#981a1a">*</span> <span style="color:#000000">rounded_length</span>);
<span style="color:#aa5500">// For convenience, both the input and output vectors on the</span>
<span style="color:#aa5500">// device are initialized to the input values. This means that</span>
<span style="color:#aa5500">// students are free to implement an in-place scan on the result</span>
<span style="color:#aa5500">// vector if desired. If you do this, you will need to keep this</span>
<span style="color:#aa5500">// in mind when calling exclusive_scan from find_repeats.</span>
<span style="color:#000000">cudaMemcpy</span>(<span style="color:#000000">device_input</span>, <span style="color:#000000">inarray</span>, (<span style="color:#000000">end</span> <span style="color:#981a1a">-</span> <span style="color:#000000">inarray</span>) <span style="color:#981a1a">*</span> <span style="color:#770088">sizeof</span>(<span style="color:#008855">int</span>), <span style="color:#000000">cudaMemcpyHostToDevice</span>);
<span style="color:#000000">cudaMemcpy</span>(<span style="color:#000000">device_result</span>, <span style="color:#000000">inarray</span>, (<span style="color:#000000">end</span> <span style="color:#981a1a">-</span> <span style="color:#000000">inarray</span>) <span style="color:#981a1a">*</span> <span style="color:#770088">sizeof</span>(<span style="color:#008855">int</span>), <span style="color:#000000">cudaMemcpyHostToDevice</span>);
<span style="color:#008855">double</span> <span style="color:#000000">startTime</span> <span style="color:#981a1a">=</span> <span style="color:#000000">CycleTimer</span>::<span style="color:#000000">currentSeconds</span>();
<span style="color:#000000">exclusive_scan</span>(<span style="color:#000000">device_input</span>, <span style="color:#000000">N</span>, <span style="color:#000000">device_result</span>);
<span style="color:#aa5500">// Wait for completion</span>
<span style="color:#000000">cudaDeviceSynchronize</span>();
<span style="color:#008855">double</span> <span style="color:#000000">endTime</span> <span style="color:#981a1a">=</span> <span style="color:#000000">CycleTimer</span>::<span style="color:#000000">currentSeconds</span>();
<span style="color:#000000">cudaMemcpy</span>(<span style="color:#000000">resultarray</span>, <span style="color:#000000">device_result</span>, (<span style="color:#000000">end</span> <span style="color:#981a1a">-</span> <span style="color:#000000">inarray</span>) <span style="color:#981a1a">*</span> <span style="color:#770088">sizeof</span>(<span style="color:#008855">int</span>), <span style="color:#000000">cudaMemcpyDeviceToHost</span>);
<span style="color:#008855">double</span> <span style="color:#000000">overallDuration</span> <span style="color:#981a1a">=</span> <span style="color:#000000">endTime</span> <span style="color:#981a1a">-</span> <span style="color:#000000">startTime</span>;
<span style="color:#770088">return</span> <span style="color:#000000">overallDuration</span>;
}
<span style="color:#aa5500">// cudaScanThrust --</span>
<span style="color:#aa5500">//</span>
<span style="color:#aa5500">// Wrapper around the Thrust library's exclusive scan function</span>
<span style="color:#aa5500">// As above in cudaScan(), this function copies the input to the GPU</span>
<span style="color:#aa5500">// and times only the execution of the scan itself.</span>
<span style="color:#aa5500">//</span>
<span style="color:#aa5500">// Students are not expected to produce implementations that achieve</span>
<span style="color:#aa5500">// performance that is competition to the Thrust version, but it is fun to try.</span>
<span style="color:#aa5500">// 这个函数实现输入的处理</span>
<span style="color:#aa5500">// 并将输入复制到GPU端</span>
<span style="color:#aa5500">// 不要修改！</span>
<span style="color:#008855">double</span> <span style="color:#0000ff">cudaScanThrust</span>(<span style="color:#008855">int*</span> <span style="color:#000000">inarray</span>, <span style="color:#008855">int*</span> <span style="color:#000000">end</span>, <span style="color:#008855">int*</span> <span style="color:#000000">resultarray</span>) {
<span style="color:#008855">int</span> <span style="color:#000000">length</span> <span style="color:#981a1a">=</span> <span style="color:#000000">end</span> <span style="color:#981a1a">-</span> <span style="color:#000000">inarray</span>;
<span style="color:#000000">thrust</span>::<span style="color:#000000">device_ptr</span><span style="color:#981a1a"><</span><span style="color:#008855">int</span><span style="color:#981a1a">></span> <span style="color:#000000">d_input</span> <span style="color:#981a1a">=</span> <span style="color:#000000">thrust</span>::<span style="color:#000000">device_malloc</span><span style="color:#981a1a"><</span><span style="color:#008855">int</span><span style="color:#981a1a">></span>(<span style="color:#000000">length</span>);
<span style="color:#000000">thrust</span>::<span style="color:#000000">device_ptr</span><span style="color:#981a1a"><</span><span style="color:#008855">int</span><span style="color:#981a1a">></span> <span style="color:#000000">d_output</span> <span style="color:#981a1a">=</span> <span style="color:#000000">thrust</span>::<span style="color:#000000">device_malloc</span><span style="color:#981a1a"><</span><span style="color:#008855">int</span><span style="color:#981a1a">></span>(<span style="color:#000000">length</span>);
<span style="color:#000000">cudaMemcpy</span>(<span style="color:#000000">d_input</span>.<span style="color:#000000">get</span>(), <span style="color:#000000">inarray</span>, <span style="color:#000000">length</span> <span style="color:#981a1a">*</span> <span style="color:#770088">sizeof</span>(<span style="color:#008855">int</span>), <span style="color:#000000">cudaMemcpyHostToDevice</span>);
<span style="color:#008855">double</span> <span style="color:#000000">startTime</span> <span style="color:#981a1a">=</span> <span style="color:#000000">CycleTimer</span>::<span style="color:#000000">currentSeconds</span>();
<span style="color:#000000">thrust</span>::<span style="color:#000000">exclusive_scan</span>(<span style="color:#000000">d_input</span>, <span style="color:#000000">d_input</span> <span style="color:#981a1a">+</span> <span style="color:#000000">length</span>, <span style="color:#000000">d_output</span>);
<span style="color:#000000">cudaDeviceSynchronize</span>();
<span style="color:#008855">double</span> <span style="color:#000000">endTime</span> <span style="color:#981a1a">=</span> <span style="color:#000000">CycleTimer</span>::<span style="color:#000000">currentSeconds</span>();
<span style="color:#000000">cudaMemcpy</span>(<span style="color:#000000">resultarray</span>, <span style="color:#000000">d_output</span>.<span style="color:#000000">get</span>(), <span style="color:#000000">length</span> <span style="color:#981a1a">*</span> <span style="color:#770088">sizeof</span>(<span style="color:#008855">int</span>), <span style="color:#000000">cudaMemcpyDeviceToHost</span>);
<span style="color:#000000">thrust</span>::<span style="color:#000000">device_free</span>(<span style="color:#000000">d_input</span>);
<span style="color:#000000">thrust</span>::<span style="color:#000000">device_free</span>(<span style="color:#000000">d_output</span>);
<span style="color:#008855">double</span> <span style="color:#000000">overallDuration</span> <span style="color:#981a1a">=</span> <span style="color:#000000">endTime</span> <span style="color:#981a1a">-</span> <span style="color:#000000">startTime</span>;
<span style="color:#770088">return</span> <span style="color:#000000">overallDuration</span>;
}
<span style="color:#aa5500">// find_repeats --</span>
<span style="color:#aa5500">//</span>
<span style="color:#aa5500">// Given an array of integers `device_input`, returns an array of all</span>
<span style="color:#aa5500">// indices `i` for which `device_input[i] == device_input[i+1]`.</span>
<span style="color:#aa5500">//</span>
<span style="color:#aa5500">// Returns the total number of pairs found</span>
<span style="color:#aa5500">// 这个就是我们要实现的第二个函数了</span>
<span style="color:#aa5500">// 可以使用前面实现的前缀和算法</span>
<span style="color:#aa5500">// 默认输入位于GPU</span>
<span style="color:#aa5500">// 输出也应当放置于GPU</span>
<span style="color:#aa5500">// 这里先用了不分奇偶的算法</span>
<span style="color:#aa5500">// 这里去了解了一下cuda内存访问机制</span>
<span style="color:#aa5500">// cuda中内存访问一定会经过L2 cache</span>
<span style="color:#aa5500">// 一个warp中的线程进行访问时，我们用的全局内存</span>
<span style="color:#aa5500">// 第一次内存访问时将数据加载至cache</span>
<span style="color:#aa5500">// 后续访问即使有冲突也能在较短的时间内恢复</span>
<span style="color:#000000">__global__</span> <span style="color:#008855">void</span> <span style="color:#0000ff">isrepeat</span>(<span style="color:#008855">int*</span> <span style="color:#000000">input</span>, <span style="color:#008855">int</span> <span style="color:#000000">length</span>, <span style="color:#008855">int*</span> <span style="color:#000000">output</span>) {
<span style="color:#008855">int</span> <span style="color:#000000">index</span> <span style="color:#981a1a">=</span> <span style="color:#000000">blockIdx</span>.<span style="color:#000000">x</span> <span style="color:#981a1a">*</span> <span style="color:#000000">blockDim</span>.<span style="color:#000000">x</span> <span style="color:#981a1a">+</span> <span style="color:#000000">threadIdx</span>.<span style="color:#000000">x</span>;
<span style="color:#770088">if</span> ( (<span style="color:#000000">index</span> <span style="color:#981a1a"><</span> <span style="color:#000000">length</span> <span style="color:#981a1a">-</span> <span style="color:#116644">1</span>) <span style="color:#981a1a">&&</span> (<span style="color:#000000">input</span>[<span style="color:#000000">index</span>] <span style="color:#981a1a">==</span> <span style="color:#000000">input</span>[<span style="color:#000000">index</span><span style="color:#981a1a">+</span><span style="color:#116644">1</span>]) )
<span style="color:#000000">output</span>[<span style="color:#000000">index</span>] <span style="color:#981a1a">=</span> <span style="color:#116644">1</span>;
<span style="color:#770088">else</span>
<span style="color:#000000">output</span>[<span style="color:#000000">index</span>] <span style="color:#981a1a">=</span> <span style="color:#116644">0</span>;
}
<span style="color:#008855">int</span> <span style="color:#0000ff">find_repeats</span>(<span style="color:#008855">int*</span> <span style="color:#000000">device_input</span>, <span style="color:#008855">int</span> <span style="color:#000000">length</span>, <span style="color:#008855">int*</span> <span style="color:#000000">device_output</span>) {
<span style="color:#aa5500">// CS149 TODO:</span>
<span style="color:#aa5500">//</span>
<span style="color:#aa5500">// Implement this function. You will probably want to</span>
<span style="color:#aa5500">// make use of one or more calls to exclusive_scan(), as well as</span>
<span style="color:#aa5500">// additional CUDA kernel launches.</span>
<span style="color:#aa5500">// </span>
<span style="color:#aa5500">// Note: As in the scan code, the calling code ensures that</span>
<span style="color:#aa5500">// allocated arrays are a power of 2 in size, so you can use your</span>
<span style="color:#aa5500">// exclusive_scan function with them. However, your implementation</span>
<span style="color:#aa5500">// must ensure that the results of find_repeats are correct given</span>
<span style="color:#aa5500">// the actual array length.</span>
<span style="color:#aa5500">// 这段和前面的意思一样：强调2的幂次处理、结果的正确性</span>
<span style="color:#770088">const</span> <span style="color:#008855">int</span> <span style="color:#000000">threadsPerBlock</span> <span style="color:#981a1a">=</span> <span style="color:#116644">512</span>;
<span style="color:#770088">const</span> <span style="color:#008855">int</span> <span style="color:#000000">blocks</span> <span style="color:#981a1a">=</span> (<span style="color:#000000">length</span> <span style="color:#981a1a">+</span> <span style="color:#000000">threadsPerBlock</span> <span style="color:#981a1a">-</span> <span style="color:#116644">1</span>) <span style="color:#981a1a">/</span> <span style="color:#000000">threadsPerBlock</span>;
<span style="color:#aa5500">// Place 1s wherever there is a repeat</span>
<span style="color:#000000">isrepeat</span><span style="color:#981a1a"><<<</span><span style="color:#000000">blocks</span>,<span style="color:#000000">threadsPerBlock</span><span style="color:#981a1a">>>></span>(<span style="color:#000000">device_input</span>, <span style="color:#000000">length</span>, <span style="color:#000000">device_output</span>);
<span style="color:#aa5500">// Use scan to count the total</span>
<span style="color:#000000">exclusive_scan</span>(<span style="color:#000000">device_input</span>, <span style="color:#000000">length</span>, <span style="color:#000000">device_output</span>);
<span style="color:#008855">int*</span> <span style="color:#000000">output</span>;
<span style="color:#000000">output</span> <span style="color:#981a1a">=</span> (<span style="color:#008855">int*</span>)<span style="color:#000000">malloc</span>(<span style="color:#770088">sizeof</span>(<span style="color:#008855">int</span>));
<span style="color:#000000">cudaMemcpy</span>(<span style="color:#000000">output</span>, <span style="color:#000000">device_output</span> <span style="color:#981a1a">+</span> <span style="color:#000000">length</span><span style="color:#981a1a">-</span><span style="color:#116644">1</span>, <span style="color:#770088">sizeof</span>(<span style="color:#008855">int</span>), <span style="color:#000000">cudaMemcpyDeviceToHost</span>);
<span style="color:#000000">printf</span>(<span style="color:#aa1111">"total %d\n"</span>, <span style="color:#000000">output</span>[<span style="color:#116644">0</span>]);
<span style="color:#008855">int</span> <span style="color:#000000">repeats</span> <span style="color:#981a1a">=</span> <span style="color:#000000">output</span>[<span style="color:#116644">0</span>];
<span style="color:#000000">free</span>(<span style="color:#000000">output</span>);
<span style="color:#770088">return</span> <span style="color:#000000">repeats</span>;
}
<span style="color:#aa5500">//</span>
<span style="color:#aa5500">// cudaFindRepeats --</span>
<span style="color:#aa5500">//</span>
<span style="color:#aa5500">// Timing wrapper around find_repeats. You should not modify this function.</span>
<span style="color:#008855">double</span> <span style="color:#0000ff">cudaFindRepeats</span>(<span style="color:#008855">int</span> <span style="color:#008855">*</span><span style="color:#000000">input</span>, <span style="color:#008855">int</span> <span style="color:#000000">length</span>, <span style="color:#008855">int</span> <span style="color:#008855">*</span><span style="color:#000000">output</span>, <span style="color:#008855">int</span> <span style="color:#008855">*</span><span style="color:#000000">output_length</span>) {
<span style="color:#008855">int</span> <span style="color:#008855">*</span><span style="color:#000000">device_input</span>;
<span style="color:#008855">int</span> <span style="color:#008855">*</span><span style="color:#000000">device_output</span>;
<span style="color:#008855">int</span> <span style="color:#000000">rounded_length</span> <span style="color:#981a1a">=</span> <span style="color:#000000">nextPow2</span>(<span style="color:#000000">length</span>);
<span style="color:#000000">cudaMalloc</span>((<span style="color:#008855">void</span> <span style="color:#008855">**</span>)<span style="color:#981a1a">&</span><span style="color:#000000">device_input</span>, <span style="color:#000000">rounded_length</span> <span style="color:#981a1a">*</span> <span style="color:#770088">sizeof</span>(<span style="color:#008855">int</span>));
<span style="color:#000000">cudaMalloc</span>((<span style="color:#008855">void</span> <span style="color:#008855">**</span>)<span style="color:#981a1a">&</span><span style="color:#000000">device_output</span>, <span style="color:#000000">rounded_length</span> <span style="color:#981a1a">*</span> <span style="color:#770088">sizeof</span>(<span style="color:#008855">int</span>));
<span style="color:#000000">cudaMemcpy</span>(<span style="color:#000000">device_input</span>, <span style="color:#000000">input</span>, <span style="color:#000000">length</span> <span style="color:#981a1a">*</span> <span style="color:#770088">sizeof</span>(<span style="color:#008855">int</span>), <span style="color:#000000">cudaMemcpyHostToDevice</span>);
<span style="color:#000000">cudaDeviceSynchronize</span>();
<span style="color:#008855">double</span> <span style="color:#000000">startTime</span> <span style="color:#981a1a">=</span> <span style="color:#000000">CycleTimer</span>::<span style="color:#000000">currentSeconds</span>();
<span style="color:#008855">int</span> <span style="color:#000000">result</span> <span style="color:#981a1a">=</span> <span style="color:#000000">find_repeats</span>(<span style="color:#000000">device_input</span>, <span style="color:#000000">length</span>, <span style="color:#000000">device_output</span>);
<span style="color:#000000">cudaDeviceSynchronize</span>();
<span style="color:#008855">double</span> <span style="color:#000000">endTime</span> <span style="color:#981a1a">=</span> <span style="color:#000000">CycleTimer</span>::<span style="color:#000000">currentSeconds</span>();
<span style="color:#aa5500">// set output count and results array</span>
<span style="color:#981a1a">*</span><span style="color:#000000">output_length</span> <span style="color:#981a1a">=</span> <span style="color:#000000">result</span>;
<span style="color:#000000">cudaMemcpy</span>(<span style="color:#000000">output</span>, <span style="color:#000000">device_output</span>, <span style="color:#000000">length</span> <span style="color:#981a1a">*</span> <span style="color:#770088">sizeof</span>(<span style="color:#008855">int</span>), <span style="color:#000000">cudaMemcpyDeviceToHost</span>);
<span style="color:#000000">cudaFree</span>(<span style="color:#000000">device_input</span>);
<span style="color:#000000">cudaFree</span>(<span style="color:#000000">device_output</span>);
<span style="color:#008855">float</span> <span style="color:#000000">duration</span> <span style="color:#981a1a">=</span> <span style="color:#000000">endTime</span> <span style="color:#981a1a">-</span> <span style="color:#000000">startTime</span>;
<span style="color:#770088">return</span> <span style="color:#000000">duration</span>;
}
<span style="color:#008855">void</span> <span style="color:#0000ff">printCudaInfo</span>()
{
<span style="color:#008855">int</span> <span style="color:#000000">deviceCount</span> <span style="color:#981a1a">=</span> <span style="color:#116644">0</span>;
<span style="color:#000000">cudaError_t</span> <span style="color:#000000">err</span> <span style="color:#981a1a">=</span> <span style="color:#000000">cudaGetDeviceCount</span>(<span style="color:#981a1a">&</span><span style="color:#000000">deviceCount</span>);
<span style="color:#000000">printf</span>(<span style="color:#aa1111">"---------------------------------------------------------\n"</span>);
<span style="color:#000000">printf</span>(<span style="color:#aa1111">"Found %d CUDA devices\n"</span>, <span style="color:#000000">deviceCount</span>);
<span style="color:#770088">for</span> (<span style="color:#008855">int</span> <span style="color:#000000">i</span><span style="color:#981a1a">=</span><span style="color:#116644">0</span>; <span style="color:#000000">i</span><span style="color:#981a1a"><</span><span style="color:#000000">deviceCount</span>; <span style="color:#000000">i</span><span style="color:#981a1a">++</span>)
{
<span style="color:#000000">cudaDeviceProp</span> <span style="color:#000000">deviceProps</span>;
<span style="color:#000000">cudaGetDeviceProperties</span>(<span style="color:#981a1a">&</span><span style="color:#000000">deviceProps</span>, <span style="color:#000000">i</span>);
<span style="color:#000000">printf</span>(<span style="color:#aa1111">"Device %d: %s\n"</span>, <span style="color:#000000">i</span>, <span style="color:#000000">deviceProps</span>.<span style="color:#000000">name</span>);
<span style="color:#000000">printf</span>(<span style="color:#aa1111">" SMs: %d\n"</span>, <span style="color:#000000">deviceProps</span>.<span style="color:#000000">multiProcessorCount</span>);
<span style="color:#000000">printf</span>(<span style="color:#aa1111">" Global mem: %.0f MB\n"</span>,
<span style="color:#000000">static_cast</span><span style="color:#981a1a"><</span><span style="color:#008855">float</span><span style="color:#981a1a">></span>(<span style="color:#000000">deviceProps</span>.<span style="color:#000000">totalGlobalMem</span>) <span style="color:#981a1a">/</span> (<span style="color:#116644">1024</span> <span style="color:#981a1a">*</span> <span style="color:#116644">1024</span>));
<span style="color:#000000">printf</span>(<span style="color:#aa1111">" CUDA Cap: %d.%d\n"</span>, <span style="color:#000000">deviceProps</span>.<span style="color:#000000">major</span>, <span style="color:#000000">deviceProps</span>.<span style="color:#000000">minor</span>);
}
<span style="color:#000000">printf</span>(<span style="color:#aa1111">"---------------------------------------------------------\n"</span>);
}</span></span>考虑进行奇偶优化减少访存过程中的碰撞
（主要是修改find_repeat
<span style="background-color:#f8f8f8"><span style="color:#333333"> <span style="color:#000000">__global__</span> <span style="color:#008855">void</span> <span style="color:#0000ff">isrepeat1</span>(<span style="color:#008855">int*</span> <span style="color:#000000">input</span>, <span style="color:#008855">int</span> <span style="color:#000000">length</span>, <span style="color:#008855">int*</span> <span style="color:#000000">output</span>) {
<span style="color:#008855">int</span> <span style="color:#000000">index</span> <span style="color:#981a1a">=</span> <span style="color:#000000">blockIdx</span>.<span style="color:#000000">x</span> <span style="color:#981a1a">*</span> <span style="color:#000000">blockDim</span>.<span style="color:#000000">x</span> <span style="color:#981a1a">+</span> <span style="color:#000000">threadIdx</span>.<span style="color:#000000">x</span>;
<span style="color:#000000">index</span> <span style="color:#981a1a">=</span> <span style="color:#000000">index</span> <span style="color:#981a1a">*</span> <span style="color:#116644">2</span>;
<span style="color:#770088">if</span> ( (<span style="color:#000000">index</span> <span style="color:#981a1a"><</span> <span style="color:#000000">length</span> <span style="color:#981a1a">-</span> <span style="color:#116644">1</span>) <span style="color:#981a1a">&&</span> (<span style="color:#000000">input</span>[<span style="color:#000000">index</span>] <span style="color:#981a1a">==</span> <span style="color:#000000">input</span>[<span style="color:#000000">index</span> <span style="color:#981a1a">+</span> <span style="color:#116644">1</span>]) )
<span style="color:#000000">output</span>[<span style="color:#000000">index</span>] <span style="color:#981a1a">=</span> <span style="color:#116644">1</span>;
<span style="color:#770088">else</span>
<span style="color:#000000">output</span>[<span style="color:#000000">index</span>] <span style="color:#981a1a">=</span> <span style="color:#116644">0</span>;
}
<span style="color:#000000">__global__</span> <span style="color:#008855">void</span> <span style="color:#0000ff">isrepeat2</span>(<span style="color:#008855">int*</span> <span style="color:#000000">input</span>, <span style="color:#008855">int</span> <span style="color:#000000">length</span>, <span style="color:#008855">int*</span> <span style="color:#000000">output</span>) {
<span style="color:#008855">int</span> <span style="color:#000000">index</span> <span style="color:#981a1a">=</span> <span style="color:#000000">blockIdx</span>.<span style="color:#000000">x</span> <span style="color:#981a1a">*</span> <span style="color:#000000">blockDim</span>.<span style="color:#000000">x</span> <span style="color:#981a1a">+</span> <span style="color:#000000">threadIdx</span>.<span style="color:#000000">x</span>;
<span style="color:#000000">index</span> <span style="color:#981a1a">=</span> <span style="color:#000000">index</span> <span style="color:#981a1a">*</span> <span style="color:#116644">2</span>;
<span style="color:#770088">if</span> ( (<span style="color:#000000">index</span> <span style="color:#981a1a">+</span> <span style="color:#116644">1</span><span style="color:#981a1a"><</span> <span style="color:#000000">length</span> <span style="color:#981a1a">-</span> <span style="color:#116644">1</span>) <span style="color:#981a1a">&&</span> (<span style="color:#000000">input</span>[<span style="color:#000000">index</span> <span style="color:#981a1a">+</span> <span style="color:#116644">1</span>] <span style="color:#981a1a">==</span> <span style="color:#000000">input</span>[<span style="color:#000000">index</span> <span style="color:#981a1a">+</span> <span style="color:#116644">2</span>]) )
<span style="color:#000000">output</span>[<span style="color:#000000">index</span> <span style="color:#981a1a">+</span> <span style="color:#116644">1</span>] <span style="color:#981a1a">=</span> <span style="color:#116644">1</span>;
<span style="color:#770088">else</span>
<span style="color:#000000">output</span>[<span style="color:#000000">index</span> <span style="color:#981a1a">+</span> <span style="color:#116644">1</span>] <span style="color:#981a1a">=</span> <span style="color:#116644">0</span>;
}
<span style="color:#008855">int</span> <span style="color:#0000ff">find_repeats</span>(<span style="color:#008855">int*</span> <span style="color:#000000">device_input</span>, <span style="color:#008855">int</span> <span style="color:#000000">length</span>, <span style="color:#008855">int*</span> <span style="color:#000000">device_output</span>) {
<span style="color:#aa5500">// CS149 TODO:</span>
<span style="color:#aa5500">//</span>
<span style="color:#aa5500">// Implement this function. You will probably want to</span>
<span style="color:#aa5500">// make use of one or more calls to exclusive_scan(), as well as</span>
<span style="color:#aa5500">// additional CUDA kernel launches.</span>
<span style="color:#aa5500">// </span>
<span style="color:#aa5500">// Note: As in the scan code, the calling code ensures that</span>
<span style="color:#aa5500">// allocated arrays are a power of 2 in size, so you can use your</span>
<span style="color:#aa5500">// exclusive_scan function with them. However, your implementation</span>
<span style="color:#aa5500">// must ensure that the results of find_repeats are correct given</span>
<span style="color:#aa5500">// the actual array length.</span>
<span style="color:#aa5500">// 这段和前面的意思一样：强调2的幂次处理、结果的正确性</span>
<span style="color:#770088">const</span> <span style="color:#008855">int</span> <span style="color:#000000">threadsPerBlock</span> <span style="color:#981a1a">=</span> <span style="color:#116644">512</span>;
<span style="color:#770088">const</span> <span style="color:#008855">int</span> <span style="color:#000000">blocks</span> <span style="color:#981a1a">=</span> (<span style="color:#000000">length</span> <span style="color:#981a1a">+</span> <span style="color:#000000">threadsPerBlock</span> <span style="color:#981a1a">-</span> <span style="color:#116644">1</span>) <span style="color:#981a1a">/</span> <span style="color:#000000">threadsPerBlock</span> <span style="color:#981a1a">/</span> <span style="color:#116644">2</span>;
<span style="color:#aa5500">// Place 1s wherever there is a repeat</span>
<span style="color:#000000">isrepeat1</span><span style="color:#981a1a"><<<</span><span style="color:#000000">blocks</span>,<span style="color:#000000">threadsPerBlock</span><span style="color:#981a1a">>>></span>(<span style="color:#000000">device_input</span>, <span style="color:#000000">length</span>, <span style="color:#000000">device_output</span>);
<span style="color:#000000">cudaDeviceSynchronize</span>();
<span style="color:#000000">isrepeat2</span><span style="color:#981a1a"><<<</span><span style="color:#000000">blocks</span>,<span style="color:#000000">threadsPerBlock</span><span style="color:#981a1a">>>></span>(<span style="color:#000000">device_input</span>, <span style="color:#000000">length</span>, <span style="color:#000000">device_output</span>);
<span style="color:#aa5500">// Use scan to count the total</span>
<span style="color:#000000">exclusive_scan</span>(<span style="color:#000000">device_input</span>, <span style="color:#000000">length</span>, <span style="color:#000000">device_output</span>);
<span style="color:#008855">int*</span> <span style="color:#000000">output</span>;
<span style="color:#000000">output</span> <span style="color:#981a1a">=</span> (<span style="color:#008855">int*</span>)<span style="color:#000000">malloc</span>(<span style="color:#770088">sizeof</span>(<span style="color:#008855">int</span>));
<span style="color:#000000">cudaMemcpy</span>(<span style="color:#000000">output</span>, <span style="color:#000000">device_output</span> <span style="color:#981a1a">+</span> <span style="color:#000000">length</span><span style="color:#981a1a">-</span><span style="color:#116644">1</span>, <span style="color:#770088">sizeof</span>(<span style="color:#008855">int</span>), <span style="color:#000000">cudaMemcpyDeviceToHost</span>);
<span style="color:#000000">printf</span>(<span style="color:#aa1111">"total %d\n"</span>, <span style="color:#000000">output</span>[<span style="color:#116644">0</span>]);
<span style="color:#008855">int</span> <span style="color:#000000">repeats</span> <span style="color:#981a1a">=</span> <span style="color:#000000">output</span>[<span style="color:#116644">0</span>];
<span style="color:#000000">free</span>(<span style="color:#000000">output</span>);
<span style="color:#770088">return</span> <span style="color:#000000">repeats</span>;
}
</span></span>总结：
主要是两部分，第一部分是识别是否满足条件，第二部分是并行统计满足条件的个数
第一部分与上课讲的奇偶排序算法十分类似
第二部分需要我们学习一下并行前缀和算法
Part3 CUDA并行渲染器
这一部分只做了题目翻译
因为我发现自己没能完成环境的配置，所以决定转战asst1
希望题目翻译部分能对下一届的学弟学妹有帮助
原题翻译：
此部分分为理论部分、实现部分、建议与帮助、思路与实现
理论部分：
初始代码的目录 /render 包含一个已经实现的用于绘制彩色圆圈的渲染器。可使用以下命令行构建并运行渲染：
<span style="background-color:#f8f8f8"><span style="color:#333333"> ./render <span style="color:#0000cc">-r</span> cpuref rgb。</span></span>程序将输出包含三个圆圈的图像文件并保存为output_0000.ppm。使用命令行运行渲染器./render -r cpuref snow。输出图像将是下雪。PPM图像可以通过预览直接在OSX上查看。Windows好像需要独立的查看器
注： 你还可以使用 -i 选项将渲染器输出直接发送到显示器而不是文件。http://atechyblog.blogspot.com/2014/12/google-cloud-compute-x11-forwarding.html（参考
初始代码包含两个版本的渲染器：在 refRenderer.cpp 中实现的顺序单线程C++参考实现，以及 cudaRenderer.cu 中不正确的并行 CUDA 实现。
建议通过检查 refRenderer.cpp 中的参考实现来熟悉渲染器代码库的结构。
在呈现第一帧之前调用方法setup，完成初始化等工作。
在您的 CUDA 加速渲染器中，setup方法可能需要包含所有渲染器的初始化代码（分配缓冲区等）。
render在每一帧都执行一次，负责将所有圆绘制到输出图像中。
渲染器的另一个主要函数 advanceAnimation 也每帧调用一次。它负责更新圆的位置和速度。
您无需在此作业中修改高级动画。
渲染器接受圆圈数组（3D 位置、速度、半径、颜色）作为输入。渲染每一帧的基本顺序算法为：
先对于每个圆更新其位置和颜色向量
然后对于每个圆，计算其边界
对于边界内的每个点，计算是否位于圆内（仅中心点位于园内才认为该点位于圆内
如果其位于园内，就计算该圆对该点像素的共享，并把它加到该像素点上
<span style="background-color:#f8f8f8"><span style="color:#333333"> Clear image
for each circle
update position and velocity
for each circle
compute screen bounding box
for all pixels in bounding box
compute pixel center point
if center point is within the circle
compute color of circle at point
blend contribution of circle into image for this pixel</span></span>下面介绍使用圆中点测试计算圆像素覆盖率的基本算法。请注意，仅当像素的中心位于圆内时，圆才会为输出像素提供颜色。
渲染器的一个重要细节是它渲染半透明的圆圈。因此，任何一个像素的颜色不是单个圆圈的颜色，而是混合所有与像素重叠的半透明圆的贡献的结果（注意上面伪代码的“混合贡献”部分）。
渲染器通过红色 （R）、绿色 （G）、蓝色 （B） 和不透明度 （alpha） 值 （RGBA） 的 4 元组表示圆的颜色。Alpha = 1 对应一个完全不透明的圆。Alpha = 0 对应于一个完全透明的圆圈。要将带有颜色 （P_r、P_g、P_b） 的像素点和一个带有颜色 （C_r、C_g、C_b C_alpha） 的半透明圆重合，则该点处的新像素值应当为：
<span style="background-color:#f8f8f8"><span style="color:#333333"> <span style="color:#000000">result_r</span> <span style="color:#981a1a">=</span> <span style="color:#000000">C_alpha</span> <span style="color:#981a1a">*</span> <span style="color:#000000">C_r</span> <span style="color:#981a1a">+</span> (<span style="color:#116644">1.0</span> <span style="color:#981a1a">-</span> <span style="color:#000000">C_alpha</span>) <span style="color:#981a1a">*</span> <span style="color:#000000">P_r</span>
<span style="color:#000000">result_g</span> <span style="color:#981a1a">=</span> <span style="color:#000000">C_alpha</span> <span style="color:#981a1a">*</span> <span style="color:#000000">C_g</span> <span style="color:#981a1a">+</span> (<span style="color:#116644">1.0</span> <span style="color:#981a1a">-</span> <span style="color:#000000">C_alpha</span>) <span style="color:#981a1a">*</span> <span style="color:#000000">P_g</span>
<span style="color:#000000">result_b</span> <span style="color:#981a1a">=</span> <span style="color:#000000">C_alpha</span> <span style="color:#981a1a">*</span> <span style="color:#000000">C_b</span> <span style="color:#981a1a">+</span> (<span style="color:#116644">1.0</span> <span style="color:#981a1a">-</span> <span style="color:#000000">C_alpha</span>) <span style="color:#981a1a">*</span> <span style="color:#000000">P_b</span></span></span>需要注意的是这种算法是不具有交换性的，我们必须保证自己的代码按照正确的顺序对图像进行渲染
接下来是一段初始代码解读，我将这部分与后续的思路放一起
实现部分：
我们实现的代码应当满足以下两种性质
原子性：
对一个像素点的RGB三个值的调整必须是原子的
即调整的过程中不应当有别的圆插入进来修改RGB
顺序性：
必须严格按照圆的顺序性进行渲染，也就是说，如果输入时，圆1比圆2先输入，我们的代码要保证圆1一定比圆2先渲染
当原子性与顺序性有不被满足时，渲染结果将类似左边的两张图片，一个表现出了错误的颜色，一个表现出了错误的层级关系
我们的目标是编写最快、最正确的 CUDA 渲染器实现。我们可以采用任何您认为合适的方法，但必须遵守上面指定的原子性和顺序要求。
建议与帮助：
入手建议：
一个好的出发点是通读 cudaRenderer.cu 并找出其中导致不符合原子性和顺序性的部分。
首先应当看看 CudaRenderer：render 如何启动 CUDA 内核内核 RenderCircles。
kernelRenderCircles是所有工作发生的地方。可使用make编译程序。然后运行 ./render -r cuda rand10k，它应该显示带有 10K 圆圈的图像。将此（不正确的）图像与运行 ./render -r cpuref rand10k 的顺序代码生成的正确图像进行比较。
实现建议：
首先重写 CUDA 启动代码实现，使其在并行运行时逻辑正确（先对了再说
然后考虑优化实现的性能问题
此时，对作业的真正思考开始了......
circleBoxTest.cu_inl中提供给您的圆相交框可用于测试代码正确性。建议您使用这些子例程。
原文此处有运行所用的命令行程序，不做翻译
以下部分为翻译的正确性检查标准
正确性检查：为了检测程序的正确性，渲染有一个方便的--check选项。此选项将参考 CPU 渲染器的顺序版本与 CUDA 渲染器一起运行，然后比较生成的图像以确保正确性。还会打印 CUDA 渲染器实现所花费的时间。
我们总共提供了五个圆圈数据集，您将根据这些数据集进行评分。您的代码必须通过我们所有的正确性测试。若要检查代码的正确性和性能分数，请在 /render 目录中运行 ./checker.py（请注意.py扩展名）。
注意：如果你的渲染器偶尔它是正确的。这不会改变当前 CUDA 渲染器通常不正确的事实。
“参考时间”是我们参考解决方案在您当前计算机上的性能。“您的时间”是当前 CUDA 渲染器解决方案的性能，其中 （F） 表示不正确的解决方案。与这些参考实施相比，您的成绩将取决于您的实施性能（请参阅评分指南）。
除了您的代码外，我们希望您提交一份清晰、高级的实现方式描述，以及您如何获得此解决方案的简要说明。具体说明您在此过程中尝试的方法，以及您如何确定如何优化代码（例如，您执行了哪些度量来指导您的优化工作？
您应该在文章中提及的工作方面包括： 复制为解决方案生成的分数表，并指定运行代码的计算机。 描述您如何分解问题以及如何将工作分配给 CUDA 线程块和线程。 描述解决方案中发生同步的位置。 您采取了哪些步骤（如果有）来降低通信要求（例如，同步或主内存带宽要求）？ 简要描述您是如何得出最终解决方案的。一路上你还尝试了哪些其他方法。是否有效？
评分准则就不翻译了0.0
在给出的几条建议中，有一部分是基于CUDA库的建议，但是那个是基于平台实现的，所以我只翻译了与平台实现无关的部分，并给出一些实现部分的理解
此作业中有两个潜在的并行轴。一个轴是跨像素的并行性，另一个轴是跨圆的并行性（前提是对重叠的圆遵守排序要求）。解决方案将需要利用并行性的两种 tpyes，可能在计算的不同部分。
Hint 1 事实上考虑我们课上学到的并行加速比的知识，如果我们按圆作为处理单位，即每次处理一个圆，其中的串行部分依赖圆的个数，不妨设为n，如果按每个像素点作为处理单位，其中的串行部分依赖单个像素点上堆叠的圆的个数，不妨设为m，易知m一定小于等于n，所以从最大的性能角度理论来看，一定是按照像素点要快于按圆
但是，由于按照像素点我们实际上需要为每个像素点都开辟一个线程，而按圆处理则为每个圆开辟一个线程，使用按照像素点处理时的访存开销要大于按圆处理（由于CUDA实现中自带的L2 Cache和warp访存机制
所以我看到一种实现是将以上两种思路并行，占据双倍的运算资源，一旦某种方法得到结果，就输出....
exclusiveScan.cu_inl 中提供的共享内存前缀和操作可能对您完成此任务很有价值（并非所有解决方案都可能选择使用它）。我们在共享内存中的两个大小的幂数组上提供了独占前缀和的实现。提供的代码不适用于非二次幂输入，并且它还要求线程块中的线程数为数组的大小。请阅读代码中的注释。
我们考虑之前优化前缀和的过程
前缀和可视为 a[n] = a[n-1] +an
此处可视为 a[n] = alpha * a[n-1] + (1-alpha) * an
如果能对该运算进行并行加速，这意味着我们可以在logn的时间内完成一个像素点的计算，且一定能保证顺序性
由于没有以原子方式执行映像更新操作逻辑的 CUDA 语言原语，您将如何确保映像更新的原子性？构造全局内存原子操作锁定是一种解决方案，但请记住，即使映像更新是原子更新，也必须按所需的顺序执行更新。我们建议您首先考虑确保并行解决方案中的顺序，然后再考虑解决方案中的原子性问题（如果它仍然存在）。
不难发现，如果能保证顺序性，原子性很大可能已经被满足
除非你的代码存在漏算的可能
最后一部分是CUDA中的bug调试帮助，可参考原文件，不做翻译
思路与实现：
先来读他的错误代码
600多行除去初始化和运行最重要的就是这个函数
<span style="background-color:#f8f8f8"><span style="color:#333333"> <span style="color:#000000">__global__</span> <span style="color:#008855">void</span> <span style="color:#0000ff">kernelRenderCircles</span>() {
<span style="color:#008855">int</span> <span style="color:#000000">index</span> <span style="color:#981a1a">=</span> <span style="color:#000000">blockIdx</span>.<span style="color:#000000">x</span> <span style="color:#981a1a">*</span> <span style="color:#000000">blockDim</span>.<span style="color:#000000">x</span> <span style="color:#981a1a">+</span> <span style="color:#000000">threadIdx</span>.<span style="color:#000000">x</span>;
<span style="color:#aa5500">//计算出自己处理第几个圆</span>
<span style="color:#770088">if</span> (<span style="color:#000000">index</span> <span style="color:#981a1a">>=</span> <span style="color:#000000">cuConstRendererParams</span>.<span style="color:#000000">numCircles</span>)
<span style="color:#770088">return</span>;
<span style="color:#008855">int</span> <span style="color:#000000">index3</span> <span style="color:#981a1a">=</span> <span style="color:#116644">3</span> <span style="color:#981a1a">*</span> <span style="color:#000000">index</span>;
<span style="color:#aa5500">// read position and radius</span>
<span style="color:#aa5500">// 读位置和输入</span>
<span style="color:#000000">float3</span> <span style="color:#000000">p</span> <span style="color:#981a1a">=</span> <span style="color:#981a1a">*</span>(<span style="color:#000000">float3</span><span style="color:#981a1a">*</span>)(<span style="color:#981a1a">&</span><span style="color:#000000">cuConstRendererParams</span>.<span style="color:#000000">position</span>[<span style="color:#000000">index3</span>]);
<span style="color:#008855">float</span> <span style="color:#000000">rad</span> <span style="color:#981a1a">=</span> <span style="color:#000000">cuConstRendererParams</span>.<span style="color:#000000">radius</span>[<span style="color:#000000">index</span>];
<span style="color:#aa5500">// compute the bounding box of the circle. The bound is in integer</span>
<span style="color:#aa5500">// screen coordinates, so it's clamped to the edges of the screen.</span>
<span style="color:#aa5500">// 计算边界</span>
<span style="color:#008855">short</span> <span style="color:#000000">imageWidth</span> <span style="color:#981a1a">=</span> <span style="color:#000000">cuConstRendererParams</span>.<span style="color:#000000">imageWidth</span>;
<span style="color:#008855">short</span> <span style="color:#000000">imageHeight</span> <span style="color:#981a1a">=</span> <span style="color:#000000">cuConstRendererParams</span>.<span style="color:#000000">imageHeight</span>;
<span style="color:#008855">short</span> <span style="color:#000000">minX</span> <span style="color:#981a1a">=</span> <span style="color:#000000">static_cast</span><span style="color:#981a1a"><</span><span style="color:#008855">short</span><span style="color:#981a1a">></span>(<span style="color:#000000">imageWidth</span> <span style="color:#981a1a">*</span> (<span style="color:#000000">p</span>.<span style="color:#000000">x</span> <span style="color:#981a1a">-</span> <span style="color:#000000">rad</span>));
<span style="color:#008855">short</span> <span style="color:#000000">maxX</span> <span style="color:#981a1a">=</span> <span style="color:#000000">static_cast</span><span style="color:#981a1a"><</span><span style="color:#008855">short</span><span style="color:#981a1a">></span>(<span style="color:#000000">imageWidth</span> <span style="color:#981a1a">*</span> (<span style="color:#000000">p</span>.<span style="color:#000000">x</span> <span style="color:#981a1a">+</span> <span style="color:#000000">rad</span>)) <span style="color:#981a1a">+</span> <span style="color:#116644">1</span>;
<span style="color:#008855">short</span> <span style="color:#000000">minY</span> <span style="color:#981a1a">=</span> <span style="color:#000000">static_cast</span><span style="color:#981a1a"><</span><span style="color:#008855">short</span><span style="color:#981a1a">></span>(<span style="color:#000000">imageHeight</span> <span style="color:#981a1a">*</span> (<span style="color:#000000">p</span>.<span style="color:#000000">y</span> <span style="color:#981a1a">-</span> <span style="color:#000000">rad</span>));
<span style="color:#008855">short</span> <span style="color:#000000">maxY</span> <span style="color:#981a1a">=</span> <span style="color:#000000">static_cast</span><span style="color:#981a1a"><</span><span style="color:#008855">short</span><span style="color:#981a1a">></span>(<span style="color:#000000">imageHeight</span> <span style="color:#981a1a">*</span> (<span style="color:#000000">p</span>.<span style="color:#000000">y</span> <span style="color:#981a1a">+</span> <span style="color:#000000">rad</span>)) <span style="color:#981a1a">+</span> <span style="color:#116644">1</span>;
<span style="color:#aa5500">// a bunch of clamps. Is there a CUDA built-in for this?</span>
<span style="color:#aa5500">// 在里面吗？</span>
<span style="color:#008855">short</span> <span style="color:#000000">screenMinX</span> <span style="color:#981a1a">=</span> (<span style="color:#000000">minX</span> <span style="color:#981a1a">></span> <span style="color:#116644">0</span>) <span style="color:#981a1a">?</span> ((<span style="color:#000000">minX</span> <span style="color:#981a1a"><</span> <span style="color:#000000">imageWidth</span>) <span style="color:#981a1a">?</span> <span style="color:#000000">minX</span> : <span style="color:#000000">imageWidth</span>) : <span style="color:#116644">0</span>;
<span style="color:#008855">short</span> <span style="color:#000000">screenMaxX</span> <span style="color:#981a1a">=</span> (<span style="color:#000000">maxX</span> <span style="color:#981a1a">></span> <span style="color:#116644">0</span>) <span style="color:#981a1a">?</span> ((<span style="color:#000000">maxX</span> <span style="color:#981a1a"><</span> <span style="color:#000000">imageWidth</span>) <span style="color:#981a1a">?</span> <span style="color:#000000">maxX</span> : <span style="color:#000000">imageWidth</span>) : <span style="color:#116644">0</span>;
<span style="color:#008855">short</span> <span style="color:#000000">screenMinY</span> <span style="color:#981a1a">=</span> (<span style="color:#000000">minY</span> <span style="color:#981a1a">></span> <span style="color:#116644">0</span>) <span style="color:#981a1a">?</span> ((<span style="color:#000000">minY</span> <span style="color:#981a1a"><</span> <span style="color:#000000">imageHeight</span>) <span style="color:#981a1a">?</span> <span style="color:#000000">minY</span> : <span style="color:#000000">imageHeight</span>) : <span style="color:#116644">0</span>;
<span style="color:#008855">short</span> <span style="color:#000000">screenMaxY</span> <span style="color:#981a1a">=</span> (<span style="color:#000000">maxY</span> <span style="color:#981a1a">></span> <span style="color:#116644">0</span>) <span style="color:#981a1a">?</span> ((<span style="color:#000000">maxY</span> <span style="color:#981a1a"><</span> <span style="color:#000000">imageHeight</span>) <span style="color:#981a1a">?</span> <span style="color:#000000">maxY</span> : <span style="color:#000000">imageHeight</span>) : <span style="color:#116644">0</span>;
<span style="color:#008855">float</span> <span style="color:#000000">invWidth</span> <span style="color:#981a1a">=</span> <span style="color:#116644">1.f</span> <span style="color:#981a1a">/</span> <span style="color:#000000">imageWidth</span>;
<span style="color:#008855">float</span> <span style="color:#000000">invHeight</span> <span style="color:#981a1a">=</span> <span style="color:#116644">1.f</span> <span style="color:#981a1a">/</span> <span style="color:#000000">imageHeight</span>;
<span style="color:#aa5500">// for all pixels in the bonding box</span>
<span style="color:#aa5500">// 计算新的像素点值</span>
<span style="color:#770088">for</span> (<span style="color:#008855">int</span> <span style="color:#000000">pixelY</span><span style="color:#981a1a">=</span><span style="color:#000000">screenMinY</span>; <span style="color:#000000">pixelY</span><span style="color:#981a1a"><</span><span style="color:#000000">screenMaxY</span>; <span style="color:#000000">pixelY</span><span style="color:#981a1a">++</span>) {
<span style="color:#000000">float4</span><span style="color:#981a1a">*</span> <span style="color:#000000">imgPtr</span> <span style="color:#981a1a">=</span> (<span style="color:#000000">float4</span><span style="color:#981a1a">*</span>)(<span style="color:#981a1a">&</span><span style="color:#000000">cuConstRendererParams</span>.<span style="color:#000000">imageData</span>[<span style="color:#116644">4</span> <span style="color:#981a1a">*</span> (<span style="color:#000000">pixelY</span> <span style="color:#981a1a">*</span> <span style="color:#000000">imageWidth</span> <span style="color:#981a1a">+</span> <span style="color:#000000">screenMinX</span>)]);
<span style="color:#770088">for</span> (<span style="color:#008855">int</span> <span style="color:#000000">pixelX</span><span style="color:#981a1a">=</span><span style="color:#000000">screenMinX</span>; <span style="color:#000000">pixelX</span><span style="color:#981a1a"><</span><span style="color:#000000">screenMaxX</span>; <span style="color:#000000">pixelX</span><span style="color:#981a1a">++</span>) {
<span style="color:#000000">float2</span> <span style="color:#000000">pixelCenterNorm</span> <span style="color:#981a1a">=</span> <span style="color:#000000">make_float2</span>(<span style="color:#000000">invWidth</span> <span style="color:#981a1a">*</span> (<span style="color:#000000">static_cast</span><span style="color:#981a1a"><</span><span style="color:#008855">float</span><span style="color:#981a1a">></span>(<span style="color:#000000">pixelX</span>) <span style="color:#981a1a">+</span> <span style="color:#116644">0.5f</span>),
<span style="color:#000000">invHeight</span> <span style="color:#981a1a">*</span> (<span style="color:#000000">static_cast</span><span style="color:#981a1a"><</span><span style="color:#008855">float</span><span style="color:#981a1a">></span>(<span style="color:#000000">pixelY</span>) <span style="color:#981a1a">+</span> <span style="color:#116644">0.5f</span>));
<span style="color:#000000">shadePixel</span>(<span style="color:#000000">index</span>, <span style="color:#000000">pixelCenterNorm</span>, <span style="color:#000000">p</span>, <span style="color:#000000">imgPtr</span>);
<span style="color:#000000">imgPtr</span><span style="color:#981a1a">++</span>;
}
}
}</span></span>很显然，这即保证不了圆的顺序性（圆和圆之间并行
也保证不了原子性（由于圆和圆之间串行，所以在修改一个点的RGB时，可能另一个圆也在修改这个点
很自然的我们想到了基于像素点的并行
一个简单的思路是按圆串行，每个圆都并行计算其影响的像素点，并使用同步机制保证满足原子性
<span style="background-color:#f8f8f8"><span style="color:#333333"> <span style="color:#000000">__global__</span> <span style="color:#008855">void</span> <span style="color:#0000ff">kernelRenderCircles</span>() {
<span style="color:#000000">__shared__</span> <span style="color:#000000">uint</span> <span style="color:#000000">circleInBox</span>[<span style="color:#000000">BLOCK_SIZE</span>];
<span style="color:#000000">__shared__</span> <span style="color:#000000">uint</span> <span style="color:#000000">circleOrder</span>[<span style="color:#000000">BLOCK_SIZE</span>];
<span style="color:#000000">__shared__</span> <span style="color:#000000">uint</span> <span style="color:#000000">circleScan</span>[<span style="color:#000000">BLOCK_SIZE</span>];
<span style="color:#000000">__shared__</span> <span style="color:#000000">uint</span> <span style="color:#000000">scratch_pad</span>[<span style="color:#116644">2</span> <span style="color:#981a1a">*</span> <span style="color:#000000">BLOCK_SIZE</span>];
<span style="color:#008855">int</span> <span style="color:#000000">imageWidth</span> <span style="color:#981a1a">=</span> <span style="color:#000000">cuConstRendererParams</span>.<span style="color:#000000">imageWidth</span>;
<span style="color:#008855">int</span> <span style="color:#000000">imageHeight</span> <span style="color:#981a1a">=</span> <span style="color:#000000">cuConstRendererParams</span>.<span style="color:#000000">imageHeight</span>;
<span style="color:#008855">float</span> <span style="color:#000000">invWidth</span> <span style="color:#981a1a">=</span> <span style="color:#116644">1.f</span> <span style="color:#981a1a">/</span> <span style="color:#000000">imageWidth</span>;
<span style="color:#008855">float</span> <span style="color:#000000">invHeight</span> <span style="color:#981a1a">=</span> <span style="color:#116644">1.f</span> <span style="color:#981a1a">/</span> <span style="color:#000000">imageHeight</span>;
<span style="color:#008855">int</span> <span style="color:#000000">tidx</span> <span style="color:#981a1a">=</span> <span style="color:#000000">threadIdx</span>.<span style="color:#000000">y</span> <span style="color:#981a1a">*</span> <span style="color:#000000">BLOCK_DIM</span> <span style="color:#981a1a">+</span> <span style="color:#000000">threadIdx</span>.<span style="color:#000000">x</span>;
<span style="color:#008855">int</span> <span style="color:#000000">num_circles</span> <span style="color:#981a1a">=</span> <span style="color:#000000">cuConstRendererParams</span>.<span style="color:#000000">numCircles</span>;
<span style="color:#008855">int</span> <span style="color:#000000">pixelX</span> <span style="color:#981a1a">=</span> <span style="color:#000000">blockIdx</span>.<span style="color:#000000">x</span> <span style="color:#981a1a">*</span> <span style="color:#000000">BLOCK_DIM</span> <span style="color:#981a1a">+</span> <span style="color:#000000">threadIdx</span>.<span style="color:#000000">x</span>;
<span style="color:#008855">int</span> <span style="color:#000000">pixelY</span> <span style="color:#981a1a">=</span> <span style="color:#000000">blockIdx</span>.<span style="color:#000000">y</span> <span style="color:#981a1a">*</span> <span style="color:#000000">BLOCK_DIM</span> <span style="color:#981a1a">+</span> <span style="color:#000000">threadIdx</span>.<span style="color:#000000">y</span>;
<span style="color:#008855">float</span> <span style="color:#000000">b_l</span> <span style="color:#981a1a">=</span> <span style="color:#000000">blockIdx</span>.<span style="color:#000000">x</span> <span style="color:#981a1a">*</span> <span style="color:#000000">BLOCK_DIM</span> <span style="color:#981a1a">*</span> <span style="color:#000000">invWidth</span>;
<span style="color:#008855">float</span> <span style="color:#000000">b_r</span> <span style="color:#981a1a">=</span> <span style="color:#000000">min</span>( <span style="color:#000000">b_l</span> <span style="color:#981a1a">+</span> <span style="color:#000000">BLOCK_DIM</span> <span style="color:#981a1a">*</span> <span style="color:#000000">invWidth</span>, <span style="color:#116644">1.0</span> );
<span style="color:#008855">float</span> <span style="color:#000000">b_b</span> <span style="color:#981a1a">=</span> <span style="color:#000000">blockIdx</span>.<span style="color:#000000">y</span> <span style="color:#981a1a">*</span> <span style="color:#000000">BLOCK_DIM</span> <span style="color:#981a1a">*</span> <span style="color:#000000">invHeight</span>;
<span style="color:#008855">float</span> <span style="color:#000000">b_t</span> <span style="color:#981a1a">=</span> <span style="color:#000000">min</span>( <span style="color:#000000">b_b</span> <span style="color:#981a1a">+</span> <span style="color:#000000">BLOCK_DIM</span> <span style="color:#981a1a">*</span> <span style="color:#000000">invHeight</span>, <span style="color:#116644">1.0</span> );
<span style="color:#000000">float4</span><span style="color:#981a1a">*</span> <span style="color:#000000">imgPtr</span>;
<span style="color:#000000">float2</span> <span style="color:#000000">pixelCenterNorm</span>;
<span style="color:#000000">float3</span> <span style="color:#000000">p</span>;
<span style="color:#008855">float</span> <span style="color:#000000">rad</span>;
<span style="color:#770088">for</span>( <span style="color:#008855">int</span> <span style="color:#000000">i</span> <span style="color:#981a1a">=</span> <span style="color:#116644">0</span>; <span style="color:#000000">i</span> <span style="color:#981a1a"><</span> <span style="color:#000000">num_circles</span>; <span style="color:#000000">i</span> <span style="color:#981a1a">+=</span> <span style="color:#000000">BLOCK_SIZE</span>) {
<span style="color:#008855">int</span> <span style="color:#000000">circle</span> <span style="color:#981a1a">=</span> <span style="color:#000000">tidx</span> <span style="color:#981a1a">+</span> <span style="color:#000000">i</span>;
<span style="color:#770088">if</span>( <span style="color:#000000">circle</span> <span style="color:#981a1a"><</span> <span style="color:#000000">num_circles</span> ) {
<span style="color:#000000">p</span> <span style="color:#981a1a">=</span> <span style="color:#981a1a">*</span>(<span style="color:#000000">float3</span><span style="color:#981a1a">*</span>)(<span style="color:#981a1a">&</span><span style="color:#000000">cuConstRendererParams</span>.<span style="color:#000000">position</span>[<span style="color:#116644">3</span> <span style="color:#981a1a">*</span> <span style="color:#000000">circle</span>]);
<span style="color:#000000">rad</span> <span style="color:#981a1a">=</span> <span style="color:#000000">cuConstRendererParams</span>. <span style="color:#000000">radius</span>[<span style="color:#000000">circle</span>];
<span style="color:#000000">circleInBox</span>[<span style="color:#000000">tidx</span>] <span style="color:#981a1a">=</span> <span style="color:#000000">circleInBoxConservative</span>(<span style="color:#000000">p</span>.<span style="color:#000000">x</span>, <span style="color:#000000">p</span>.<span style="color:#000000">y</span>, <span style="color:#000000">rad</span>, <span style="color:#000000">b_l</span>, <span style="color:#000000">b_r</span>, <span style="color:#000000">b_t</span>, <span style="color:#000000">b_b</span>);
} <span style="color:#770088">else</span> {
<span style="color:#000000">circleInBox</span>[<span style="color:#000000">tidx</span>] <span style="color:#981a1a">=</span> <span style="color:#116644">0</span>;
}
<span style="color:#000000">__syncthreads</span>();
<span style="color:#000000">sharedMemExclusiveScan</span>(<span style="color:#000000">tidx</span>, <span style="color:#000000">circleInBox</span>, <span style="color:#000000">circleScan</span>, <span style="color:#000000">scratch_pad</span>, <span style="color:#000000">BLOCK_SIZE</span>);
<span style="color:#770088">if</span>( <span style="color:#000000">circleInBox</span>[<span style="color:#000000">tidx</span>] <span style="color:#981a1a">==</span> <span style="color:#116644">1</span>) {
<span style="color:#000000">circleOrder</span>[<span style="color:#000000">circleScan</span>[<span style="color:#000000">tidx</span>]] <span style="color:#981a1a">=</span> <span style="color:#000000">tidx</span> <span style="color:#981a1a">+</span> <span style="color:#000000">i</span>;
}
<span style="color:#000000">__syncthreads</span>();
<span style="color:#770088">if</span>( <span style="color:#000000">pixelX</span> <span style="color:#981a1a"><</span> <span style="color:#000000">imageWidth</span> <span style="color:#981a1a">&&</span> <span style="color:#000000">pixelY</span> <span style="color:#981a1a"><</span> <span style="color:#000000">imageHeight</span>) {
<span style="color:#008855">int</span> <span style="color:#000000">totalCircles</span> <span style="color:#981a1a">=</span> <span style="color:#000000">circleScan</span>[<span style="color:#000000">BLOCK_SIZE</span> <span style="color:#981a1a">-</span> <span style="color:#116644">1</span>] <span style="color:#981a1a">+</span> <span style="color:#000000">circleInBox</span>[<span style="color:#000000">BLOCK_SIZE</span> <span style="color:#981a1a">-</span> <span style="color:#116644">1</span>];
<span style="color:#770088">for</span>( <span style="color:#008855">int</span> <span style="color:#000000">j</span> <span style="color:#981a1a">=</span> <span style="color:#116644">0</span>; <span style="color:#000000">j</span> <span style="color:#981a1a"><</span> <span style="color:#000000">totalCircles</span>; <span style="color:#000000">j</span><span style="color:#981a1a">++</span> ) {
<span style="color:#000000">circle</span> <span style="color:#981a1a">=</span> <span style="color:#000000">circleOrder</span>[<span style="color:#000000">j</span>];
<span style="color:#000000">p</span> <span style="color:#981a1a">=</span> <span style="color:#981a1a">*</span>(<span style="color:#000000">float3</span><span style="color:#981a1a">*</span>)(<span style="color:#981a1a">&</span><span style="color:#000000">cuConstRendererParams</span>.<span style="color:#000000">position</span>[<span style="color:#116644">3</span><span style="color:#981a1a">*</span><span style="color:#000000">circle</span>]);
<span style="color:#000000">rad</span> <span style="color:#981a1a">=</span> <span style="color:#000000">cuConstRendererParams</span>.<span style="color:#000000">radius</span>[<span style="color:#000000">circle</span>];
<span style="color:#000000">imgPtr</span> <span style="color:#981a1a">=</span> (<span style="color:#000000">float4</span><span style="color:#981a1a">*</span>)(<span style="color:#981a1a">&</span><span style="color:#000000">cuConstRendererParams</span>.<span style="color:#000000">imageData</span>[<span style="color:#116644">4</span> <span style="color:#981a1a">*</span> (<span style="color:#000000">pixelY</span> <span style="color:#981a1a">*</span> <span style="color:#000000">imageWidth</span> <span style="color:#981a1a">+</span> <span style="color:#000000">pixelX</span>)]);
<span style="color:#000000">pixelCenterNorm</span> <span style="color:#981a1a">=</span> <span style="color:#000000">make_float2</span>(<span style="color:#000000">invWidth</span> <span style="color:#981a1a">*</span> (<span style="color:#000000">static_cast</span><span style="color:#981a1a"><</span><span style="color:#008855">float</span><span style="color:#981a1a">></span>(<span style="color:#000000">pixelX</span>) <span style="color:#981a1a">+</span> <span style="color:#116644">0.5f</span>),
<span style="color:#000000">invHeight</span> <span style="color:#981a1a">*</span> (<span style="color:#000000">static_cast</span><span style="color:#981a1a"><</span><span style="color:#008855">float</span><span style="color:#981a1a">></span>(<span style="color:#000000">pixelY</span>) <span style="color:#981a1a">+</span> <span style="color:#116644">0.5f</span>));
<span style="color:#000000">shadePixel</span>(<span style="color:#000000">circle</span>, <span style="color:#000000">pixelCenterNorm</span>, <span style="color:#000000">p</span>, <span style="color:#000000">imgPtr</span>);
}
}
}
}</span></span>
1万+
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
