CUDA中严格别名规则与通用拷贝核的实现及问题咨询

我需要拷贝多个不同类型的数组，示例代码如下：

cudaMemcpy(dst0, src0, num0 * sizeof(int), cudaMemcpyDeviceToDevice);
cudaMemcpy(dst1, src1, num1 * sizeof(float), cudaMemcpyDeviceToDevice);
cudaMemcpy(dst2, src2, num2 * sizeof(int16_t), cudaMemcpyDeviceToDevice);

由于单个数组规模不大，使用单个拷贝核函数比多次调用cudaMemcpy效率更高，核函数框架如下：

__global__ void copy_kernel(void **dst, void **src, int num_arrays, const int* array_sizes, const int* type_sizes) {...}

int main() {
    ... // prepare data etc.
    copy_kernel<<<num_arrays, 1024>>>(dst, src, num_arrays, array_sizes, type_sizes);
    ...
}

为简化说明，后续仅考虑单个float类型数组（假设sizeof(float)=sizeof(int32_t)=4）的情况，以下是三种实现方案：

三种实现方案 #

1. 专用float拷贝核 #

简单高效，但仅支持float类型，通用性差

__global__ void copy_kernel_float(float* dst, const float* src, int num) {
    for (int i = threadIdx.x; i < num; i += blockDim.x) {
        dst[i] = src[i];
    }
}

2. 通用字符拷贝核 #

支持任意类型，但性能较低，推测原因是32字节内存事务比128字节事务慢

__global__ void copy_kernel_any(void* dst, const void* src, int num,
                                int type_size) {
    int num_char = num * type_size;
    auto my_dst = reinterpret_cast<char*>(dst);
    auto my_src = reinterpret_cast<const char*>(src);
    for (int i = threadIdx.x; i < num_char; i += blockDim.x) {
        my_dst[i] = my_src[i];
    }
}

3. 通用int32_t打包拷贝核 #

支持任意类型且性能与专用核相当，但疑似违反严格别名规则（假设num * type_size % sizeof(int32_t) = 0）

__global__ void copy_kernel_any_packed(void* dst, const void* src, int num,
                                       int type_size) {
    int num_int32 = num * type_size / sizeof(int32_t);
    auto my_dst = reinterpret_cast<int32_t*>(dst);
    auto my_src = reinterpret_cast<const int32_t*>(src);
    for (int i = threadIdx.x; i < num_int32; i += blockDim.x) {
        my_dst[i] = my_src[i];
    }
}

技术问题 #

1. 第三种实现是否合法？已知在C++中reinterpret_cast<float*>到int32_t*属于未定义行为，且reinterpret_cast<int32_t*>(reinterpret_cast<char*>(dst))也无法规避该问题？
2. 关于内存事务性能的推测是否正确？我在1070Ti上测试拷贝1000万个float值的两个合法核函数，float拷贝速度与cudaMemcpy相当（约0.4ms），char拷贝则更慢（约0.6ms）？
3. 若第三种实现违反严格别名规则，是否存在既合法又高效的实现方案（例如单指令拷贝4个字符）？

若我的表述存在错误，欢迎指正。

问题解答 #

1. 第三种实现的合法性分析 #

严格来说，该实现违反了C++严格别名规则。根据C++标准，除char、unsigned char、std::byte外，不同类型的指针不能用来访问对象的内存——即使通过char指针中转再强转，最终用int32_t指针操作实际为float类型的对象内存，仍属于未定义行为。

虽然CUDA编译器（nvcc）对严格别名规则的处理有一定宽松度，实际运行可能不会触发问题，但这并不代表代码合法。一旦编译器优化策略调整或移植到其他标准编译器环境，可能出现不可预测的行为。

2. 内存事务性能推测的正确性 #

你的推测基本正确。以1070Ti所属的Pascal架构GPU为例：

• 专用float拷贝核每次访问4字节（一个float），32线程的线程束一次访问32*4=128字节，刚好匹配Pascal架构L1缓存的事务粒度，能充分利用内存带宽。
• 字符拷贝核每次仅访问1字节，线程束一次仅处理32字节，属于小粒度事务，GPU需要额外的拆分/合并操作，降低了带宽利用率，导致性能变慢。

你的测试数据也验证了这一点：float拷贝的带宽接近硬件峰值，而char拷贝的带宽明显更低。

3. 合法且高效的替代实现 #

可以通过基于char/std::byte的批量拷贝结合编译器自动向量化，或使用CUDA内置工具，实现合法且高效的拷贝：

方案A：编译器自动向量化的批量拷贝

调整核函数让线程每次拷贝多个字节，编译器会自动优化为向量指令：

__global__ void copy_kernel_optimized(void* dst, const void* src, int num, int type_size) {
    size_t total_bytes = static_cast<size_t>(num) * type_size;
    char* dst_char = static_cast<char*>(dst);
    const char* src_char = static_cast<const char*>(src);

    // 按4字节批量拷贝，剩余字节单独处理
    size_t i = threadIdx.x * 4;
    for (; i + 4 <= total_bytes; i += blockDim.x * 4) {
        *reinterpret_cast<uint32_t*>(dst_char + i) = *reinterpret_cast<const uint32_t*>(src_char + i);
    }
    // 处理剩余1-3字节
    for (; i < total_bytes; i += blockDim.x) {
        dst_char[i] = src_char[i];
    }
}

这种方式通过char指针计算偏移后强转，本质是对内存字节流的批量操作，符合严格别名规则，且能触发编译器的向量优化，提升性能。

方案B：使用CUDA内置向量类型

利用uint4等向量类型一次处理16字节，最大化内存带宽利用率：

__global__ void copy_kernel_vectorized(void* dst, const void* src, int num, int type_size) {
    size_t total_bytes = static_cast<size_t>(num) * type_size;
    char* dst_char = static_cast<char*>(dst);
    const char* src_char = static_cast<const char*>(src);

    // 按16字节批量拷贝
    size_t i = threadIdx.x * 16;
    for (; i + 16 <= total_bytes; i += blockDim.x * 16) {
        *reinterpret_cast<uint4*>(dst_char + i) = *reinterpret_cast<const uint4*>(src_char + i);
    }
    // 处理剩余字节：先4字节，再1字节
    for (; i + 4 <= total_bytes; i += blockDim.x * 4) {
        *reinterpret_cast<uint32_t*>(dst_char + i) = *reinterpret_cast<const uint32_t*>(src_char + i);
    }
    for (; i < total_bytes; i += blockDim.x) {
        dst_char[i] = src_char[i];
    }
}

方案C：使用CUDA内置__memcpy函数

直接调用硬件加速的__memcpy，由编译器优化为最优拷贝指令：

__global__ void copy_kernel_builtin(void** dst, void** src, int num_arrays, const int* array_sizes, const int* type_sizes) {
    int idx = blockIdx.x;
    if (idx >= num_arrays) return;
    size_t total_bytes = static_cast<size_t>(array_sizes[idx]) * type_sizes[idx];
    __memcpy(dst[idx], src[idx], total_bytes);
}

该方法既符合C++标准，又能达到接近cudaMemcpy的性能。

内容的提问来源于stack exchange，提问作者courage