Reverse Array

原始题目：LeetGPU - Reverse Array

题目描述

编写一个 GPU 程序，原地反转 32 位浮点数数组。程序应对输入数组执行原地反转——即第 $i$ 个元素与第 $N - 1 - i$ 个元素交换。

实现要求

• 不允许使用外部库。
• solve 函数签名必须保持不变。
• 最终结果必须存储回 input 数组中。

示例

示例 1

Input:  [1.0, 2.0, 3.0, 4.0]
Output: [4.0, 3.0, 2.0, 1.0]

示例 2

Input:  [1.5, 2.5, 3.5]
Output: [3.5, 2.5, 1.5]

约束条件

• $1 \le N \le 100{,}000{,}000$。
• 性能测试在 $N = 25{,}000{,}000$ 的规模下进行。

解题思路

反转操作的本质是交换对称位置上的元素对。每个线程可以负责交换一对元素 $(i, N-1-i)$，需要 $N/2$ 个线程。注意处理 $N$ 为奇数时中间元素无需交换的情况。内存访问模式是非合并的（前半部分访问低地址，后半部分访问高地址），但交换操作使得总访问是平衡的。

代码实现

CUDA

#include <cuda_runtime.h>

__global__ void reverse_kernel(float* data, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int half = N / 2;
    if (i < half) {
        int j = N - 1 - i;
        float tmp = data[i];
        data[i] = data[j];
        data[j] = tmp;
    }
}

extern "C" void solve(float* data, int N) {
    int half = N / 2;
    int threadsPerBlock = 256;
    int blocksPerGrid = (half + threadsPerBlock - 1) / threadsPerBlock;
    reverse_kernel<<<blocksPerGrid, threadsPerBlock>>>(data, N);
    cudaDeviceSynchronize();
}

Triton

import triton
import triton.language as tl

@triton.jit
def reverse_kernel(data_ptr, N: tl.constexpr, BLOCK_SIZE: tl.constexpr):
    i = tl.program_id(0) * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
    mask = i < N // 2
    j = N - 1 - i
    a = tl.load(data_ptr + i, mask=mask)
    b = tl.load(data_ptr + j, mask=mask)
    tl.store(data_ptr + i, b, mask=mask)
    tl.store(data_ptr + j, a, mask=mask)