CUDAで動的に割り当てられた配列コピー

Question

CUDAカーネルで配列をコピーする方法はありますか？私はいくつかの条件を満たして出力に入力配列の一部からデータをコピーしてもnumber_elementにコピーされた要素の数を取得したい

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

：たとえば

ありがとうございます。

Answer 1

はい、作成することができます。

たとえば、答えがBest way to copy global into shared memoryの場合は、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) 

Answer 2

あなたが本当にストリームコンパクションと呼ばれる記述されているもの：いずれの場合で

は、グリッドの動作のために、あなたはループを変更する必要があります。推力ライブラリには、カーネル内から呼び出すことができる一連のストリーム圧縮関数が組み込まれています。コンパイルし、そのように動作します

#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 

これは、少量のデータにカーネル内ストリームの圧縮を実行する迅速かつ簡単な方法かもしれません。大量のデータがある場合は、ホストからの推力を使用して、あなたの代わりにスラストランカーネルを使用するのがおそらくより理にかなっています。