float* x; float* y; // … int block_size = 256; int grid_size = (N + block_size - 1) / block_size; add_gpu<<<grid_size, block_size>>>(N, x, y); // … } Ok, about the kernel parameters += stride) y[i] = x[i] + y[i]; } TEST_CASE("cppcon-1", "[CUDA]") { // … add_gpu<<<grid_size, block_size>>>(N, x, y); // … } Ok, about the kernel parameters 10 |Memory“In a typical PC

blockDim • 当前板块的编号： blockIdx • 总的板块数量： gridDim • 线程 (thread) ：并行的最小单位 • 板块 (block) ：包含若干个线程 • 网格 (grid) ：指整个任务，包含若干个板块 • 从属关系：线程＜板块＜网格 • 调用语法： <<>> 区分板块和线程有点麻烦？“扁平化”他们！ • 你可能觉得纳闷，既然已经有线程可以并行了 刚刚的 for 循环是串行的，我们可以把线 程数量调为 n ，然后用 threadIdx.x 作为 i 索引。这样就实现了，每个线程负责给数 组中一个元素的赋值。 小技巧：网格跨步循环（ grid-stride loop ） • 无论调用者指定了多少个线程 （ blockDim ），都能自动根据给定的 n 区间循环，不会越界，也不会漏掉几个元 素。 • 这样一个 for 循环非常符合 gridDim ，看起来非常方便。 本方法出自英伟达官方博客： https://developer.nvidia.com/blog/cuda-pro-tip-write-flexible-kernels-grid-stride-loops/ 第 4 章： C++ 封装 std::vector 的秘密：第二模板参数 • 你知道吗？ std::vector 作为模板类，其实有两个模板参数： std::vector

0 码力 | 142 页 | 13.52 MB | 1 年前

3

被实际分配。比如这里我们分配了 16GB 内 存，但是只访问了他的前 4KB ，这样只有一个页被分配，所以非常快。 • 其实操作系统惰性分配的特性，也是 SPGrid （ Sparsely-Paged-Grid ）得以实现的基础 ，他利用 mmap 分配比机器大得多的内存（比如 2048*2028*1024 的三维网格），然后 在里面索引，这样就相当于利用硬件的分页机制实现了稀疏数据结构，既能高效利用内存 插桩这个操作，就是说在结构网格中，从一个点往周围固定范围读取值， 并根据一定权重累加，然后用于修改自身的值。 • 比如求一个场的梯度、散度、旋度、拉普拉斯。如果场是用结构网格 （ structured grid ）表示，那就是一个插桩操作。 • 插桩的内核（ kernel ）指的就是这个“周围范围”的形状（如右图三个例子） 和每个地方读取到值对修改自身值的权重等信息。 • 个人认为，图像处理中的模

npm downloads) ? • 10) ? nativescript-paypal (3,852) • Marcel Kloubert • 9) ? nativescript-grid-view (4,006) • Peter Staev • 8) ⛅ nativescript-floatingactionbutton (4,916) • Brad Martin

with “native” Label TextField Repeater SegmentedBar Layouts (Traditional) Absolut e Dock Grid Stack Wrap Layouts (Flexbox) GridLayout GridLayout Label TextField StackLayout Platform-Specific

提前释放一段页面。 • 除此之外， mmap 还有一个好处，他会保证其内存（被读 取访问时）是零初始化的。 配合莫顿分块， AOSOA 等第七课的技术，就得到 SPGrid(sparse-paged grid) SPGrid 还支持自适应的网格 SPGrid 的利弊 • 优点：平坦直观，适合插桩，顺序访问，自适应网格 。 • 缺点：尺寸受限，操作系统挂钩，依赖 x86 硬件机 制。 • 顺便一提，

* 3) % 4 号线程。 这样总体来看每个线程分到的块的位置是随机的， 从而由于正太分布数量越大方差越小的特点，每个 线程分到的总工作量大概率是均匀的。 • GPU 上称为网格跨步循环（ grid-stride loop ）。 1 1 1 1 2 2 2 3 3 4 3 3 4 4 2 4 tbb::static_partitioner 创建了 4 个线程 4