万物一齐，而百兽率舞。

本文简略介绍CUDA编程的相关背景

异构计算

首先，异构计算就是CPU+X，这个X可以是GPU，也可以是NPU，FPGA，DSP等等，在CPU上运行的代码称为Host主机代码，X上运行的代码称为Device设备代码

其次，异构计算和同构计算的一个不同点在于，异构的处理要比同构复杂，因为涉及同构下自动处理的计算，控制，传输都要人工进行干预

GPU

CPU vs GPU

先看cpu的架构和gpu架构的不同：

先看几个共同的部分：

• 控制器Control
• 缓存器Cache
• 随机存取存储器DRAM
• 总线PCIe Bus
• 算术逻辑单元ALU

不需要过多的计组原理知识，直观来看，GPU弱化了控制器的能力，大大堆砌了ALU的数量，且为一组ALU（SM，简单理解为线程束）分配了单独的Control和cache，所以直觉上，对了逻辑简单，数据量大的任务，GPU更高效

低并行逻辑复杂的程序适合用CPU
高并行逻辑简单的大数据计算适合GPU

CPU和GPU线程的区别：

1. CPU线程是重量级实体，操作系统交替执行线程，线程上下文切换花销很大
2. GPU线程是轻量级的，GPU应用一般包含成千上万的线程，多数在排队状态，线程之间切换基本没有开销。
3. CPU的核被设计用来尽可能减少一个或两个线程运行时间的延迟，而GPU核则是大量线程，最大幅度提高吞吐量

GPU简介

对于Nvida，GPU分为4个系列：

• Tegra
• Geforce
• Quadro
• Tesla

针对不同的应用场景，比如Tegra用于嵌入式，Geforce是平时打游戏用到，Tesla主要用于计算，但目前看Geforce和Tesla貌似也没那么分得很清楚，只是某些系列上有容错的硬件部分

GPU也存在版本号，即其计算能力，但这和性能不是成正比的，只代表架构的更新，对于CUDA来说，新版的cuda只支持最近几个版本的GPU，参考：https://en.wikipedia.org/wiki/CUDA#GPUs_supported

真正评价GPU能力的分为两种，一种是计算性能指标，一种是容量指标：

• 计算性能主要看FLOPS，一般单位在T级别（TFLOPS），称每秒浮点数运算次数，分单精度双精度，根据不同的架构，双精度一般是单精度的1/N
• 容量指标上有：显存，带宽，核心数量等，特别的如果一个CUDA程序爆显存了，则不使用统一内存的情况下是无法运行程序的

CUDA

CUDA平台不是单单指软件或者硬件，而是建立在Nvidia GPU上的一整套平台，并扩展出多语言支持

对于CUDA API来说，分为driver和运行时两种API，二者互斥不能同时使用，性能没啥差距

• driver API 低级，复杂，灵活，一般用于配合其他语言调用CUDA
• 运行时简单，基于driver API，稍微高级一点

CUDA编程环境架构如下：

nvcc

nvcc之于cuda相当于普通c++之于g++。因为cuda程序中host和device的代码是混在一起的，nvcc会分离这两个部分分别处理：

其中，PTX相当于CUDA的汇编语言

nvdia-smi

可以通过nvidia-smi(Nvidia's system management interface)程序检查与设置设备。它包含在CUDA开发工具套装内。

CUDA的程序结构

一般CUDA程序分成下面这些步骤：

1. 分配GPU内存
2. 拷贝内存到设备
3. 调用CUDA内核函数来执行计算
4. 把计算完成数据拷贝回主机端
5. 内存销毁

参考

1. https://face2ai.com/program-blog/#GPU%E7%BC%96%E7%A8%8B%EF%BC%88CUDA%EF%BC%89

万物化作，萌区有状，盛衰之杀，变化之流也。

本文以helloworld为例，简述CUDA的线程结构

基本概念

即使是HelloWorld，Cuda程序也引入了一些不同的概念，确保了解《简略CUDA和GPU相关概念》中的内容在进行下去

从HelloWorld开始

下面是一个最简单的HelloWorld程序：

/**
 运行命令: cl ./hello_host_only.cpp
**/

#include <cstdio>
using namespace std;

int main() {
    printf("Host: Hello wolrd!");
    return 0;
}

在windows可在x64 Native Tools Command Prompt for VS 2022用cl来编译它

只有Host代码的CUDA版本HelloWolrd

/**
 运行命令: nvcc -arch=sm_62 hello_host_only.cu
 sm_62的来源:
 https://docs.nvidia.com/cuda/archive/12.9.0/cuda-compiler-driver-nvcc/index.html
 https://en.wikipedia.org/wiki/CUDA#GPUs_supported
*/

#include <cstdio>
using namespace std;

int main() {
    printf("Host: Hello wolrd!");
    return 0;
}

代码完全没有变化，只有文件后缀名发生了变化

但这里性质发生了变化，代码分为了Host代码部分和Device代码部分由nvcc分别处理

nvcc指定运算架构

注意到注释部分，nvcc可以指定虚拟和真实架构的编号，可以通过链接查看本地设备可以使用的标识

如果需要匹配多个架构，可以使用多个-gencode -arch 虚拟架构 -code 真实架构来编译产生不同的输出，最后真实架构不能落后与虚拟架构

使用CUDA线程

从这里正式开始添加Device代码：

/**
 运行命令: nvcc  XX.cu
*/
#include<cuda.h>
#include<cuda_runtime.h>
#include<cstdio>

__global__ void hello_from_gpu()
{
    printf("Device: Hello World.");
}
int main()
{
    const int grid_size=1;
    const int block_size=1;
    hello_from_gpu<<<grid_size,block_size>>>();

    cudaDeviceSynchronize();
    return 0;
}

注意四点：

• include包含了两个cuda相关头文件，这不是必须要添加的，nvcc会自动引入，其中包含了cstdlib的内容，但没有包含cstdio的内容，最后device程序是不认iostream的
• __global___就是device代码，即核函数的标志
• cudaDeviceSynchronize的功能是同步host和device，也就是说gpu上的代码和cpu上的代码是异步运行的
• <<<grid_size,block_size>>>指定了用多少线程运行该核函数

使用多线程的CUDA程序

再进一步改进程序

/**
 运行命令: nvcc  XX.cu
*/
#include<cuda.h>
#include<cuda_runtime.h>
#include<cstdio>

__global__ void hello_from_gpu()
{
    const int bid = blockIdx.x;
    const int tid = threadIdx.x;
    printf("Device: Hello World from block %d and thread %d\n",bid,tid);
}
int main()
{
    const int grid_size=2;
    const int block_size=4;
    hello_from_gpu<<<grid_size,block_size>>>();

    cudaDeviceSynchronize();
    return 0;
}

CUDA将线程组织为网格，网格被划分为线程块，即网格由线程块组成，线程块由线程组成，三者之间并不是简单的包含关系（即一般论的包含关系，网格有N个线程块，每个线程块又有M个线程），但是目前示例中确实只是单纯的包含关系，可简单认为有网格两个线程块，每个线程块四个线程共八个线程

下图展示了多维线程的布局，这也是网格的真正分布：

可以认为网格是一个三维立方体，被划分出的小格就是线程块，而小格又可以按三维继续分割，分割出的就是线程。

像上面简单的以2和4指定的线程块个数和线程个数，三维上可以分别认为是(2,1,1)和(4,1,1)，只按x轴延伸的2个大方块，每个大方块又是由按x轴延展的4个小方块组成

多维线程的CUDA程序

既然有了三维建构，每个线程块和线程自然有其三维坐标，按照简单的数学逻辑也可以转换为1维坐标，这里直接把这些概念在代码中展示：

/**
 运行命令: nvcc  XX.cu
*/
#include<cuda.h>
#include<cuda_runtime.h>
#include<cstdio>

__global__ void hello_from_gpu()
{
    const int bid = blockIdx.x;
    const int tid_x = threadIdx.x;
    const int tid_y = threadIdx.y;
    const int tid_z = threadIdx.z;
    //1D坐标:
    const int bid_1D = blockIdx.x+blockIdx.y*gridDim.x+blockIdx.z*gridDim.x*gridDim.y;
    const int tid_1D = threadIdx.x + blockDim.x*threadIdx.y + blockDim.x*blockDim.y*threadIdx.z;
    printf("Device(%d)(%d): Hello World from block %d and thread (%d-%d-%d)\n",bid_1D,tid_1D,bid,tid_x,tid_y,tid_z);
}
int main()
{
    const int grid_size=2;
    const dim3 block_size(3,3,3);
    hello_from_gpu<<<grid_size,block_size>>>();

    cudaDeviceSynchronize();
    return 0;
}

注意到：

• dim3类似于vec3，表3维向量
• gridDim和blockDim是dim3固有变量，内容是当初规定的各个轴的上限
• blockIdx和threadIdx也是dim3固有变量，包含了当前线程的坐标

另外，线程块在3个轴上的个数和一个线程块的线程数是有限制的：

• 线程块在xyz三轴上分别最多：2^31-1,65535,65535
• 一个线程块线程总数最多1024，在三维上最多:1024,1024,64

下面是超过了限制的程序：

/**
 运行命令: nvcc  XX.cu
*/
#include<cuda.h>
#include<cuda_runtime.h>
#include<cstdio>

__global__ void hello_from_gpu()
{
    const int bid = blockIdx.x;
    const int tid_x = threadIdx.x;
    const int tid_y = threadIdx.y;
    const int tid_z = threadIdx.z;
    //1D坐标:
    const int bid_1D = blockIdx.x+blockIdx.y*gridDim.x+blockIdx.z*gridDim.x*gridDim.y;
    const int tid_1D = threadIdx.x + blockDim.x*threadIdx.y + blockDim.x*blockDim.y*threadIdx.z;
    printf("Device(%d)(%d): Hello World from block %d and thread (%d-%d-%d)\n",bid_1D,tid_1D,bid,tid_x,tid_y,tid_z);
}
int main()
{
    const dim3 grid_size(1,65535,1); //大于规定的2^31-1,65535,65535范围虽然能通过编译,但没有输出
    const dim3 block_size(1024,1,1); //大于1024虽然能通过编译,但没有输出
    hello_from_gpu<<<grid_size,block_size>>>();

    cudaDeviceSynchronize();
    return 0;
}

线程束

还有一个没有在代码中体现出的概念：线程束，简单理解为同一线程块中相邻的N个线程，一般按架构为32个一组

夫道未始有封，言未始有常，为是而有畛也。

作为《深度学习入门》的阅读笔记，本文简略概述感知机，并辅以python实现

Python基础

再简单的描述离不开基本的语言基础，确保了解《蜻蜓点水-python》中的内容在进行下去

逻辑运算和感知机

如上所述，python中有逻辑运算，电路中也有逻辑运算，并以“门”的形式呈现，有电流则为1（True），无电流则为0（True）。

如下图所示，三种门的符号，真值表就不再赘述，可以用python快速验证

感知机的历史1957提出，是神经网络的起源的算法之一。

感知机接受多个信号，并输出一个信号。

感知机实现

先说实现的思路，上述三种逻辑运算都可以用同一种逻辑表达，区别只在于特定参数的不同

b称为偏置，w称为权重，二者控制了函数的不同逻辑表示

下面直接给出代码实现：

• 初始化函数设定了权重和偏置
• forward函数在深度学习中代表正向传播，这里借用作为感知机的处理函数
• 逻辑如公式，但将等号给了1

class Perceptron:
    def __init__(self, w=np.array([0,0]),delta=-1.):
        self.w=w
        self.delta = delta
    def forward(self,x):
        tmp=self.w*x
        if np.sum(tmp) + self.delta < 0:
            return 0
        else:
            return 1

andclass = Perceptron(w=np.array([0.5,0.5]))
print(andclass.forward(np.array([0,0])))
print(andclass.forward(np.array([1,0])))
print(andclass.forward(np.array([0,1])))
print(andclass.forward(np.array([1,1])))
# 0
# 0
# 0
# 1

orclass = Perceptron(w=np.array([1,1]))
print(orclass.forward(np.array([0,0])))
print(orclass.forward(np.array([1,0])))
print(orclass.forward(np.array([0,1])))
print(orclass.forward(np.array([1,1])))

# 0
# 1
# 1
# 1

nandclass = Perceptron(w=np.array([-0.5,-0.5]),delta=0.9)
print(nandclass.forward(np.array([0,0])))
print(nandclass.forward(np.array([1,0])))
print(nandclass.forward(np.array([0,1])))
print(nandclass.forward(np.array([1,1])))

# 1
# 1
# 1
# 0

目前感知机的局限

可以对当前感知机的输入和输出作图，可得：

总之，有办法做一条直线将0和1的结果划分开，能这么表示的称为线性感知机

但是这样的感知机无法表示XOR操作，也就是异或操作：

XOR感知机的实现

该问题由1969年被提出，直到1986年被解决，这段时间人工智能的发展是停滞的。

解决的方法就是使用多层感知机，即单独的“门”无法解决的问题，就用多个连接一起解决，直接给方法：

用代码表示就是:

def XOR(x1, x2):
    s1 = NAND(x1, x2)
    s2 = OR(x1, x2)
    y = AND(s1, s2)
    return y

进一步的实现：

#def XOR(x1, x2):
#    s1 = NAND(x1, x2)
#    s2 = OR(x1, x2)
#    y = AND(s1, s2)
#    return y
class XOR:
    def __init__(self):
        self.nandclass = Perceptron(w=np.array([-0.5,-0.5]),delta=0.9)
        self.orclass = Perceptron(w=np.array([1,1]))
        self.andclass = Perceptron(w=np.array([0.5, 0.5]))
    def forward(self,x):
        s1 = self.nandclass.forward(x)
        s2 = self.orclass.forward(x)
        tmp = self.andclass.forward(np.array([s1,s2]))
        return tmp

print("XOR: ")
print(XOR().forward(np.array([0,0])))
print(XOR().forward(np.array([1,0])))
print(XOR().forward(np.array([0,1])))
print(XOR().forward(np.array([1,1])))

#XOR: 
#0
#1
#1
#0

其他问题

多层感知机的层数

按权重算，2层，因为有两套权重参数

按节点（神经元）算，3层，从输入，到中间结果，到输出，有三个阶段

多层感知机的进一步

感知机通过叠加层能够进行非线性的表示，理论上还可以表示计算机进行的处理

多层感知机可以认为是一种神经网络

道在屎溺。作为《深度学习入门》的阅读笔记，本文简略概述python

算数计算

print(1 - 2)
print(4 * 5)
print(7 / 5)
print(3 ** 2) # 乘方

-1
20
1.4
9

数据类型和Type函数

type(10)

int

type(2.718)

float

type("hello")

str

type(True)

bool

变量

x = 100
y = 3.14
print(x * y)
type(x * y)
# 是注释标识

314.0

float

列表

a = [1,2,3,4,5]
print(a)
len(a)

[1, 2, 3, 4, 5]

5

print(a[4])
a[4]=99
print(a)

5
[1, 2, 3, 4, 99]

# 切片
a[0:2] # 获取索引为0到2（不包括2！）的元素

[1, 2]

a[1:]  # 获取从索引为1的元素到最后一个元

[2, 3, 4, 99]

a[:3]  # 获取从第一个元素到索引为3（不包括3！）的元素

[1, 2, 3]

a[:-1] # 获取从第一个元素到最后一个元素的前一个元素之间的元素

[1, 2, 3, 4]

a[:-2] # 获取从第一个元素到最后一个元素的前二个元素之间的元素

[1, 2, 3]

字典

me = {'height':180} # 生成字典
print(me)

{'height': 180}

me['height']=70
print(me['height'])

70

布尔型

hungry = True
sleepy = False
type(hungry)

bool

not hungry

False

hungry and sleepy

False

hungry or sleepy

True

if语句

if hungry:
    print('hungry')
elif sleepy:
    print('sleepy')
else:
    print('not hungry')

hungry

for 语句

for i in [1,2,3]:
    print(i)

1
2
3

函数

def hello(object):
    print("Hello " + object + "!")
hello("cat!")

Hello cat!!

类

class Man:
    def __init__(self, name):
        self.name = name
        print("Initialized!")
    def hello(self):
        print("Hello " + self.name + "!")
    def goodbye(self):
        print("Good-bye " + self.name + "!")

m = Man("David")
m.hello()
m.goodbye()

Initialized!
Hello David!
Good-bye David!

Numpy 速成

import numpy as np

x = np.array([1.,2.,3.])
print(x)
type(x)
y = np.array([2.0, 4.0, 6.0])

[1. 2. 3.]

以下是对应元素的运算,如果元素个数不同，程序就会报错，所以元素个数保持一致非常重要

print(x+y)
print(x-y)
print(x/y)
print(x*y)

[3. 6. 9.]
[-1. -2. -3.]
[0.5 0.5 0.5]
[ 2.  8. 18.]

NumPy数组不仅可以进行element-wise运算，也可以和单一的数值（标量）组合起来进行运算

x/2.0

array([0.5, 1. , 1.5])

Numpy的N维数组

A = np.array([[1, 2], [3, 4]])
B = np.array([[3, 0],[0, 6]])
print(A)

[[1 2]
 [3 4]]

A.shape

(2, 2)

A.dtype

dtype('int32')

print(A+B)

[[ 4  2]
 [ 3 10]]

print(A*B)

[[ 3  0]
 [ 0 24]]

print(A*10)

[[10 20]
 [30 40]]

广播

A = np.array([[1, 2], [3, 4]])
B = np.array([10, 20])
A * B

array([[10, 40],
       [30, 80]])

访问元素

X = np.array([[51, 55], [14, 19], [0, 4]])
X[0][1]

55

X[0]

array([51, 55])

for row in X:
    print(row)

[51 55]
[14 19]
[0 4]

X = X.flatten()
print(X)

[51 55 14 19  0  4]

还有一些特殊的访问方法:

X[np.array([0, 2, 4])]

array([51, 14,  0])

X > 15

array([ True,  True, False,  True, False, False])

X[X>15]

array([51, 55, 19])

Matplolib 速成

import matplotlib.pyplot as plt
 # 生成数据
x = np.arange(0, 6, 0.1) # 以0.1为单位，生成0到6的数据
y1 = np.sin(x)
y2 = np.cos(x)
# 绘制图形
plt.plot(x, y1)
plt.show()

​
​

plt.plot(x, y1, label="sin")
plt.plot(x, y2, linestyle = "--", label="cos") # 用虚线绘制
plt.xlabel("x") # x轴标签
plt.ylabel("y") # y轴标签
plt.title('sin & cos') # 标题
plt.legend() #图例
plt.show()

​

​

from matplotlib.image import imread
img = imread('../dataset/lena.png') # 读入图像（设定合适的路径！）
plt.imshow(img)
plt.show()

​

​

原因

在调整spark配置问题时把集群内部通信接口暴露给了公网，提供了攻击的机会

攻击者通过端口扫描的方式，用暴露的接口【通过某种方式】进入集群，集群内部可无密码ssh访问，直接访问ssh目录生成密钥，并修改authorized_keys权限

添加redis和程序包（注入+进程隐藏）

• https://github.com/arget13/memdlopen/blob/main/README.md
• https://github.com/gianlucaborello/libprocesshider/blob/master/README.md

这导致了top无法发现其真正进程

发现

cpu占用异常，单核100%，和top无法对应

netstat -alntp 发现异常tcp链接：5.180.174.211:80，检查为境外乌克兰ip，lastb发现有大量异常失败登录，且ip大多来自境外，疑似撞库

处理

关掉集群其他机器

立即修改安全组，关闭除22以外的公网端口

试图删除authorized_keys known_host中的攻击者数据，发现并无权限，但查看权限并无异常，使用lsattr查看文件属性发现被添加ia标记，无法修改无法删除，使用chattr 修改无效且返回结果看起来比较诡异，查看/usr/bin/chattr 发现已被修改，使用ls -l |grep Jan 过滤出最近修改的所有文件，发现其他新增命令工具

但此时需要chattr 来修改标识才能进行操作，从https://github.com/posborne/linux-programming-interface-exercises/blob/master/15-file-attributes/chattr.c下载文件gcc编译，重命名为chattr，利用其来修改标识 chattr -ia，删除原本的chattr 替换之，然后就可以对其他文件进行操作

https://c.biancheng.net/view/874.html

查看/etc/ld.so.preload文件，此文件原先应为空或不存在，设法清空或删除该文件（和文件里的文件路径），top的问题就可以恢复

参考：https://www.jianshu.com/p/d4b3d87fdc51

也可通过busybox top工具直接发现真正的进程，定位到程序所在的目录（我的是/var/tmp/.11/，里面都是编译后的二进制文件和上述github的bash脚本），先不急删

使用systemctl status 对应pid，发现有对应的service（于我情况如此，也可能会使用定时程序），清理掉就可以，然后附带检查service而文件的修改日期，发现有其他攻击者新增但未启动的service一并处理

另外本环境没有redis服务，故连带清理掉，这样cpu和内存就应该能降下来

最后删除病毒文件即可

然后关闭该机器，对其他机器依次操作（因为集群是免密互通的，保险起见）

后续改进

关闭所有不必要的公网端口

内部web服务使用nginx反代并管理登录

Nginx 配置 安全认证 反向代理 HDFS web 页面_nginx如何代理一个web页面-CSDN博客