摘要：▶ C/C++ 扩展 阅读全文

摘要：▶ 动态并行。 ● 动态并行直接从 GPU 上创建工作，可以减少主机和设备间数据传输，在设备线程中调整配置。有数据依赖的并行工作可以在内核运行时生成，并利用 GPU 的硬件调度和负载均衡。动态并行要求算法和程序要提前改进，消除递归、不规则的循环、结构或其他不适合并行的情况。 ● 动态并行的经典图 ● 阅读全文

摘要：▶ 纹理内存访问补充（见纹理内存博客 http://www.cnblogs.com/cuancuancuanhao/p/7809713.html） ▶ 计算能力 ● 不同计算能力的硬件对计算特性的支持。 ● 不同计算能力的硬件技术特性（重要）。 ● 浮点运算技术标准描述（原文） ■ All comp 阅读全文

摘要：▶ 升级到 Visual Studio 2017 和 CUDA 9.1 之后，直接编译以前的 CUDA C 程序出现了如下报错： ● 那两个“应输入表达式”是读不了 <<< 和 >>> 操作符； ● 最后那个错很直接，在 host_config.h 中有下列对 Visual Studio 的版本筛选 阅读全文

摘要：▶ 数学函数 ● 舍入函数，考虑被舍入参数有双精度浮点和单精度浮点，舍入方式有区别，舍入结果有整形、长整形和长长整形，所以共有以下舍入函数。 ● 推荐单精度浮点使用函数 rint() 系列（单指令计算），而不用函数 round() 系列（8 指令计算）。其他几个用于单精度浮点的舍入函数也是单指令计算 阅读全文

摘要：▶ 协作组，要求 cuda ≥ 9.0，一个简单的例子见 http://www.cnblogs.com/cuancuancuanhao/p/7881093.html ● 灵活调节需要进行通讯的线程组合（不一定是线程块或是线程束）的尺寸，在更多粒度上进行线程协作。 ● 协作组功能支持 CUDA 的各种 阅读全文

摘要：▶ 线程束表决函数（Warp Vote Functions） ● 用于同一线程束内各线程通信和计算规约指标。 ● 在设备代码的一个线程中调用 _all(predicate)，__any(mask, predicate)，__ballot(mask, predicate) 时，该线程所在的线程束中所有 阅读全文

摘要：▶ 可缓存只读操作（Read-Only Data Cache Load Function），定义在 sm_32_intrinsics.hpp 中。从地址 adress 读取类型为 T 的函数返回，T 可以是 char，short，int，long longunsigned char，unsigned 阅读全文

摘要：▶ 纹理内存读取函数 ▶ 表面内存读写函数 阅读全文

摘要：附录 A，CUDA计算设备 附录 B，C语言扩展 ▶ 函数的标识符 ● __device__，__global__ 和 __host__ ● 宏 __CUDA_ARCH__ 可用于区分代码的运行位置. ● __noinline__ 和 __forceinlie__ ■ __device__ 函数由便 阅读全文

摘要：▶ 图形互操作性，OpenGL 与 Direct3D 相关。（没学过，等待填坑） ▶ 版本号与计算能力 ● 计算能力（Compute Capability）表征了硬件规格，CUDA版本号表征了驱动接口和运行时接口的（软件）规格。 ● Driver API 的版本号定义在在驱动头文件中的宏 CUDA_ 阅读全文

摘要：▶ 表面内存使用 ● 创建 cuda 数组时使用标志 cudaArraySurfaceLoadStore 来创建表面内存，可以用表面对象（surface object）或表面引用（surface reference）来对其进行读写。 ● 使用 Surface Object API ■ 涉及的结构定义 阅读全文

摘要：▶ 纹理内存使用 ● 纹理内存使用有两套 API，称为 Object API 和 Reference API 。纹理对象（texture object）在运行时被 Object API 创建，同时指定了纹理单元。纹理引用（Tezture Reference）在编译时被 Reference API 创 阅读全文

摘要：▶ 编程接口。参考 http://chenrudan.github.io/ ▶ Runtime API 为高层级管理接口，提供申请和释放设备内存，数据迁移，多 GPU 管理等。Driver API 为较低层级的控制接口，提供 CUDA 上下文（模拟设备主机进程），CUDA 模块（模拟设备动态加载库） 阅读全文

摘要：▶ 使用函数 cudaMallocPitch() 和配套的函数 cudaMemcpy2D() 来使用二维数组。C 中二维数组内存分配是转化为一维数组，连贯紧凑，每次访问数组中的元素都必须从数组首元素开始遍历；而 cuda 中这样分配的二维数组内存保证了数组每一行首元素的地址值都按照 256 或 51 阅读全文

摘要：10. 执行配置优化 ● 一个 SM中，占用率 = 活动线程的数量 / 最大可能活动线程的数量。后者保存在设备属性的 maxThreadsPerMultiProcessor 分量中（GTX1070为例，该值为2048）。较高的占用率不代表计算效率很高，但是较低的占用率意味着内存延迟严重，还有改进的空 阅读全文

摘要：0. APOD过程 ● 评估。分析代码运行时间的组成，对瓶颈进行并行化设计。了解需求和约束条件，确定应用程序的加速性能改善的上限。 ● 并行化。根据原来的代码，采用一些手段进行并行化，例如使用现有库，或加入一些预处理指令等。同时需要代码重构来暴露它们固有的并行性。 ● 优化。并行化完成后，需要通过优 阅读全文

摘要：关于cuBLAS库中矩阵乘法相关的函数及其输入输出进行详细讨论。 ▶ 涨姿势： ● cuBLAS中能用于运算矩阵乘法的函数有4个，分别是 cublasSgemm（单精度实数）、cublasDgemm（双精度实数）、cublasCgemm（单精度复数）、cublasZgemm（双精度复数），它们的定义 阅读全文

摘要：三种不同的方法计算前缀和，并与CPU的结果进行了对比。 ▶ 结果如下图。第一种方法存在不可重现的bug，且仅当输入数组规模大于512时开始出现，原因未知。其他几种方法计算结果均正确，再计算较短的响亮的时候第三种方法（改良的收集 - 分发树法）效率最高，当向量长度远大于自己设定的阈值1024（单个线程 阅读全文

摘要：展示了三种不同的GPU一维卷积方法，分别为简单（全局内存）卷积，含光环元素的共享内存方法，不含光环元素的共享内存方法。并且改进了CPU的一维卷积方案（不需要分边界情况单独处理）。 ▶ 输出结果如下图，计时部分有点问题（如何使用同一个 cudaEvent _t start, stop 对多个事件进行计 阅读全文