动态分配的数组副本在CUDA中

Question

我们可以使用任何方法复制cuda内核中的数组？

例如：

__device__ int number_element; __device__ void copyData(float* input, float* output){}

我想将某些输入数组中的数据复制到满足某些条件并将复制元素数量的数量的 nubly_element的数量 em>

谢谢。

Answer 1

您真正描述的内容称为流压实。推力库具有内置的一系列流压实功能，可以从内核中调用。作为一个琐碎的例子：

#include <iostream>
#include <thrust/copy.h>
#include <thrust/execution_policy.h>
struct op
{
  __host__ __device__
  bool operator()(const int x) { return (x % 3) == 0; }
};
__global__ void kernel(int* input, int* output, int Nin, int* Nout)
{
    auto output_end = thrust::copy_if(thrust::device, input, input + Nin, output, op());
    *Nout = output_end - output;
}
int main()
{
    const int N = 10;
    const size_t sz = sizeof(int) * size_t(N);
    int* in;
    cudaMallocManaged((void **)&in, sz);
    int* out;
    cudaMallocManaged((void **)&out, sz);
    int* Nout;
    cudaMallocManaged((void **)&Nout, sizeof(int));
    for(int i=0; i<N; i++) {
        in[i] = 1+i;
        out[i] = -1;
    }
    kernel<<<1,1>>>(in, out, N, Nout);
    cudaDeviceSynchronize();
    for(int i=0; i < *Nout; i++) {
        std::cout << i << " " << out[i] << std::endl;
    }
    return 0;
}

汇编和运行类似：

$ nvcc -std=c++11 -arch=sm_52 thrust_device_compact.cu 
$ ./a.out 
0 3
1 6
2 9

这可能是在少量数据上执行内核中流动压实的快速简便方法。如果您有很多数据，那么使用主机的推力并代表您的推力运行核可能更有意义。

Answer 2

是的，您可以通过编写一个。

您可以在此答案中执行以下方式：将全局复制到共享内存中，只需跳过smem零件。

//assumes sizeof(T) is multiple of sizeof(int) and is aligned to at least alignof(int)
//assumes single-dimention kernel
//assumes it is launched for all threads in block
template <typename T>
__device__ void memCopy(T* dest, T* src, size_t size) {
    int* iDest = (int*)dest;
    int* iSrc = (int*)src;
    for(size_t i = threadIdx.x; i<size*sizeof(T)/sizeof(int); i+=blockDim.x)
        iDest[i] = iSrc[i];
    __syncthreads();
}

这是一个单个块操作，旨在用于该特定块。如果您想要一个整个网格，则可以证明您可以将其作为单独的内核启动，以确保所有其他块都可以看到所有写作。在这种情况下，cudaMemcpy可能比内核调用更好。

无论如何，对于网格操作，您需要更改循环：

for(size_t i = threadIdx.x+blockIdx.x*blockDim.x;
    i<size*sizeof(T)/sizeof(int);
    i+=blockDim.x*gridDim.x)