摘要：
▶ 编程接口。参考 http://chenrudan.github.io/ ▶ Runtime API 为高层级管理接口，提供申请和释放设备内存，数据迁移，多 GPU 管理等。Driver API 为较低层级的控制接口，提供 CUDA 上下文（模拟设备主机进程），CUDA 模块（模拟设备动态加载库） 阅读全文
|
摘要：
▶ 编程接口。参考 http://chenrudan.github.io/ ▶ Runtime API 为高层级管理接口，提供申请和释放设备内存，数据迁移，多 GPU 管理等。Driver API 为较低层级的控制接口，提供 CUDA 上下文（模拟设备主机进程），CUDA 模块（模拟设备动态加载库） 阅读全文
摘要：
▶ 使用函数 cudaMallocPitch() 和配套的函数 cudaMemcpy2D() 来使用二维数组。C 中二维数组内存分配是转化为一维数组，连贯紧凑，每次访问数组中的元素都必须从数组首元素开始遍历；而 cuda 中这样分配的二维数组内存保证了数组每一行首元素的地址值都按照 256 或 51 阅读全文
摘要：
10. 执行配置优化 ● 一个 SM中，占用率 = 活动线程的数量 / 最大可能活动线程的数量。后者保存在设备属性的 maxThreadsPerMultiProcessor 分量中（GTX1070为例，该值为2048）。较高的占用率不代表计算效率很高，但是较低的占用率意味着内存延迟严重，还有改进的空 阅读全文
摘要：
0. APOD过程 ● 评估。分析代码运行时间的组成，对瓶颈进行并行化设计。了解需求和约束条件，确定应用程序的加速性能改善的上限。 ● 并行化。根据原来的代码，采用一些手段进行并行化，例如使用现有库，或加入一些预处理指令等。同时需要代码重构来暴露它们固有的并行性。 ● 优化。并行化完成后，需要通过优 阅读全文
摘要：
▶ 学习回调函数的基本概念，并在CUDA的任务流中插入基于CPU的主机函数，作为回调函数使用。 ▶ 源代码（合并了 3 个源文件，删掉了没有用到的部分） ● 输出结果 ▶ 涨姿势 ● 回调函数的使用：首先在 cuda_runtime_api.h 中给出了能作为回调函数的主机函数格式，然后给出了回调函 阅读全文
摘要：
原子操作。并且在静态代码和运行时编译两种条件下使用。 ▶ 源代码：静态使用 ▶ 源代码：即时编译 ▶ 输出结果： ▶ 涨姿势 ● 一个有趣的数列：命 x0 = 0，xn = xn-1 XOR n，则有 x4n == 4n，x4n+1 = 1, x4n+2 == 4n+3, x4n+3 == 0。当改 阅读全文
摘要：
在核函数中使用强制终止函数 assert()。并且在静态代码和运行时编译两种条件下使用。 ▶ 源代码：静态使用 即时编译版： ▶ 输出结果： ▶ 涨姿势： ● 在核函数中使用 assert( condition ) 来检查各线程中是否满足某条件。 若不满足条件 condition，则强制终止该线程， 阅读全文
摘要：
使用CUDA的 Driver API 来计算矩阵乘法。 ▶ 源代码： ▶ 输出结果： ▶ 涨姿势： ● 头文件 matrixMul.h 的内容： ● C++ 中 string 类的基本使用方法 ● class StopWatchInterface ，定义于 helper_timer.h 中用于计时的 阅读全文
摘要：
关于cuBLAS库中矩阵乘法相关的函数及其输入输出进行详细讨论。 ▶ 涨姿势： ● cuBLAS中能用于运算矩阵乘法的函数有4个，分别是 cublasSgemm（单精度实数）、cublasDgemm（双精度实数）、cublasCgemm（单精度复数）、cublasZgemm（双精度复数），它们的定义 阅读全文
摘要：
使用CUDA的线性代数库cuBLAS来计算矩阵乘法。这里主要记录调用规则，关于乘法函数中详细的参数说明和调用规则见另一篇随笔。 ▶ 源代码： ▶ 输出结果： ▶ 涨姿势： ● 代码依然很烂。多用了一个 h_C 根本没有用上，其作用被 h_CUBLAS 取代了，而且源代码中有free(h_C)却没有f 阅读全文
摘要：
矩阵乘法，使用一维线程块和共享内存。并且在静态代码和运行时编译两种条件下使用。 ▶ 源代码：静态使用 ▶ 源代码：运行时编译 ▶ 输出结果： ▶ 涨姿势： ● 程序写得很烂，各种声明、初始化杂糅。 ● 一个根据cuda错误种类返回错误描述的函数 ● 预编译命令展开循环 等价于 #pragma unr 阅读全文
摘要：
在核函数代码中加入并行线程执行（Parallel Thread eXecution，PTX），通过汇编指令获取得有关线程束的信息。并且在静态代码和运行时编译两种条件下使用。 ▶ 源代码：静态使用 ▶ 源代码：运行时编译 ▶ 输出结果： ▶ 涨姿势： ● 获取当前线程在线程束中的编号，即同意先乘数中的 阅读全文
摘要：
▶ 使用cuda内置无符号整数结构（__half2）及其汇编函数，计算两个向量的内积。 ▶ 源代码 ● 输出结果 ▶ 涨姿势 ● CUDA 无符号半精度整数，就是用 unsigned short 对齐到 2 Byte 来封装的 ● 关于 __inline__ 和 __forceinline__ 参考 阅读全文
摘要：
▶ 在OpenMP的多线程程序中，各线程分别调用CUDA进行计算。OpenMP的简单示例。 ▶ 源代码，OpenMP 出了点问题，没有正确输出结果 阅读全文
摘要：
▶ 使用 cuda 内置结构 cudaFuncAttributes 来观察核函数的共享内存、寄存器数量 ▶ 源代码 ● 输出结果： ▶ 涨姿势： ● cuda 使用扩展名为 .cuh 的头文件 ● cuda内置结构 cudaFuncAttributes 的定义： ● 通过使用cuda的内置结构和函数 阅读全文
摘要：
▶ 分离编译【留坑，在 Linux 上用命令行试一下】 ▶ 源代码： ● 输出结果： ▶ 涨姿势： ● cuda 内置的 int2 类型，整数有序对。涉及的定义如下： ● 警告函数和错误检查函数 阅读全文
摘要：
▶ 使用 clock() 函数在CUDA核函数内部进行计时，将核函数封装为PTX并在另外的代码中读取和使用。 ▶ 源代码：文件内建核函数计时 ● 输出结果，比较不同的 blockDim.x 和 threadDim.x 情况结果如下图表所示。 ▶ 涨姿势： ● 在核函数中也能使用 time.h 中的 阅读全文
摘要：
▶ CUDA 动态并行实现快排算法（单线程的递归调用） ▶ 源代码：动态并行递归调用线程块。要点：添加 -rdc=true 选项（生成 relocatable device code，相当于执行分离编译），以及链接库 cudadevrt.lib （用于动态并行，不同于运行时库 cudart.lib） 阅读全文
摘要：
▶ CPU - GPU 异步操作 ▶ 源代码 ● 输出结果： ▶ 新姿势： ● 调用主函数时的第0个参数作为程序名字符串，可以用于输出。 ● 在没有附加 flag 的情况下申请主机内存，注意使用cudaFreeHost释放 ● 记录 CPU 调用 CUDA 所用的时间 ● 查看GPU队列状态的函数 阅读全文
摘要：
三种不同的方法计算前缀和，并与CPU的结果进行了对比。 ▶ 结果如下图。第一种方法存在不可重现的bug，且仅当输入数组规模大于512时开始出现，原因未知。其他几种方法计算结果均正确，再计算较短的响亮的时候第三种方法（改良的收集 - 分发树法）效率最高，当向量长度远大于自己设定的阈值1024（单个线程 阅读全文
摘要：
展示了三种不同的GPU一维卷积方法，分别为简单（全局内存）卷积，含光环元素的共享内存方法，不含光环元素的共享内存方法。并且改进了CPU的一维卷积方案（不需要分边界情况单独处理）。 ▶ 输出结果如下图，计时部分有点问题（如何使用同一个 cudaEvent _t start, stop 对多个事件进行计 阅读全文
摘要：
▶ 采用两种方法计算任意长度的数组规约加法。第一种是将原数组收缩到一个固定长度的共享内存数组中，再利用二分规约计算求和，这里测试了固定长度分别取2，4，8……1024的情形；第二种方法是将原数组分段，每段使用各自的共享内存进行二分规约，结果汇总后再次分段，规约，直至数组长度减小到1，输出结果。 ● 阅读全文
摘要：
▶ 计算矩阵矩阵乘法 Am×n Bn×p == Cm×p 过程。 ▶ 原始矩阵乘法，一个线程计算结果矩阵中的一个元素。 1 #include <stdio.h> 2 #include <stdlib.h> 3 #include <malloc.h> 4 #include <time.h> 5 #in 阅读全文
摘要：
▶ P46。SPMD (Single-Program Multiple-Data) 单程序多数据，CUDA使用的并行编程风格。并行处理单元在数据的多个部分执行相同程序，但这些处理单元不用同时执行限购通的指令；SIMD (Single-Instruction Multiple-Data) 单指令多数据 阅读全文
摘要：
▶ 直接的矩阵加法，没有优化 ▶输出结果，对于超大型矩阵的加法计算，左下图为一维跳转，右下图为二维跳转。矩阵较大时初始化矩阵费时较多。超大矩阵的计算上GPU相对CPU产生了一定的优势，一维跳转比二维跳转稍快，因为省去了复杂的下标映射。 ▶ 对于较小的矩阵，GPU没有发挥出优势， 因为内存拷贝等方面耗 阅读全文
|