Reverse-CUDA学习笔记1
CPU vs GPU
学cuda肯定得先了解为什么用GPU，和CPU的区别在哪里
对于处理器（CPU/GPU）有两个指标要考虑：
- 延迟：发出指令到收到结果的时间间隔
- 吞吐量：单位时间之间处理的指令数
CPU特点：
- 多级高速缓存结构：处理运算速度远大于访问存储的速度，遵循空间换时间思想，设计了该结构，把经常访问的内容放到低级缓存（访问更快），不经常访问的放到高级缓存中，提高整体访问存储速度；L1>L2>L3访问速度
- 很多控制单元（PU）：分支预测机制、流水线前传机制
- 运算单元强大：整型浮点型复杂运算速度快
GPU特点：
- 虽有缓存结构但数量少：GPU减少指令访问缓存次数
- 控制单元非常简单：没有CPU的两种机制，复杂指令运算较慢
- 运算单元非常多：采用长延时流水线以实现高吞吐量，每一行运算单元控制器只有一个，即每一行运算单元使用的指令是相同的，不同的是数据内容；这种方式提高了效率
GPU设计原则：增加简单指令的吞吐（吞吐导向设计）
总结下二者的对比：
- CPU：连续计算、延迟优先、CPU比GPU单条复杂指令延迟快10倍以上
- GPU：并行计算、吞吐优先、GPU比CPU单位时间执行指令数量大10倍以上
Q：所以什么地方适合用GPU
A：计算密集、数据并行的时候，计算不复杂，可以拆分指令
CUDA概念扫盲
啥是CUDA
CUDA——Compute Unified Device Architecture 统一计算架构，英伟达NVIDIA 2007年推出的一种软硬件集成技术
CUDA可以干啥
通过CUDA，使用者可以利用NVIDIA的GPU进行图像处理之外的运算，同时可以利用GPU作为C编译器的开发环境
CUDA Toolkit：将自家的CUDA C语言执行于GPU的部分编译成PTX中间语言或是特定NVIDIA GPU架构的机器代码；执行于CPU部分的C/C++代码仍依赖外部编译器
CUDA由什么组成，怎么结合软硬件
CPU-主机、GPU-设备，各自独立内存；一般来说，在CPU上运行串行工作负载，将并行计算卸载到GPU上
根据上面的概念，CUDA将其架构分为软件层和硬件层：
- 软件层：线程、线程块、网格
- 硬件层：CUDA核心（流处理器，SP）、流式多处理器（SM，streaming multiprocessor）、GPU
- 线程束：连接硬件层和软件层
线程和CUDA核心：最小的基本运算单位（单指令多线程）
线程：在CUDA中，线程按束调度：同束线程共享相同的程序计数器，同步地执行相同的指令，作用域各自寄存器存储的数据
CUDA核心：CUDA中执行标量运算指令的基本单元，核心组件是整数运算单元和浮点数运算单元；同样也是按束调度
- 二者联系：单个CUDA核心执行来自单个线程的指令
线程块、网格、kernel：CUDA的概念模型
线程->线程块->网格，如下图左半部分
threadId.x = blockidx.x*blockDim.x+threadidx.x
threadId.y = blockidx.x*blockDim.y+threadidx.y
核函数：由主机调用，在设备上运行的函数，指示了网格内所有线程的行为。 单个 kernel 可能由多个线程块执行，线程块内的线程将 共享某块内存，并在必要时同步。
流式多处理器：并行计算核心
单个完整的SM包括：多个SP、指令缓存、线程束调度器、分派单元、寄存器文件、加载/存储单元队列、特殊功能单元队列、共享内存/L1缓存、统一缓存等
线程束：连接软件层和硬件层的纽带
CUDA编程
一个CUDA程序要做的：
- CPU申请GPU内存，把要计算的内容从系统内存拷贝到GPU内存
- GPU核函数进行计算
- 从GPU内存拷贝到系统内存，并释放GPU显存和内存
实验室windows本机没gpu，只好现在远程服务器linux上配置了一个cuda环境，配置比较顺利，先前跑深度学习已经装了好几版cuda toolkit
配置参考教程：https://zhuanlan.zhihu.com/p/79059379
跑了一个示例代码，作用是查找设备和获取设备信息
#include <stdio.h>
#include <cuda_runtime.h>
int main() {
int device_count;
cudaGetDeviceCount(&device_count); // 1. 获取设备数量
printf("Found %d CUDA devices\n", device_count);
cudaSetDevice(0); // 2. 选择 0 号设备
int dev;
cudaGetDevice(&dev); // 3. 获取当前的设备编号
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0); // 4. 获取设备属性
printf("Device Name: %s\n", prop.name);
printf("Compute Capability: %d.%d\n", prop.major, prop.minor);
return 0;
}编译指令为nvcc -std=c++11 -o getdevice getDevice.cpp，输出如下
Found 2 CUDA devices
Device Name: NVIDIA GeForce RTX 3090
Compute Capability: 8.6找了个涉及计算的代码
CPU如下
#include <iostream>
#include <cstdlib>
#include <sys/time.h>
using namespace std;
void vecAdd(float* A, float* B, float* C, int n) {
for (int i = 0; i < n; i++) {
C[i] = A[i] + B[i];
}
}
int main(int argc, char *argv[]) {
int n = atoi(argv[1]);
cout << n << endl;
size_t size = n * sizeof(float);
// host memery
float *a = (float *)malloc(size);
float *b = (float *)malloc(size);
float *c = (float *)malloc(size);
for (int i = 0; i < n; i++) {
float af = rand() / double(RAND_MAX);
float bf = rand() / double(RAND_MAX);
a[i] = af;
b[i] = bf;
}
struct timeval t1, t2;
gettimeofday(&t1, NULL);
vecAdd(a, b, c, n);
gettimeofday(&t2, NULL);
double timeuse = (t2.tv_sec - t1.tv_sec) + (double)(t2.tv_usec - t1.tv_usec)/1000000.0;
cout << timeuse << endl;
free(a);
free(b);
free(c);
return 0;
}GPU如下
#include <iostream>
#include <stdio.h>
#include <cstdlib>
#include <sys/time.h>
#include <cuda_runtime.h>
using namespace std;
__global__
void vecAddKernel(float* A_d, float* B_d, float* C_d, int n)
{
int i = threadIdx.x + blockDim.x * blockIdx.x;
if (i < n) C_d[i] = A_d[i] + B_d[i];
}
int main(int argc, char *argv[]) {
int n = atoi(argv[1]);
cout << n << endl;
size_t size = n * sizeof(float);
// host memery
float *a = (float *)malloc(size);
float *b = (float *)malloc(size);
float *c = (float *)malloc(size);
for (int i = 0; i < n; i++) {
float af = rand() / double(RAND_MAX);
float bf = rand() / double(RAND_MAX);
a[i] = af;
b[i] = bf;
}
float *da = NULL;
float *db = NULL;
float *dc = NULL;
cudaMalloc((void **)&da, size);
cudaMalloc((void **)&db, size);
cudaMalloc((void **)&dc, size);
cudaMemcpy(da,a,size,cudaMemcpyHostToDevice);
cudaMemcpy(db,b,size,cudaMemcpyHostToDevice);
cudaMemcpy(dc,c,size,cudaMemcpyHostToDevice);
struct timeval t1, t2;
int threadPerBlock = 256;
int blockPerGrid = (n + threadPerBlock - 1)/threadPerBlock;
printf("threadPerBlock: %d \nblockPerGrid: %d \n",threadPerBlock,blockPerGrid);
gettimeofday(&t1, NULL);
vecAddKernel <<< blockPerGrid, threadPerBlock >>> (da, db, dc, n);
gettimeofday(&t2, NULL);
cudaMemcpy(c,dc,size,cudaMemcpyDeviceToHost);
//for (int i = 0; i < 10; i++)
// cout << vecA[i] << " " << vecB[i] << " " << vecC[i] << endl;
double timeuse = (t2.tv_sec - t1.tv_sec) + (double)(t2.tv_usec - t1.tv_usec)/1000000.0;
cout << timeuse << endl;
cudaFree(da);
cudaFree(db);
cudaFree(dc);
free(a);
free(b);
free(c);
return 0;
}在循环次数少的情况下GPU不如CPU，正好和前面讲GPU优先吞吐量，延迟较长，循环少的情况下优势反倒无法体现出来
在循环次数多的情况下GPU速度快很多
ida反编译GPU编译文件可以看到已经把vecAddKernel编译成了PTX代码，无法直接查看
int __fastcall main(int argc, const char **argv, const char **envp)
{
__m128d v3; // xmm0
__int64 v4; // rax
__m128d v5; // xmm0
__int64 v6; // rax
struct timeval v8; // [rsp+20h] [rbp-A0h] BYREF
timeval tv; // [rsp+30h] [rbp-90h] BYREF
float *v10; // [rsp+48h] [rbp-78h] BYREF
float *v11; // [rsp+50h] [rbp-70h] BYREF
float *v12; // [rsp+58h] [rbp-68h] BYREF
__int64 v13; // [rsp+60h] [rbp-60h] BYREF
unsigned int v14; // [rsp+68h] [rbp-58h]
__int64 v15; // [rsp+70h] [rbp-50h] BYREF
unsigned int v16; // [rsp+78h] [rbp-48h]
double v17; // [rsp+80h] [rbp-40h]
unsigned int v18; // [rsp+88h] [rbp-38h]
unsigned int v19; // [rsp+8Ch] [rbp-34h]
float v20; // [rsp+90h] [rbp-30h]
float v21; // [rsp+94h] [rbp-2Ch]
void *v22; // [rsp+98h] [rbp-28h]
void *v23; // [rsp+A0h] [rbp-20h]
void *ptr; // [rsp+A8h] [rbp-18h]
size_t size; // [rsp+B0h] [rbp-10h]
int v26; // [rsp+B8h] [rbp-8h]
int i; // [rsp+BCh] [rbp-4h]
v26 = atoi(argv[1]);
v4 = std::ostream::operator<<(&std::cout, (unsigned int)v26);
std::ostream::operator<<(v4, &std::endl<char,std::char_traits<char>>);
size = 4LL * v26;
ptr = malloc(size);
v23 = malloc(size);
v22 = malloc(size);
for ( i = 0; i < v26; ++i )
{
v3.m128d_f64[0] = (double)rand() / 2147483647.0;
v5 = _mm_unpacklo_pd(v3, v3);
v21 = v5.m128d_f64[0];
v5.m128d_f64[0] = (double)rand() / 2147483647.0;
v3 = _mm_unpacklo_pd(v5, v5);
v20 = v3.m128d_f64[0];
*((float *)ptr + i) = v21;
*((float *)v23 + i) = v20;
}
v12 = 0LL;
v11 = 0LL;
v10 = 0LL;
cudaMalloc(&v12, size);
cudaMalloc(&v11, size);
cudaMalloc(&v10, size);
cudaMemcpy(v12, ptr, size, 1LL);
cudaMemcpy(v11, v23, size, 1LL);
cudaMemcpy(v10, v22, size, 1LL);
v19 = 256;
v18 = (v26 + 255) / 256;
printf("threadPerBlock: %d \nblockPerGrid: %d \n", 256, v18);
gettimeofday(&tv, 0LL);
dim3::dim3((dim3 *)&v13, v19, 1u, 1u);
dim3::dim3((dim3 *)&v15, v18, 1u, 1u);
if ( !(unsigned int)_cudaPushCallConfiguration(v15, v16, v13, v14, 0LL, 0LL) )
vecAddKernel(v12, v11, v10, v26);
gettimeofday(&v8, 0LL);
cudaMemcpy(v22, v10, size, 2LL);
v17 = (double)(LODWORD(v8.tv_usec) - LODWORD(tv.tv_usec)) / 1000000.0
+ (double)(LODWORD(v8.tv_sec) - LODWORD(tv.tv_sec));
v6 = std::ostream::operator<<(&std::cout, v17);
std::ostream::operator<<(v6, &std::endl<char,std::char_traits<char>>);
cudaFree(v12);
cudaFree(v11);
cudaFree(v10);
free(ptr);
free(v23);
free(v22);
return 0;
}
__int64 __fastcall vecAddKernel(float *a1, float *a2, float *a3, int a4)
{
return __device_stub__Z12vecAddKernelPfS_S_i(a1, a2, a3, a4);
}
char __fastcall __device_stub__Z12vecAddKernelPfS_S_i(float *a1, float *a2, float *a3, int a4)
{
char result; // al
int v5; // [rsp+14h] [rbp-6Ch] BYREF
float *v6; // [rsp+18h] [rbp-68h] BYREF
float *v7; // [rsp+20h] [rbp-60h] BYREF
float *v8; // [rsp+28h] [rbp-58h] BYREF
_QWORD v9[4]; // [rsp+30h] [rbp-50h] BYREF
__int64 v10; // [rsp+50h] [rbp-30h] BYREF
__int64 v11; // [rsp+58h] [rbp-28h] BYREF
__int64 v12; // [rsp+60h] [rbp-20h] BYREF
int v13; // [rsp+68h] [rbp-18h]
__int64 v14; // [rsp+70h] [rbp-10h] BYREF
int v15; // [rsp+78h] [rbp-8h]
int v16; // [rsp+7Ch] [rbp-4h]
v8 = a1;
v7 = a2;
v6 = a3;
v5 = a4;
v9[0] = &v8;
v9[1] = &v7;
v9[2] = &v6;
v9[3] = &v5;
v16 = 4;
__device_stub__Z12vecAddKernelPfS_S_i(float *,float *,float *,int)::__f = (__int64)vecAddKernel;
v14 = 0x100000001LL;
v15 = 1;
v12 = 0x100000001LL;
v13 = 1;
result = (unsigned int)_cudaPopCallConfiguration(&v14, &v12, &v11, &v10) != 0;
if ( !result )
return cudaLaunchKernel<char>((unsigned int)vecAddKernel, v14, v15, v12, v13, (unsigned int)v9, v11, v10);
return result;
}cuobjdump -ptx ./VectorSumGPU 可以获取ptx代码
Fatbin elf code:
================
arch = sm_52
code version = [1,7]
producer = <unknown>
host = linux
compile_size = 64bit
Fatbin elf code:
================
arch = sm_52
code version = [1,7]
producer = <unknown>
host = linux
compile_size = 64bit
Fatbin ptx code:
================
arch = sm_52
code version = [7,3]
producer = <unknown>
host = linux
compile_size = 64bit
compressed
.version 7.3
.target sm_52
.address_size 64
.visible .entry _Z12vecAddKernelPfS_S_i(
.param .u64 _Z12vecAddKernelPfS_S_i_param_0,
.param .u64 _Z12vecAddKernelPfS_S_i_param_1,
.param .u64 _Z12vecAddKernelPfS_S_i_param_2,
.param .u32 _Z12vecAddKernelPfS_S_i_param_3
)
{
.reg .pred %p<2>;
.reg .f32 %f<4>;
.reg .b32 %r<6>;
.reg .b64 %rd<11>;
ld.param.u64 %rd1, [_Z12vecAddKernelPfS_S_i_param_0];
ld.param.u64 %rd2, [_Z12vecAddKernelPfS_S_i_param_1];
ld.param.u64 %rd3, [_Z12vecAddKernelPfS_S_i_param_2];
ld.param.u32 %r2, [_Z12vecAddKernelPfS_S_i_param_3];
mov.u32 %r3, %tid.x;
mov.u32 %r4, %ctaid.x;
mov.u32 %r5, %ntid.x;
mad.lo.s32 %r1, %r5, %r4, %r3;
setp.ge.s32 %p1, %r1, %r2;
@%p1 bra $L__BB0_2;
cvta.to.global.u64 %rd4, %rd1;
mul.wide.s32 %rd5, %r1, 4;
add.s64 %rd6, %rd4, %rd5;
cvta.to.global.u64 %rd7, %rd2;
add.s64 %rd8, %rd7, %rd5;
ld.global.f32 %f1, [%rd8];
ld.global.f32 %f2, [%rd6];
add.f32 %f3, %f2, %f1;
cvta.to.global.u64 %rd9, %rd3;
add.s64 %rd10, %rd9, %rd5;
st.global.f32 [%rd10], %f3;
$L__BB0_2:
ret;
}这里就是正常的逆向分析了，目前还没看到有什么辅助分析的工具，下次学习笔记里学习下
小结
初步学习了下CUDA相关知识概念，理解了GPU和CPU区别，然后配置了下Linux CUDA环境，目前能够编译cuda程序，此外学习了如何使用cuobjdump
下次试着用相关api学习下如何编写些实际应用，比如文件加密、图像处理之类的
