高级内存优化

note

本篇文档的数字仅做示例参考。具体数据，请以实际情况为准。

本文档帮你深入理解 MUSA 高级内存特性：系统内存优化、L2 缓存管理、Cluster 内存和异步执行。

页锁定内存

分配和释放

#include <musa.h>
#include <stdio.h>

int main() {
    size_t size = 1000000 * sizeof(float);
    
    // 分配页锁定内存
    float* h_pinned;
    musaMallocHost(&h_pinned, size);
    
    // 初始化数据
    for (int i = 0; i < 1000000; i++) {
        h_pinned[i] = float(i);
    }
    
    // 使用完成后释放
    musaFreeHost(h_pinned);
    
    return 0;
}

异步内存传输

void asyncTransfer(float* h_pinned, float* d_data, size_t size) {
    musaStream_t stream;
    musaStreamCreate(&stream);
    
    // 异步 H2D 拷贝
    musaMemcpyAsync(d_data, h_pinned, size, 
                    musaMemcpyHostToDevice, stream);
    
    // 启动 kernel（与传输重叠）
    processKernel<<<grid, block, 0, stream>>>(d_data);
    
    // 异步 D2H 拷贝
    musaMemcpyAsync(h_pinned, d_data, size,
                    musaMemcpyDeviceToHost, stream);
    
    // 等待完成
    musaStreamSynchronize(stream);
    musaStreamDestroy(stream);
}

流水线传输

void pipelineTransfer(float* h_input, float* h_output, int n) {
    float *d_input, *d_output;
    musaMalloc(&d_input, n * sizeof(float));
    musaMalloc(&d_output, n * sizeof(float));
    
    // 分配 pinned buffers
    float *h_pinned_in, *h_pinned_out;
    musaMallocHost(&h_pinned_in, n * sizeof(float));
    musaMallocHost(&h_pinned_out, n * sizeof(float));
    
    musaStream_t stream;
    musaStreamCreate(&stream);
    
    // 拷贝输入数据到 pinned buffer
    memcpy(h_pinned_in, h_input, n * sizeof(float));
    
    // 异步传输 + 计算流水线
    musaMemcpyAsync(d_input, h_pinned_in, n, 
                    musaMemcpyHostToDevice, stream);
    processKernel<<<grid, block, 0, stream>>>(d_input, d_output);
    musaMemcpyAsync(h_pinned_out, d_output, n,
                    musaMemcpyDeviceToHost, stream);
    
    // 等待完成
    musaStreamSynchronize(stream);
    
    // 拷贝结果
    memcpy(h_output, h_pinned_out, n * sizeof(float));
    
    // 清理
    musaFreeHost(h_pinned_in);
    musaFreeHost(h_pinned_out);
    musaFree(d_input);
    musaFree(d_output);
    musaStreamDestroy(stream);
}

可移植页锁定内存（零拷贝，Zero-Copy）

// 分配可移植 pinned memory（支持多 GPU）
float* h_mapped;
musaMallocHost(&h_mapped, size, musaHostAllocMapped);

// 获取设备指针
float* d_mapped;
musaHostGetDevicePointer(&d_mapped, h_mapped, 0);

// 直接在 kernel 中使用（零拷贝）
__global__ void zeroCopyKernel(float* data) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < 1000000) {
        data[idx] = data[idx] * 2.0f;
    }
}

zeroCopyKernel<<<grid, block>>>(d_mapped);
musaDeviceSynchronize();

musaFreeHost(h_mapped);

---

零拷贝内存

零拷贝基础

零拷贝允许 GPU 直接访问主机内存，无需显式拷贝：

#include <musa.h>
#include <stdio.h>

#define N 1000000

__global__ void zeroCopyKernel(float* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        data[idx] = data[idx] * 2.0f + 1.0f;
    }
}

int main() {
    size_t size = N * sizeof(float);
    
    // 分配 mapped pinned memory
    float* h_data;
    musaMallocHost(&h_data, size, musaHostAllocMapped);
    
    // 获取设备指针
    float* d_data;
    musaHostGetDevicePointer(&d_data, h_data, 0);
    
    // CPU 初始化
    for (int i = 0; i < N; i++) {
        h_data[i] = float(i);
    }
    
    // GPU 处理（直接访问主机内存）
    dim3 blockSize(256);
    dim3 gridSize((N + blockSize.x - 1) / blockSize.x);
    zeroCopyKernel<<<gridSize, blockSize>>>(d_data, N);
    musaDeviceSynchronize();
    
    // CPU 读取结果（无需拷贝）
    printf("Result[0] = %f\n", h_data[0]);
    
    musaFreeHost(h_data);
    return 0;
}

零拷贝性能考虑

场景	推荐方式	原因
大数据集	Pinned Memory + 异步拷贝	PCIe 带宽最大化

---

写合并内存（Write-Combined Memory）

分配 Write-Combined Memory

// 分配 write-combined memory（优化写入带宽）
float* h_wc;
musaMallocHost(&h_wc, size, musaHostAllocWriteCombined);

// 写入数据（高带宽）
for (int i = 0; i < N; i++) {
    h_wc[i] = float(i);
}

// 拷贝到设备
musaMemcpy(d_data, h_wc, size, musaMemcpyHostToDevice);

musaFreeHost(h_wc);

性能对比

void compareMemoryTypes(int N) {
    size_t size = N * sizeof(float);
    float *d_data, *d_result;
    musaMalloc(&d_data, size);
    musaMalloc(&d_result, size);
    
    // 方式 1: Pageable Memory
    float* h_pageable = (float*)malloc(size);
    for (int i = 0; i < N; i++) h_pageable[i] = float(i);
    
    musaMemcpy(d_data, h_pageable, size, musaMemcpyHostToDevice);
    // 带宽：~3 GB/s（受限于 PCIe + 临时 buffer 拷贝）
    
    // 方式 2：锁定内存（Pinned Memory）
    float* h_pinned;
    musaMallocHost(&h_pinned, size);
    for (int i = 0; i < N; i++) h_pinned[i] = float(i);
    
    musaMemcpy(d_data, h_pinned, size, musaMemcpyHostToDevice);
    // 带宽：~6-8 GB/s（直接 PCIe 传输）
    
    // 方式 3：写合并内存（Write-Combined Memory）
    float* h_wc;
    musaMallocHost(&h_wc, size, musaHostAllocWriteCombined);
    for (int i = 0; i < N; i++) h_wc[i] = float(i);
    
    musaMemcpy(d_data, h_wc, size, musaMemcpyHostToDevice);
    // 带宽：~8-10 GB/s（优化的写入路径）
    
    // 清理
    free(h_pageable);
    musaFreeHost(h_pinned);
    musaFreeHost(h_wc);
    musaFree(d_data);
    musaFree(d_result);
}

---

L2 缓存管理

L2 缓存是 GPU 内存层次结构中的关键组件，位于全局内存和更高层缓存（如 L1/共享内存）之间。

L2 缓存预留（L2 Set-aside）

预留部分 L2 缓存给特定数据，减少关键数据被驱逐：

#include <musa.h>
#include <stdio.h>

int main() {
    // 查询 L2 缓存大小
    musaDeviceProp_t prop;
    int device;
    musaGetDevice(&device);
    musaGetDeviceProperties(&prop, device);
    
    printf("L2 cache size: %d bytes\n", prop.l2CacheSize);
    
    // 分配设备内存
    size_t size = 1000000 * sizeof(float);
    float* important_data;
    musaMalloc(&important_data, size);
    
    // 设置 L2 set-aside（预留 25% L2 缓存给该数据）
    size_t reserve_size = prop.l2CacheSize / 4;
    musaSetAttribute(important_data, musaAttributeL2CacheReserve, 
                     &reserve_size, sizeof(reserve_size));
    
    // 执行 kernel，important_data 更可能保留在 L2 中
    kernel1<<<grid, block>>>(important_data);
    
    musaFree(important_data);
    return 0;
}

多内核共享数据

void multiKernelAccess(float* shared_data, size_t size) {
    int device;
    musaGetDevice(&device);
    
    // 查询 L2 缓存大小
    int l2_size;
    musaDeviceGetAttribute(&l2_size, musaDeviceAttributeL2CacheSize, device);
    
    // 预留 50% L2 缓存给共享数据
    size_t reserve = l2_size / 2;
    musaSetAttribute(shared_data, musaAttributeL2CacheReserve, 
                     &reserve, sizeof(reserve));
    
    // 多个 kernel 顺序执行，共享数据保留在 L2 中
    kernel1<<<grid, block>>>(shared_data);
    kernel2<<<grid, block>>>(shared_data);
    kernel3<<<grid, block>>>(shared_data);
    
    // 无需担心中间数据被驱逐
}

访问策略窗口（Access Policy Window）

对于只访问一次的数据，使用流式策略避免污染 L2 缓存：

// 设置流式访问策略
void streamingAccess(float* streaming_data, size_t size) {
    musaAccessPolicyWindow policy;
    policy.base_ptr = streaming_data;
    policy.num_bytes = size;
    policy.hitRatio = 0.0f;  // 0% 命中率（只用一次）
    policy.hitProp = musaAccessPropertyStreaming;
    policy.missProp = musaAccessPropertyStreaming;
    
    musaSetAttribute(streaming_data, musaAttributeAccessPolicyWindow,
                     &policy, sizeof(policy));
    
    // 执行 kernel，数据不会污染 L2 缓存
    streamingKernel<<<grid, block>>>(streaming_data);
}

重用数据策略

对于频繁重用的数据，设置重用策略：

// 设置重用策略
void reuseAccess(float* reuse_data, size_t size) {
    musaAccessPolicyWindow policy;
    policy.base_ptr = reuse_data;
    policy.num_bytes = size;
    policy.hitRatio = 1.0f;  // 100% 命中率（频繁访问）
    policy.hitProp = musaAccessPropertyReuse;
    policy.missProp = musaAccessPropertyReuse;
    
    musaSetAttribute(reuse_data, musaAttributeAccessPolicyWindow,
                     &policy, sizeof(policy));
    
    // 多个 kernel 重复访问
    for (int i = 0; i < 10; i++) {
        processKernel<<<grid, block>>>(reuse_data);
    }
}

混合访问模式

void mixedAccessPattern(float* lookup_table, float* streaming_input, 
                        float* output, int N) {
    size_t table_size = 256 * sizeof(float);
    size_t input_size = N * sizeof(float);
    
    // Lookup table：重用策略
    musaAccessPolicyWindow table_policy;
    table_policy.base_ptr = lookup_table;
    table_policy.num_bytes = table_size;
    table_policy.hitRatio = 1.0f;
    table_policy.hitProp = musaAccessPropertyReuse;
    table_policy.missProp = musaAccessPropertyReuse;
    musaSetAttribute(lookup_table, musaAttributeAccessPolicyWindow,
                     &table_policy, sizeof(table_policy));
    
    // 输入数据：流式策略
    musaAccessPolicyWindow input_policy;
    input_policy.base_ptr = streaming_input;
    input_policy.num_bytes = input_size;
    input_policy.hitRatio = 0.0f;
    input_policy.hitProp = musaAccessPropertyStreaming;
    input_policy.missProp = musaAccessPropertyStreaming;
    musaSetAttribute(streaming_input, musaAttributeAccessPolicyWindow,
                     &input_policy, sizeof(input_policy));
    
    // 执行 kernel
    lookupKernel<<<grid, block>>>(lookup_table, streaming_input, output, N);
    musaDeviceSynchronize();
}

缓存行优化

// 查询 L2 缓存行大小
int cache_line_size;
musaDeviceGetAttribute(&cache_line_size, 
                       musaDeviceAttributeL2CacheLineSize, device);
printf("L2 cache line size: %d bytes\n", cache_line_size);

// 按缓存行对齐数据结构
struct alignas(128) CacheAlignedData {
    float data[32];   // 128 bytes，对齐缓存行
};

---

Cluster 内存

Cluster 内存（分布式共享内存）允许不同线程块（Thread Block）的线程共享数据并通信。

Cluster Group 基础

#include <musa.h>
#include <cooperative_groups.h>

using namespace cooperative_groups;

__global__ void clusterExample(float* data, int N) {
    // 获取 cluster group
    cluster_group g = this_cluster();
    
    // Cluster 信息
    int cluster_rank = g.thread_rank();      // 在 cluster 中的排名
    int cluster_size = g.size();             // cluster 总线程数
    int num_clusters = g.num_clusters();     // cluster 数量
    
    // Cluster 内线程 ID
    int cluster_tid = g.thread_index();
    
    if (cluster_rank == 0) {
        // 只有 cluster 0 执行
        data[0] = float(cluster_size);
    }
    
    // Cluster 内同步
    g.sync();
}

Cluster 共享内存

// 使用 clusterS__musa_cluster_sync() 进行 Cluster 同步
extern "C" __device__ void clusterS__musa_cluster_sync(unsigned int clusterSize);

__global__ void clusterSyncExample() {
    int clusterId = blockIdx.x;
    int threadId = threadIdx.x;
    
    // 每个线程初始化数据
    float localData = float(threadId) * float(clusterId);
    
    // 使用 clusterS__musa_cluster_sync() 同步
    // 参数为 cluster 内的 block 数量
    clusterS__musa_cluster_sync(gridDim.x);
    
    // 同步后可以访问其他 block 的数据
}

线程块间通信

使用 Cluster 共享内存实现线程块间数据交换：

__global__ void blockCommunication(float* input, float* output, int N) {
    clusterShared float cluster_buffer[1024];
    
    int cluster_idx = threadIdx.x + blockIdx.x * blockDim.x;
    
    // 步骤 1：写入本 block 数据
    cluster_buffer[cluster_idx] = input[cluster_idx];
    
    // Cluster 同步
    __cluster_sync();
    
    // 步骤 2：读取相邻 block 数据
    int left_idx = cluster_idx - blockDim.x;
    int right_idx = cluster_idx + blockDim.x;
    
    float left_value = (blockIdx.x > 0) ? cluster_buffer[left_idx] : 0.0f;
    float right_value = (blockIdx.x < 3) ? cluster_buffer[right_idx] : 0.0f;
    
    // 计算（使用邻居数据）
    output[cluster_idx] = left_value + cluster_buffer[cluster_idx] + right_value;
}

直方图示例（Cluster 优化版）

#include <musa.h>
#include <cooperative_groups.h>

using namespace cooperative_groups;

#define NUM_BINS 256
#define BLOCK_SIZE 256
#define CLUSTER_SIZE 512  // 2 blocks per cluster

__global__ void histogramCluster(unsigned int* input, 
                                 unsigned int* hist, int N) {
    cluster_group g = this_cluster();
    
    clusterShared unsigned int cluster_hist[NUM_BINS];
    
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int cluster_base = (blockIdx.x / 2) * CLUSTER_SIZE;
    int local_idx = threadIdx.x + (blockIdx.x % 2) * blockDim.x;
    
    // 初始化
    for (int i = threadIdx.x; i < NUM_BINS; i += blockDim.x) {
        cluster_hist[i] = 0;
    }
    g.sync();
    
    // 累加
    if (idx < N) {
        unsigned int bin = input[idx] % NUM_BINS;
        cluster_hist[bin] += 1;
    }
    g.sync();
    
    // 合并（每个 cluster 一个 thread 负责）
    if (local_idx == 0) {
        for (int i = 0; i < NUM_BINS; i++) {
            atomicAdd(&hist[i], cluster_hist[i]);
        }
    }
}

---

异步执行

流并发

使用多个流实现 Kernel 并发执行：

void concurrentKernels(float* data1, float* data2, float* data3, int n) {
    musaStream_t stream1, stream2, stream3;
    musaStreamCreate(&stream1);
    musaStreamCreate(&stream2);
    musaStreamCreate(&stream3);
    
    // 三个 kernel 在不同流中并发执行
    kernel1<<<grid, block, 0, stream1>>>(data1, n);
    kernel2<<<grid, block, 0, stream2>>>(data2, n);
    kernel3<<<grid, block, 0, stream3>>>(data3, n);
    
    // 等待所有完成
    musaStreamSynchronize(stream1);
    musaStreamSynchronize(stream2);
    musaStreamSynchronize(stream3);
    
    musaStreamDestroy(stream1);
    musaStreamDestroy(stream2);
    musaStreamDestroy(stream3);
}

依赖管理

使用事件实现流间依赖：

void streamDependency(float* data) {
    musaEvent_t event;
    musaEventCreate(&event);
    
    musaStream_t stream1, stream2;
    musaStreamCreate(&stream1);
    musaStreamCreate(&stream2);
    
    // stream1 执行 kernel
    kernel1<<<grid, block, 0, stream1>>>(data);
    
    // 在 stream1 中记录事件
    musaEventRecord(event, stream1);
    
    // stream2 等待事件完成后执行
    musaStreamWaitEvent(stream2, event, 0);
    kernel2<<<grid, block, 0, stream2>>>(data);
    
    musaStreamDestroy(stream1);
    musaStreamDestroy(stream2);
    musaEventDestroy(event);
}

异步内存操作

// 异步内存拷贝（必须在非默认流中）
musaStream_t stream;
musaStreamCreate(&stream);

musaMemcpyAsync(d_dst, h_src, size, musaMemcpyHostToDevice, stream);

// 异步内存设置
musaMemsetAsync(d_data, value, size, stream);

// 异步 2D 内存拷贝
musaMemcpy2DAsync(d_data, d_pitch, h_data, h_pitch,
                  width * sizeof(float), height,
                  musaMemcpyHostToDevice, stream);

完整示例：多流并发处理

#include <musa.h>
#include <stdio.h>

#define NUM_STREAMS 4
#define DATA_SIZE (1 << 20)  // 1MB

__global__ void processKernel(float* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        data[idx] = data[idx] * 2.0f + 1.0f;
    }
}

int main() {
    // 分配主机内存
    float* h_data = (float*)malloc(DATA_SIZE * sizeof(float));
    for (int i = 0; i < DATA_SIZE; i++) {
        h_data[i] = float(i);
    }
    
    // 分配设备内存
    float* d_data;
    musaMalloc(&d_data, DATA_SIZE * sizeof(float));
    
    // 创建多个流
    musaStream_t streams[NUM_STREAMS];
    for (int i = 0; i < NUM_STREAMS; i++) {
        musaStreamCreate(&streams[i]);
    }
    
    // 计算每个流的工作量
    int chunk_size = DATA_SIZE / NUM_STREAMS;
    int block_size = 256;
    
    // 异步拷贝到设备
    for (int i = 0; i < NUM_STREAMS; i++) {
        musaMemcpyAsync(d_data + i * chunk_size, 
                       h_data + i * chunk_size,
                       chunk_size * sizeof(float),
                       musaMemcpyHostToDevice,
                       streams[i]);
    }
    
    // 并发执行 kernel
    for (int i = 0; i < NUM_STREAMS; i++) {
        int grid_size = (chunk_size + block_size - 1) / block_size;
        processKernel<<<grid_size, block_size, 0, streams[i]>>>(
            d_data + i * chunk_size, chunk_size);
    }
    
    // 异步拷贝回主机
    for (int i = 0; i < NUM_STREAMS; i++) {
        musaMemcpyAsync(h_data + i * chunk_size,
                       d_data + i * chunk_size,
                       chunk_size * sizeof(float),
                       musaMemcpyDeviceToHost,
                       streams[i]);
    }
    
    // 等待所有完成
    musaStreamSynchronize(streams[0]);
    
    printf("Processing completed\n");
    
    // 清理
    for (int i = 0; i < NUM_STREAMS; i++) {
        musaStreamDestroy(streams[i]);
    }
    musaFree(d_data);
    free(h_data);
    
    return 0;
}

---

性能优化建议

锁定内存（Pinned Memory）使用原则

• 仅在需要异步传输时使用
• 用完立即释放，避免耗尽系统内存
• 复用 pinned buffer，避免频繁分配/释放

零拷贝选择指南

场景	推荐方式
一次性大数据传输	musaMemcpy + 锁定内存
频繁小数据传输	零拷贝
只读数据	零拷贝 + musaMemAdvise

L2 缓存优化

• 对频繁访问的关键数据使用 L2 预留缓存
• 对流式数据使用流式策略，避免污染缓存
• 对齐数据结构到缓存行边界

Cluster 内存适用场景

• 多 block 协作的算法（归约、扫描）
• 需要 block 间交换数据的 kernel
• 直方图统计等多 block 累加场景

---

常见问题

Q1: 锁定内存（Pinned Memory）有什么缺点？

A:

• 占用物理内存，减少系统可用内存
• 可能导致系统内存碎片化
• 过度分配可能影响系统性能

Q2: L2 预留缓存（L2 Set-aside）有什么限制？

A:

• 预留的 L2 缓存不能超过总 L2 的 75%
• 过度预留可能导致其他数据性能下降

Q3: Cluster 内存和共享内存（Shared Memory）有什么区别？

A:

• 共享内存（Shared Memory）：单个 block 内线程共享
• Cluster 共享内存：多个 block 共享
• Cluster 支持跨 block 同步（clusterSync()）

Q4: 异步内存拷贝有什么要求？

A:

• 必须使用非默认流
• 主机内存应该是锁定内存（Pinned Memory）
• 需要检查设备是否支持

---

相关文档

• 内存层次结构：内存层次基础
• Host/Device 模型：Host/Device 编程
• 线程层次结构：线程层次
• 执行模型：执行模型