背景 最近在调研各种hashmap.. 发现ska::flat hash map性能优秀。。于是来看看代码。。 发现最大的特点是,ska::flat_hash_map使用了带probe count上限的robin hood hashing 相关概念 Distance_from_desired 对于采用了open addressing的hash实现，当插入发生冲突时，会以一定方式(如线性探测、平方探测等)来探测下一个可以插入的slot. 因而实际插入的slot位置与理想的slot位置通常不相同，这段距离定义为distance_from_desired 在没有冲突的理想情况下，所有distance_from_desired的值应该都为0 …

背景 起因是同事在实现int4的功能，结果流水线有一条死活过不了(gcc版本为4.8.5),一直core dump 经过初步排查，找出了如下最小可以复现的代码: 1 2#include <immintrin.h>3 4class Test{ 5 public: 6 Test(){ 7 tmp = _mm256_set_epi32(0,0,0,0,0,0,0,0); 8 } 9 private: 10 __m256i tmp; 11}; 12int main(){ 13 auto *tmp = new Test(); 14 return 0; 15} gcc版本为4.8.5 其中编译选项为 1g++ -std=c++11 …

名词说明 CUDA. 一般来说指的是CUDA SDK. 目前经常使用的是CUDA 8.0和CUDA 10.1两个版本. 8.0和10.1都是SDK的版本号. CUDNN. The NVIDIA CUDA® Deep Neural Network library (cuDNN). 是一个可以为神经网络提供GPU加速的库 compute capability. 是GPU的固有参数,可以理解为GPU的版本.越新的显卡该数值往往越高. tensorRT.NVIDIA TensorRT™ is an SDK for high-performance deep learning inference. 是一个深度学习推理库,旨在提供高性能的推理 …

**Halide is a programming language designed to make it easier to write high-performance image and array processing code on modern machines. ** halide有两个特性比较吸引人。一个是对于各种平台架构的支持。 * CPU architectures: X86, ARM, MIPS, Hexagon, PowerPC * Operating systems: Linux, Windows, macOS, Android, iOS, Qualcomm QuRT * GPU …

迫于生计，最近要学习halide 先去学习/复习一下常见的编译优化技巧。 loop unrolling，也就是循环展开，顾名思义，就是把循环展开来写。 normal loop: int x; for (x = 0; x < 100; x++) { delete(x); } after loop unrolling: int x; for (x = 0; x < 100; x += 5 ) { delete(x); delete(x + 1); delete(x + 2); delete(x + 3); delete(x + 4); } 循环展开是一种优化，可以手动实现也可以编译器自动实现。 为什么要将循环展开？ * …

CUDA 编程涉及到在不同的平台上同时运行代码:包含CPU的host 和包含GPU的device. 所以了解host和device的对性能优化是非常重要的。 2.1. Differences between Host and Device Threading resources host 上能同时运行的线程数目比较少（比如24个） device上能同时运行的线程数目比较多（数量级通常为1E3，1E4等） Threads 操作系统必须交换CPU执行通道上和下的线程以提供多线程功能。因此，上下文切换(当交换两个线程时)既慢又昂贵。 相比之下，GPU上的线程非常轻量级。在典型的系统中，成千上万的线程排队等待工作(每个线程有32个线程)。如 …

APOD指的是Assess, Parallelize, Optimize, Deploy 如图所示，APOD过程是一个循环的过程，每次只进行一部分，从A到P到O到D,然后再进行下一轮的APOD Assess 对于一个项目，第一步要做的是接触(Assess)项目，得到项目代码中每部分的执行时间。有了这部分内容，开发者就可以找到并行优化的瓶颈所在，并开始尝试用GPU加速。 根据Amdahl's and Gustafson's laws，开发者可以确定并行优化的性能上界。 Parallelize 找到瓶颈所在并确定了优化的目标和期望，开发者就可以优化代码了。调用一些如cuBLAS, cuFFT, or Thrust …

可以了解成并行版的STL(? 过了一遍nvidia的官方网文档 发现如果熟悉STL的话,thrust没什么太多好说的,看起来很简单... 不过还是开一篇记录一下,一段时间内估计要和cuda c++ 打交道,就当记录使用过程中遇到的问题吧.

由于发现cuda c++ 的 debug方式和c++ 差别很大,因此打算再开一篇,专门记录一些和error checking 以及debug有关的内容. Error checks in CUDA code can help catch CUDA errors at their source. There are 2 sources of errors in CUDA source code: 1. Errors from CUDA **API** calls. For example, a call to `cudaMalloc()` might fail. 2. Errors from CUDA **kernel** …

uodate:有毒吧。kernel中出问题原来是不会报错的。。。。 请教了组里的hust学长orz..、 学到了cuda-memcheck命令和cudaGetLastError来查看问题。。可以参考What is the canonical way to check for errors using the CUDA runtime API? 先放一波资料。 * <del>[An Even Easier Introduction to CUDA](https://devblogs.nvidia.com/even-easier-introduction-cuda/)</del> * …