Posts

硬件限制 这是容易量化的方面。目前CUDA编程指南的附录F列出了一些硬件限制，这些限制限制了内核启动时每块可以有多少个线程。如果你超过了这些限制，你的内核将无法运行。这些限制可以粗略地概括为： 每个区块不能超过 $512$ / $1024$ 个线程（分别是计算能力1.x或2.x及以后的计算能力 每个块的最大尺寸限制在 $[512, 512, 64]$ / $[1024, 1024, 64]$（计算能力1.x/2.x及以后的计算能力 每个块消耗的寄存器总数不能超过 $8k/16k/32k/64k/32k/64k/32k/64k$ （计算能力 $1.0,1.1/1.2,1.3/2.x-3.0/3.2/3.5-5.2/5.3/6-6.1/6.2/7.0$ 每个块不能消耗超过 $16kb/48kb/96kb$ 的共享内存（计算能力 $1.x/2.x-6.2/7.0$ 如果你保持在这些限制之内，任何你能成功编译的内核都会无错误地启动。 性能调教 这是需要经验的一部分。在上述的硬件约束条件下，你选择的每块线程数可以而且确实影响到硬件上运行的代码性能。每个代码的表现都是不同的，唯一真正的方法是通过仔细的基准测试和剖析来量化它。但还是那句话，非常粗略地总结一下： 每个区块的线程数应该是wrap大小的整数倍，在目前所有的硬件上都是 $32$ GPU上的每个流式多处理器必须有足够的active wraps来充分隐藏架构的的所有不同内存和指令流水线延迟，以实现最大吞吐量。这里的正确做法是尝试实现最佳的硬件占用率 CUDA内置函数 上述指出了块的大小是如何影响性能的，并提出了一种基于占用率最大化的通用启发式选择方法。在不想提供选择块大小的标准的情况下，值得一提的是，CUDA 6.5+包括几个新的运行时函数来帮助占用率的计算和启动配置1。 其中一个有用的函数是cudaOccupancyMaxPotentialBlockSize，它启发式地计算了一个能达到最佳占用率的块大小。该函数提供的值可以作为手动优化参数的起点。下面是一个例子： /************************/ /* TEST KERNEL FUNCTION */ /************************/ __global__ void MyKernel(int *a, int *b, int *c, int N) { int idx = threadIdx.x + blockIdx.x * blockDim.x; if (idx < N) { c[idx] = a[idx] + b[idx]; } } /********/ /* MAIN */ /********/ void main() { const int N = 1000000; int blockSize; // The launch configurator returned block size int minGridSize; // The minimum grid size needed to achieve the maximum occupancy for a full device launch int gridSize; // The actual grid size needed, based on input size int* h_vec1 = (int*) malloc(N*sizeof(int)); int* h_vec2 = (int*) malloc(N*sizeof(int)); int* h_vec3 = (int*) malloc(N*sizeof(int)); int* h_vec4 = (int*) malloc(N*sizeof(int)); int* d_vec1; cudaMalloc((void**)&d_vec1, N*sizeof(int)); int* d_vec2; cudaMalloc((void**)&d_vec2, N*sizeof(int)); int* d_vec3; cudaMalloc((void**)&d_vec3, N*sizeof(int)); for (int i=0; i<N; i++) { h_vec1[i] = 10; h_vec2[i] = 20; h_vec4[i] = h_vec1[i] + h_vec2[i]; } cudaMemcpy(d_vec1, h_vec1, N*sizeof(int), cudaMemcpyHostToDevice); cudaMemcpy(d_vec2, h_vec2, N*sizeof(int), cudaMemcpyHostToDevice); float time; cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop); cudaEventRecord(start, 0); cudaOccupancyMaxPotentialBlockSize(&minGridSize, &blockSize, MyKernel, 0, N); // Round up according to array size gridSize = (N + blockSize - 1) / blockSize; cudaEventRecord(stop, 0); cudaEventSynchronize(stop); cudaEventElapsedTime(&time, start, stop); printf("Occupancy calculator elapsed time: %3....

使用多线程提高并行性 同步的代价 我们研究对一列整数 $0, …, n - 1$ 求和，我们将序列划分成 $t$ 个不相交的的区域，给 $t$ 个线程每个分配一个区域。将线程的和放入一个变量中，并且我们使用互斥锁来保护这个变量。 use std::{ sync::{Arc, Mutex}, thread, time::Instant, }; fn main() { let args = std::env::args().collect::<Vec<String>>(); if args.len() != 3 { panic!("Usage: {} <nthreads> <log_nelems>", args[0]); } let nthreads: usize = args[1].parse().unwrap(); let log_nelems: usize = args[2].parse().unwrap(); let nelems = 1_usize << log_nelems; let nelems_per_thread = nelems / nthreads; let gsum = Arc::new(Mutex::new(0)); let now = Instant::now(); let mut handlers = vec!...

不正确地使用new和delete 无论我们如何努力，要释放所有动态分配的内存是非常困难的。即使我们能做到这一点，也往往不能安全地避免出现异常。让我们看一个简单的例子。 void SomeMethod() { ClassA *a = new ClassA; SomeOtherMethod(); // it can throws an execption delete a; } 如果SomeOtherMethod抛出了异常，那么a对象永远不会被删除。下面的例子展示了一个更加安全同时又更简洁的实现，使用了在C++11提出的std::unique_ptr。 void SomeMethod() { std::unique_ptr<ClassA> a(new ClassA); SomeOtherMethod(); } 无论发生什么，当a退出作用域的时候，它会被释放。 然而，这仅仅是C++中这种错误最简单的例子，还有很多例子表明delete应该在其他地方调用，也许是在外层函数或者另一个线程中。这就是为什么应该避免使用new和delete，而应该使用适当的智能指针。 被忘记的虚析构函数 这是最常见的错误之一，如果派生类中有动态内存分配，将会导致派生类的内存泄漏。这里有一些例子，当一个类不打算用于继承，并且它的大小和性能是至关重要的。虚析构函数或任何其他虚函数在类在类中引入了额外的数据，即指向虚函数表的指针，这使得类的任何实例的大小变大。 然而，在大多数情况下，类可以被继承，即使它的初衷并非如此。因此，在声明一个类的时候，添加一个虚析构函数是一个非常好的做法。否则，如果一个类由于性能的原因必须不包含虚函数，那么在类的声明文件里面加上一个注释，说明这个类不应该被继承，是一个很好的做法。避免这个问题的最佳选择之一是使用一个支持在创建类时创建虚析构函数的IDE。 关于这个问题，还有一点是来自标准库的类或模板。它们不是用来继承的，也没有一个虚析构函数。例如，如果我们创建了一个公开继承自std::string的新的增强字符串类，就可能有人错误地使用它与std::string的指针或引用，从而导致内存泄漏。 class MyString : public std::string { ~MyString() {} }; int main() { std::string *s = new MyString(); delete s; // May not invoke the destructor defined in MyString } 为了避免这样的问题，重用标准库中的类或模板的一个更安全的方法是使用私有继承1或组合。 用delete或智能指针删除一个数组 创建动态大小的临时数组往往是必要的。当它们不再需要时，释放分配的内存是很重要的。这里的问题是，C++需要带有[]括号的特殊删除操作符，这一点很容易被遗忘。delete[]操作符不仅会删除分配给数组的内存，而且会首先调用数组中所有对象的析构函数。对原始类型使用不带[]括号的删除操作符也是不正确的，尽管这些类型没有析构函数，每个编译器都不能保证一个数组的指针会指向数组的第一个元素，所以使用不带[]括号的delete也会导致未定义的行为。 在数组中使用智能指针，如unique_ptr<T>, shared_ptr，也是不正确的。当这样的智能指针从作用域中退出时，它将调用不带[]括号的删除操作符，这将导致上面描述的同样问题。如果需要对数组使用智能指针，可以使用unique_ptr<T[]>的特殊化。 如果不需要引用计数的功能，主要是数组的情况，最优雅的方法是使用STL向量来代替。它们不只是负责释放内存，而且还提供额外的功能。...

为什么要有完美转发 下面是一个类工厂函数： template <typename T, typename Arg> std::shared_ptr<T> factory(Arg arg) { return std::shared_ptr<T>( new T(arg)); } 参数对象arg在上面的例子中是传值方式传递，这带来了生成额外临时对象1的代价，所以我们改成引用传递： template <typename T, typename Arg> std::shared_ptr<T> factory(Arg &arg) { return std::shared_ptr<T>( new T(arg)); } 但这种实现的问题是不能绑定右值实参。如factory<X>(42)将编译报错，进一步的，我们按常量引用来传递： template <typename T, typename Arg> std::shared_ptr<T> factory(const Arg &arg) { return std::shared_ptr<T>( new T(arg)); } 这种实现的问题是不能支持移动语义，形参使用右值引用可以解决完美转发问题。 引用折叠 在C++11之前，我们不能对一个引用类型继续引用，但C++由于右值引用的出现而放宽2了这一做法，从而产生了引用折叠规则，允许我们对引用进行引用，既能左引用，又能右引用。但是却遵循如下规则： 函数形参类型 实参类型 推导后函数形参类型 T& 左引用 T& T& 右引用 T& T&& 左引用 T& T&& 右引用 T&& 模板参数类型推导 对函数模板template<typename T>void foo(T&&);，应用上述引用折叠规则，可总结出以下结论： 如果实参是类型A的左值，则模板参数T的类型为A&，形参类型为A&； 如果实参是类型A的右值，则模板参数T的类型为A&&，形参类型为A&&。 这同样适用于类模板的成员函数模板的类型推导：...

Kernels CUDA C++对C++进行了扩展，允许程序员定义C++函数，称为内核，当被调用时，由 $N$ 个不同的CUDA线程并行执行 $N$ 次，而不是像普通C++函数那样只执行一次。 kernel是使用__global__声明定义的，对于特定的内核调用，执行该内核的CUDA线程数量是使用<<<...>>>执行配置语法指定的（C++语言扩展）。每个执行内核的线程都有一个唯一的线程ID，可以在内核内通过内置变量访问。 作为说明，下面的示例代码，使用内置变量threadIdx，将两个大小为 $N$ 的向量 $A$ 和 $B$ 相加，并将结果存入向量 $C$ 。 // Kernel definition __global__ void VecAdd(float* A, float* B, float* C) { int i = threadIdx.x; C[i] = A[i] + B[i]; } int main() { ... // Kernel invocation with N threads VecAdd<<<1, N>>>(A, B, C); ... } 在这里，执行VecAdd()的 $N$ 个线程中的每一个都执行了一次加法。 线程体系 为方便起见，threadIdx是一个 $3$ 分量的向量，因此可以用一维、二维或三维的线程索引来识别线程，形成一个一维、二维或三维的线程块，称为线程块。这提供了一种自然的方式来调用域中的元素进行计算，如矢量、矩阵或体积。 一个例子，下面的代码将两个大小为 $N\times N$ 的矩阵 $A$ 和 $B$ 相加，并将结果存入矩阵 $C$ 。...