ECE408 Lecture 7 Convolution and Constant Memory
这次回顾ECE408 Lecture 7,这次介绍了卷积和Constant Memory。
课程主页:
https://wiki.illinois.edu/wiki/display/ECE408
搬运视频:
https://www.youtube.com/playlist?list=PL6RdenZrxrw-UKfRL5smPfFFpeqwN3Dsz
本讲目标
学习卷积,一种重要的并行计算模式;
广泛应用于信号、图像和视频处理;
许多科学和工程应用中的计算基础;
卷积神经网络(CNN)的关键组成部分;
重要的GPU技术;
利用高速缓存;
卷积计算公式连续和离散情形:
\begin{aligned}
f(x) * g(x)&=\int_{-\infty}^{\infty} f(\tau) \cdot g(x-\tau) d \tau \\
f[x] * g[x]&=\sum_{k=-\infty}^{\infty} f[k] \cdot g[x-k]
\end{aligned}卷积应用
一种在信号处理、数字记录、图像处理、视频处理、计算机视觉和机器学习中以各种形 ...
ECE408 Lecture 6 Generalized Tiling and DRAM Bandwidth
这次回顾ECE408 Lecture 6,这次介绍了Tiling的一般形式以及DRAM带宽。
课程主页:
https://wiki.illinois.edu/wiki/display/ECE408
搬运视频:
https://www.youtube.com/playlist?list=PL6RdenZrxrw-UKfRL5smPfFFpeqwN3Dsz
处理任意大小的矩阵
第5讲中的分块矩阵乘法内核只能处理维度是分块维度的倍数的矩阵;
但是,实际应用程序需要处理任意大小的矩阵;
我们可以将行和列填充(添加元素)成块大小的倍数,但是会有很大的空间和数据传输时间开销;
我们可以在代码中添加显式检查来处理边界;
例子考虑TILE_WIDTH = 2, width = 3的矩阵乘法。
考虑载入Block(0, 0)的第2个tile,边界情况需要特殊处理:
计算:
考虑载入Block(1, 1)的第1个tile,边界情况需要特殊处理:
计算:
2x2例子中的主要情形
计算有效$P$元素但使用无效输入的线程
Block(0, 0)的第2个tile;
计算无效$P$元素的线程 ...
ECE408 Lecture 5 Locality and Tiled Matrix Multiply
这次回顾ECE408 Lecture 5,这次介绍了Tiled矩阵乘法。
课程主页:
https://wiki.illinois.edu/wiki/display/ECE408
搬运视频:
https://www.youtube.com/playlist?list=PL6RdenZrxrw-UKfRL5smPfFFpeqwN3Dsz
本讲目标
学习评估全局内存访问的性能影响;
为MP3做准备
tiled矩阵乘法;
了解tiling的好处;
在内存带宽为150 GB/s的设备上性能如何?
所有线程访问其输入矩阵元素所在的全局内存
每个浮点乘加 (2 fp ops) 两次内存访问(8字节);
每个FLOP使用4B内存;
150 GB/s将代码限制在37.5 GFLOPS;
150 / 4 = 37.5;
实际代码运行速度约为25 GFLOPS;
需要大幅减少内存访问以接近超过1,000 GFLOPS的峰值;
即主要性能瓶颈是全局内存的读取:
通过重用避免带宽(BW)瓶颈
$M$和$N$的每个元素在计算$P$时使用$\mathrm{Width}$次;
为避免BW限 ...
ECE408 Lecture 4 CUDA Memory Model
这次回顾ECE408 Lecture 4,这次介绍了CUDA内存模型。
课程主页:
https://wiki.illinois.edu/wiki/display/ECE408
搬运视频:
https://www.youtube.com/playlist?list=PL6RdenZrxrw-UKfRL5smPfFFpeqwN3Dsz
本讲内容
了解CUDA线程可访问内存的基本特征;
准备MP-2
基本矩阵乘法;
学习评估全局内存访问的性能影响;
内容回顾控制(分支)divergence
分支的主要性能问题是divergence
单个warp中的线程采用不同的路径;
不同的执行路径在当前的GPU中被序列化;
常见情况:当分支条件是线程ID的函数时出现divergence
if (threadIdx.x % 2) { }
这为warp中的线程创建了两个不同的控制路径;
有divergence(50%的线程什么都不做);
if ((threadIdx.x / WARP_SIZE) % 2) { }
还创建了两个不同的控制路径;
但 ...
ECE408 Lecture 3 CUDA Parallel Execution Model
这次回顾ECE408 Lecture 3,这次的介绍了CUDA并行执行模型。
课程主页:
https://wiki.illinois.edu/wiki/display/ECE408
搬运视频:
https://www.youtube.com/playlist?list=PL6RdenZrxrw-UKfRL5smPfFFpeqwN3Dsz
本讲内容
了解更多关于CUDA线程的多维逻辑组织;
学习使用控制结构,例如kernel中的循环;
学习线程调度、延迟容限和硬件占用的概念;
内容回顾这里回顾向量加法。
基本情形矩阵加法,假设$n=1000$,block size为256:
vecAdd<<<ceil(N/256.0), 256>>>(...)
i = blockIdx.x * blockDim.x + threadIdx.x;
if (i<N) C[i] = A[i] + B[i];
图示:
红色部分解释了为什么要增加i<N的判断:为了在数组大小不整除block size时结果依然正确。
该代码对应一个线程处理1个元素,那么处理两个 ...
ECE408 Lecture 2 Introduction to Parallel Computing and CUDA
这次回顾ECE408 Lecture 2,这次的内容是并行计算核CUDA的介绍。
课程主页:
https://wiki.illinois.edu/wiki/display/ECE408
搬运视频:
https://www.youtube.com/playlist?list=PL6RdenZrxrw-UKfRL5smPfFFpeqwN3Dsz
本讲内容
数据并行计算的基本概念;
CUDA C/C++ 编程接口的基本特性;
数据并行计算示例:将彩色图像转换为灰度图像
CUDA执行模型
典型的计算密集型C/C++ 代码
串行或适度并行的部分;
高度并行的部分;
串行部件->CPU(或Host)
高度并行部分->GPU(或Device)
并行线程数组CUDA内核作为线程网格(grid/数组)执行
网格中的所有线程都运行相同的内核代码;
单程序多数据(SPMD 模型);
每个线程都有一个唯一的索引,用于计算内存地址和做出控制决策;
CUDA的逻辑执行模型每个CUDA 内核
由grid执行;
grid是thread blocks的3D数组;
thread b ...
ECE408 Lecture 1 Introduction to the Course
这次回顾ECE408 Lecture 1,这次的内容是课程介绍。
课程主页:
https://wiki.illinois.edu/wiki/display/ECE408
搬运视频:
https://www.youtube.com/playlist?list=PL6RdenZrxrw-UKfRL5smPfFFpeqwN3Dsz
课程目标
学习对大规模并行处理器进行编程并实现
高性能;
功能性和可维护性;
后续的可扩展性;
技术相关
并行编程基础;
并行算法的原理和模式;
编程API、工具和技术;
处理器架构特性和限制;
杀手级应用;
课本/笔记课程的配套书有两本:
D. Kirk and W. Hwu, “Programming MassivelyParallel Processors – A Hands-on Approach,”Morgan Kaufman Publisher, 3rd edition, 2016,ISBN 978-012381472
NVIDIA, NVidia CUDA C Programming Guide,version 7.5 or ...
Chapter 2 Asymptotic Analysis笔记 Part 2 (Stanford Machine Learning Theory)
这里回顾第二章的第二部分,这部分对上一部分的定理举了一些实例。
课程主页:
https://web.stanford.edu/class/stats214/
课程视频:
https://www.youtube.com/playlist?list=PLoROMvodv4rP8nAmISxFINlGKSK4rbLKh
课件:
https://github.com/tengyuma/cs229m_notes/blob/main/master.pdf
Chapter 2 Asymptotic Analysis2.1 经验风险最小化的渐近例子定理2.4
在定理2.1的假设之外,假设存在参数模型$P(y \mid x ; \theta), \theta \in \Theta$,使得对于某个$\theta_{\star} \in \Theta$,$\left\{y^{(i)} \mid x^{(i)}\right\}_{i=1}^n \sim P\left(y^{(i)} \mid x^{(i)} ; \theta_{\star}\right)$。假设我们的损失函数为$\ell\lef ...
Chapter 2 Asymptotic Analysis笔记 Part 1 (Stanford Machine Learning Theory)
这里回顾第二章的第一部分,这部分主要讨论训练集无限的情况下ERM的理论上界。
课程主页:
https://web.stanford.edu/class/stats214/
课程视频:
https://www.youtube.com/playlist?list=PLoROMvodv4rP8nAmISxFINlGKSK4rbLKh
课件:
https://github.com/tengyuma/cs229m_notes/blob/main/master.pdf
Chapter 2 Asymptotic Analysis在本章中,我们使用渐近方法(即假设训练样本数$n\to \infty$)来得到ERM的上界。 然后,我们将这些结果实例化为损失函数为最大似然的情况,并讨论渐近的局限性。(在以后的章节中,我们将假设$n$有限并提供非渐近分析。)
2.1 经验风险最小化的渐近对于ERM的渐进分析,我们希望证明其excess risk有如下上界(其中$c$是依赖于问题的常数):
L(\hat{\theta})-\inf _{\theta \in \Theta} L(\theta) ...
Chapter 1 Supervised Learning Formulations笔记 (Stanford Machine Learning Theory)
最近开始学习一些机器学习理论,目前从斯坦福的Machine Learning Theory (CS229M/STATS214)入手,课程的优点是课件,视频比较全,缺点是作业无法公开获得。这里对第一章做一个总结,这一章介绍了监督学习的范式,引入了population risk,excess risk和empirical risk等概念。
课程主页:
https://web.stanford.edu/class/stats214/
课程视频:
https://www.youtube.com/playlist?list=PLoROMvodv4rP8nAmISxFINlGKSK4rbLKh
课件:
https://github.com/tengyuma/cs229m_notes/blob/main/master.pdf
Chapter 1 Supervised Learning Formulations在本章中,我们将建立有监督学习的标准理论公式,并介绍经验风险最小化 (ERM) 范式。
1.1 有监督学习基本概念在监督学习中,我们有输入和输出,输入属于输入空间$\mathcal ...
ECE408 环境配置以及Lab 0
最近开始学习Cuda,找了一些资料,比较下来发现ECE408资料最全,后续应该会把资料整理一下,这次先介绍环境配置和实验0。
课程主页:
https://wiki.illinois.edu/wiki/display/ECE408
搬运视频:
https://www.youtube.com/playlist?list=PL6RdenZrxrw-UKfRL5smPfFFpeqwN3Dsz
参考资料:
https://wiki.illinois.edu/wiki/display/ECE408/Labs+and+Project
环境配置环境配置包括Cuda环境以及课程自带的依赖库。
Cuda环境这部分略过,网上的资料很多。
课程依赖库课程依赖库需要下载并编译:
git clone https://github.com/abduld/libwb.git
cd libwb
make all
后续会介绍如何使用该库。
Lab 0下载作业:
git clone https://github.com/aschuh703/ECE408.git
接着是依赖库的使用。
依赖库实际上是头文件,在编 ...
Softmax极值项关于温度的导数
最近研究Softmax时发现一个有意思的结论,这里简单记录下。
Softmax和温度给定:
\mathbf x =[x_1,\ldots, x_d]^{\top} \in \mathbb R^{d}Softmax定义为:
\mathrm{Softmax}(\mathbf x)_i =\frac{\exp(x_i)}{\sum_j\exp(x_j)}在实际中还有温度参数$\tau$,主要是控制Softmax的陡峭程度,计算公式为:
\mathrm{Softmax}(\mathbf x;\tau )_i =\frac{\exp(x_i/\tau)}{\sum_j\exp(x_j/\tau)}Softmax极值项关于$\tau$有一个单调性的结论:
Softmax最大分量关于$\tau$递减;
Softmax最小分量关于$\tau$递增;
关于温度的导数为了证明结论,考虑如下函数的导数:
f_i(\tau)\triangleq \mathrm{Softmax}(\mathbf x;\tau )_i =\frac{\exp(x_i/\tau)}{\sum_j\exp(x_j/\t ...