如何将使用std::array的C++代码便捷移植到CUDA Thrust？

Great question! You're absolutely right that std::array isn't usable in CUDA device code—it lacks __device__-qualified member functions and isn't designed for the GPU's memory model. thrust::device_vector is an excellent replacement here, and we can adjust your code to keep that clean, unified macro (or even improve it with type-safe aliases) while making it CUDA-compatible. Let's break this down step by step:

1. Replace std::array with thrust::device_vector in your macro #

First, swap out the header and update your macro to use Thrust's device-side container. Unlike std::array, device_vector manages GPU memory automatically and works seamlessly with both host and device code:

#include <thrust/device_vector.h>
const unsigned int arraySize = 1024;
#define ARRAY_DEF thrust::device_vector<int>

int main() {
    ARRAY_DEF x(arraySize); // Initialize with your fixed size
    // Replace std::array::fill() with Thrust's fill algorithm
    thrust::fill(x.begin(), x.end(), 1);
    // Or use device_vector::assign() if you prefer:
    // x.assign(arraySize, 1);
    return 0;
}

2. Using the container in CUDA kernels #

If you need to operate on this data in a __global__ kernel, you can get a raw device pointer to the underlying memory with x.data(), which works just like a regular GPU pointer:

__global__ void processArray(int* arr, size_t size) {
    const int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < size) {
        // Example: Double each element
        arr[idx] *= 2;
    }
}

int main() {
    ARRAY_DEF x(arraySize);
    thrust::fill(x.begin(), x.end(), 1);
    
    // Launch kernel with enough threads to cover the array
    const dim3 blockSize(256);
    const dim3 gridSize((arraySize + blockSize.x - 1) / blockSize.x);
    processArray<<<gridSize, blockSize>>>(x.data(), x.size());
    
    // Wait for kernel to finish before accessing results
    cudaDeviceSynchronize();
    return 0;
}

3. Optional: Keep compile-time fixed size (like std::array) #

If you specifically need the compile-time size guarantee of std::array, you can use a native CUDA device array with conditional compilation to split host/device types. Note that this requires manual memory management (copying between host and device):

#include <array>
#include <thrust/device_vector.h>

const unsigned int arraySize = 1024;

// Use std::array on host, native device array on GPU
#ifdef __CUDA_ARCH__
#define ARRAY_DEF int[arraySize]
#else
#define ARRAY_DEF std::array<int, arraySize>
#endif

__global__ void processArray(ARRAY_DEF arr) {
    const int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < arraySize) {
        arr[idx] *= 2;
    }
}

int main() {
    // Host-side array
    ARRAY_DEF hostArr;
    hostArr.fill(1);
    
    // Allocate and copy to device
    ARRAY_DEF* devArr;
    cudaMalloc(&devArr, sizeof(hostArr));
    cudaMemcpy(devArr, hostArr.data(), sizeof(hostArr), cudaMemcpyHostToDevice);
    
    // Launch kernel
    const dim3 blockSize(256);
    const dim3 gridSize((arraySize + blockSize.x - 1) / blockSize.x);
    processArray<<<gridSize, blockSize>>>(*devArr);
    
    // Copy results back to host
    cudaMemcpy(hostArr.data(), devArr, sizeof(hostArr), cudaMemcpyDeviceToHost);
    cudaFree(devArr);
    
    return 0;
}

4. Pro tip: Replace macros with type-safe aliases #

For better maintainability and type safety, ditch the macro and use a using declaration instead. It does the same job but lets the compiler catch type errors early:

using ARRAY_TYPE = thrust::device_vector<int>;
// Then use ARRAY_TYPE x(arraySize); instead of ARRAY_DEF x;

Final notes #

thrust::device_vector is the most straightforward drop-in replacement because it handles memory allocation/deallocation, supports STL-like algorithms, and bridges host/device code smoothly. If you need more control over memory (like using pinned host memory), Thrust also has tools for that—thrust::pinned_vector can speed up host-device transfers significantly.

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