第4章因子数据并行化

第4章：因子数据并行化

好，咱们进入正题。因子挖掘上了GPU，第一件事就是——数据怎么摆？

很多人觉得，把数据从CPU搬到GPU不就完事了吗？嗯，没那么简单。数据在GPU上的组织方式，直接决定了你的计算效率。我见过不少项目，代码写得挺漂亮，一跑起来慢得像蜗牛，问题就出在数据布局上。

4.1 因子数据在GPU上的表示

先问个问题：你有1000只股票，每只股票有500个时间点的因子数据，你怎么存？

两种主流方式：结构体数组（AoS）和数组结构体（SoA）。名字听着绕，其实说白了就是——

AoS：先存第一只股票的所有因子，再存第二只的。像一本本独立的档案。
SoA：把所有股票的因子1放一起，因子2放一起。像按科目分类的成绩单。

我个人习惯在因子挖掘中用SoA。为什么？因为GPU的并行本质是SIMT（单指令多线程）。你想想看，如果我要计算所有股票的某个因子，用SoA的话，相邻线程访问的是连续内存地址——这叫合并访问，带宽利用率直接拉满。

核心原则：GPU喜欢连续访问，讨厌随机跳转。SoA天然适配这个特性。

举个例子。假设你有N只股票，每只股票有M个因子值：

// AoS方式（不推荐）
struct StockData {
    float factor1;
    float factor2;
    float factor3;
};
StockData stocks[N];

// SoA方式（推荐）
struct StockDataSoA {
    float* factor1;  // 长度为N的数组
    float* factor2;
    float* factor3;
};

我在项目中遇到过一件事。有个同事用AoS存了5000只股票的日频因子，跑一个简单的均值计算，结果比CPU还慢。一分析，发现全局内存访问效率只有12%。改成SoA后，直接飙到85%以上。差距就这么大。

小技巧：如果你实在要用AoS（比如从某个库继承的数据结构），可以用cudaMemcpy2D做转置，或者用共享内存做重排。但说实话，不如一开始就设计成SoA。

4.2 数据从CPU到GPU的传输

数据布局搞定了，接下来就是搬数据。最基础的API就是cudaMemcpy。

cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, 
                       enum cudaMemcpyKind kind);

四个参数，简单粗暴。但有几个坑，我踩过，你们别踩。

坑一：传输方向别搞反。kind参数有四种：

枚举值	含义
cudaMemcpyHostToDevice	CPU → GPU
cudaMemcpyDeviceToHost	GPU → CPU
cudaMemcpyDeviceToDevice	GPU内部拷贝
cudaMemcpyHostToHost	CPU内部拷贝（基本不用）

坑二：别用默认流做大量传输。默认流是同步的，也就是说，cudaMemcpy没完成，CPU就卡在那等着。你想想看，几百兆的数据传输，CPU干等着，多浪费。

警告：我曾经在一个回测系统里，用默认流传输了2GB的因子数据，结果每次启动要等3秒。后来改成异步流，时间降到0.4秒。记住——传输和计算要重叠。

4.3 异步传输与流（Stream）

这才是今天的重头戏。流（Stream）是CUDA里实现并行的核心机制。你可以把流想象成一条流水线——

一个流里，操作是顺序执行的
不同流之间，操作可以并行

异步传输的套路是这样的：

cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

// 异步传输，不阻塞CPU
cudaMemcpyAsync(d_data1, h_data1, size, cudaMemcpyHostToDevice, stream1);
cudaMemcpyAsync(d_data2, h_data2, size, cudaMemcpyHostToDevice, stream2);

// 在stream1上启动kernel
kernel1<<<grid, block, 0, stream1>>>(d_data1);
// 在stream2上启动kernel
kernel2<<<grid, block, 0, stream2>>>(d_data2);

cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);

cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);

这里的关键是cudaMemcpyAsync。它不会等传输完成就返回，CPU可以继续干别的事。比如准备下一批数据，或者启动其他计算。

为什么会这样？因为GPU有专门的DMA引擎，传输数据不需要占用计算单元。你想想看，传输和计算可以同时进行，这不就是白嫖性能吗？

性能调优黄金法则：用多个流把数据传输和kernel计算重叠起来。理想情况下，传输时间被完全隐藏，你只看到计算时间。

我建议的做法是——把因子数据分成多个块。比如1000只股票，分成4组，每组250只。用4个流分别处理：

流1传输第1组，同时CPU准备第2组
流1传输完，立即启动kernel计算第1组
流2传输第2组，与流1的kernel计算并行
以此类推，形成流水线

这种模式叫流水线并行。我在做高频因子计算时，用4个流把吞吐量提升了2.3倍。注意，不是所有场景都能线性提升，因为PCIe带宽是共享的。一般来说，2-4个流效果最好，再多了反而会因为竞争带宽而收益递减。

避坑指南：我曾经用16个流想榨干性能，结果发现传输时间反而变长了。原因是太多流争抢PCIe带宽，每个流都在等。后来改成4个流，配合cudaStreamCreateWithPriority设置优先级，效果最好。

最后说一个容易被忽略的点——固定内存（Pinned Memory）。用cudaHostAlloc分配的内存，传输速度比普通malloc快很多。因为GPU可以直接通过DMA访问固定内存，不需要CPU介入。

float* h_data;
cudaHostAlloc(&h_data, size, cudaHostAllocDefault);
// 用这个指针做cudaMemcpyAsync，速度能快30%-50%

但注意，固定内存是稀缺资源，别滥用。我一般只给传输缓冲区用固定内存，其他数据还是用普通malloc。

嗯，这一章的内容就这些。数据布局选SoA，传输用异步，流要适量，内存用固定。把这几点做到位，你的因子挖掘GPU化就成功了一半。

📚 课程目录