fern::Interval (y, out.y_start, out.y_start + out.y_len, l fern::Compute( fern::Producer(Convolution Input Filters Convolution 65 }) )) template void gemm(Matrix A,Matrix B,Matrix fern::Interval void conv(image input, image filter, int StrideArg, image out);Convolution Input Filters Convolution 66 }) )) template void gemm(Matrix A,Matrix B,Matrix fern::Interval void conv(image input, image filter, int StrideArg, image out);Convolution Input Filters Convolution 67 }) )) template void gemm(Matrix A,Matrix B,Matrix fern::Interval

problems • Intra-register sorting • Fast linear median-of-seven filter • Fast small-kernel convolution • Faster (?) UTF-8 to UTF-32 conversion (with AVX2) • No heavy code, but lots of pictures • Small-Kernel Convolution 3 CppCon 2020 - Adventures in SIMD ThinkingCopyright © 2020 Bob Steagall K E W B C O M P U T I N G Convolution • f is a signal • g is a kernel • Output f*g is the convolution • Every CppCon 2020 - Adventures in SIMD Thinking 4Copyright © 2020 Bob Steagall K E W B C O M P U T I N G Convolution CppCon 2020 - Adventures in SIMD Thinking 5 S = s0 s1 s2 s3 s4 s5 s6

如果你试图分离声明和定义，调用另一个文件里 的 __device__ 或 __global__ 函数，就会出错 。 分离 __device__ 函数的声明和定义：解决 • 开启 CMAKE_CUDA_SEPARABLE_COMPILATION 选 项（设为 ON ），即可启用分离声明和定义的支持。 • 不过我还是建议把要相互调用的 __device__ 函数放在 同一个文件，这样方便编译器自动内联优化（第四课讲 glDispatchComputeIndirect 的 API 和这个很像，但毕竟没有 CUDA 可以直接在核函数里调用核函数并指定参数这么方便…… 不过，这个功能同样需要开启 CUDA_SEPARABLE_COMPILATION 。 第 2 章：内存管理 如何从核函数里返回数据？ • 我们试着把 kernel 的返回类型声明为 int ，试 图从 GPU 返回数据到 CPU 。 • 但发现这样做会在编译期出错，为什么？

problems • Intra-register sorting • Fast linear median-of-seven filter • Fast small-kernel convolution • Faster (?) UTF-8 to UTF-32 conversion (with AVX2) • No heavy code, but lots of pictures •

波操作，就属于插桩。有的 插桩内核各轴向是对称的（比如高斯模糊），有的是单单往一个方向延伸 很长（比如径向模糊），有的内核是正方形（箱滤波）。 • 人工智障圈子里好像管这个叫卷积（ convolution ）什么的。 X 方向插桩：手动预取下一缓存行 • X 方向的插桩，因为最内层有另一个长度仅为 nblur*2+1 的顺序读取，让 CPU 误以为 t++ 是我 们想要的顺序读取，然而