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的概念模型

线程->线程块->网格,如下图左半部分

Kernel 批处理(左)与 CUDA 内存模型(右)

threadId.x = blockidx.x*blockDim.x+threadidx.x

threadId.y = blockidx.x*blockDim.y+threadidx.y

核函数:由主机调用,在设备上运行的函数,指示了网格内所有线程的行为。 单个 kernel 可能由多个线程块执行,线程块内的线程将 共享某块内存,并在必要时同步

流式多处理器:并行计算核心

单个完整的SM包括:多个SP、指令缓存、线程束调度器、分派单元、寄存器文件、加载/存储单元队列、特殊功能单元队列、共享内存/L1缓存、统一缓存等

线程束:连接软件层和硬件层的纽带

CUDA编程

一个CUDA程序要做的:

  1. CPU申请GPU内存,把要计算的内容从系统内存拷贝到GPU内存
  2. GPU核函数进行计算
  3. 从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学习下如何编写些实际应用,比如文件加密、图像处理之类的

参考文章

最后修改:2025 年 06 月 04 日
如果觉得我的文章对你有用,请随意赞赏