untitled

#include <cuda_runtime.h>
#define BLOCK_DIM 256

__global__ void reverse_array_basic(float* input, int N) {
    int idx = blockDim.x * blockIdx.x + threadIdx.x;
    if (idx <= N/2) {
        int rev_idx = N - 1 - idx;
        float tmp = input[idx];
        input[idx] = input[rev_idx];
        input[rev_idx] = tmp;
    }
}

__global__ void reverse_array_smem(float* input, int N) {
    int idx = blockDim.x * blockIdx.x + threadIdx.x;
    int rev_idx = N - 1 - idx;
    __shared__ float left[BLOCK_DIM], right[BLOCK_DIM];

    if (idx <= N/2) {
        left[threadIdx.x] = input[idx];
        right[threadIdx.x] = input[rev_idx];
    }
    __syncthreads();

    if (idx <= N/2) {
        input[rev_idx] = left[threadIdx.x];
        input[idx] = right[threadIdx.x];
    }

}

// input is device pointer
extern "C" void solve(float* input, int N) {
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;

    //reverse_array_basic<<<blocksPerGrid, threadsPerBlock>>>(input, N);
    //reverse_array_smem<<<blocksPerGrid, threadsPerBlock>>>(input, N);
    
    cudaDeviceSynchronize();
}