Reverse-CUDA学习笔记1

博主： Xherlock
发布时间：2025 年 06 月 04 日
169 次浏览
暂无评论
2885字数
分类： Reverse

# 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程序要做的：

1. CPU申请GPU内存，把要计算的内容从系统内存拷贝到GPU内存
2. GPU核函数进行计算
3. 从GPU内存拷贝到系统内存，并释放GPU显存和内存

实验室windows本机没gpu，只好现在远程服务器linux上配置了一个cuda环境，配置比较顺利，先前跑深度学习已经装了好几版cuda toolkit

配置参考教程：https://zhuanlan.zhihu.com/p/79059379

跑了一个示例代码，作用是查找设备和获取设备信息

~~~c++
#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如下

~~~c++
#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如下

~~~c++
#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代码，无法直接查看

~~~c++
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 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学习下如何编写些实际应用，比如文件加密、图像处理之类的

## 参考文章

* https://lazypool-blog.netlify.app/2025/02/23/cuda-programming/
* https://zhuanlan.zhihu.com/p/442052342
* https://zhuanlan.zhihu.com/p/23762077736

最后修改：2025 年 06 月 04 日

如果觉得我的文章对你有用，请随意赞赏

发表评论取消回复
使用cookie技术保留您的个人信息以便您下次快速评论，继续评论表示您已同意该条款

评论 *

私密评论

名称 *

🎲

邮箱 *

地址

Reverse（十一）
浏览次数: 11610
Reverse（十三）
浏览次数: 4255
Win11新机鼓捣记录
浏览次数: 2835
js逆向爬虫1
浏览次数: 2203
欢迎使用 Typecho
浏览次数: 2072

mars
博主你好，请问怎么知道的alloc函数属于Buffer下的方法...
lj
该评论仅登录用户及评论双方可见
Mr.Guan
:@(赞一个)
Typecho
欢迎加入 Typecho 大家族

Reverse-CUDA学习笔记1

Xherlock • 2025 年 06 月 04 日

# Reverse-CUDA学习笔记1

## CPU vs GPU

> 学cuda肯定得先了解为什么用GPU，和CPU的区别在哪里

对于处理器（CPU/GPU）有两个**指标**要考虑：

* 延迟：发出指令到收到结果的时间间隔
* 吞吐量：单位时间之间处理的指令数

CPU特点：

GPU特点：

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

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

单个完整的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

跑了一个示例代码，作用是查找设备和获取设备信息

~~~c++
#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. 获取当前的设备编号

编译指令为`nvcc -std=c++11 -o getdevice getDevice.cpp`，输出如下

~~~
Found 2 CUDA devices
Device Name: NVIDIA GeForce RTX 3090
Compute Capability: 8.6
~~~

找了个涉及计算的代码

CPU如下

~~~c++
#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如下

~~~c++
#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);

cudaFree(da);
    cudaFree(db);
    cudaFree(dc);

free(a);
    free(b);
    free(c);
    return 0;
}
~~~

在循环次数少的情况下GPU不如CPU，正好和前面讲GPU优先吞吐量，延迟较长，循环少的情况下优势反倒无法体现出来

在循环次数多的情况下GPU速度快很多

ida反编译GPU编译文件可以看到已经把vecAddKernel编译成了PTX代码，无法直接查看

cuobjdump -ptx ./VectorSumGPU 可以获取ptx代码

~~~

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

$L__BB0_2:
ret;

}
~~~

这里就是正常的逆向分析了，目前还没看到有什么辅助分析的工具，下次学习笔记里学习下

## 小结

初步学习了下CUDA相关知识概念，理解了GPU和CPU区别，然后配置了下Linux CUDA环境，目前能够编译cuda程序，此外学习了如何使用cuobjdump

下次试着用相关api学习下如何编写些实际应用，比如文件加密、图像处理之类的

## 参考文章

* https://lazypool-blog.netlify.app/2025/02/23/cuda-programming/
* https://zhuanlan.zhihu.com/p/442052342
* https://zhuanlan.zhihu.com/p/23762077736

Reverse-CUDA学习笔记1

发表评论取消回复
使用cookie技术保留您的个人信息以便您下次快速评论，继续评论表示您已同意该条款

Reverse（十一）

Reverse（十三）

Win11新机鼓捣记录

js逆向爬虫1

欢迎使用 Typecho

数据结构与算法~图

信息安全概论

数据库学习（一）

美化和集成控制台

计算机网络和因特网（一）

Reverse-CUDA学习笔记1

发表评论 取消回复 使用cookie技术保留您的个人信息以便您下次快速评论，继续评论表示您已同意该条款

Reverse-CUDA学习笔记1

发表评论取消回复
使用cookie技术保留您的个人信息以便您下次快速评论，继续评论表示您已同意该条款