CUDA - デバイスアレイポインタを交換するには？

Question

私は便利なデバイス配列ポインタをスワップするために関数swapを書いたが、それは動作していない、私はスワップ機能でローカル配列ポインタを交換していると私はそれに渡しているものではないと仮定します。

__global__ void device_add_one(float *A, float *B) 
{ 
    for (int index = blockIdx.x * blockDim.x + threadIdx.x; 
     index < N; 
     index += blockDim.x * gridDim.x) 
    { 
     // just for the example 
     B[index] = A[index] + 1; 
    { 
} 

void swap(float *a, float *b) 
{ 
    float *temp = a; 
    a = b; 
    b = temp; 
} 

void loop(float *host_array, int size, int loops) 
{ 
    cudaMalloc(&A, (size * sizeof(float)); 
    cudaMalloc(&B, (size * sizeof(float)); 

    cudaMemcpy(A, host_array, (size * sizeof(float), cudaMemcpyHostToDevice); 

    for (int i = 0; i < loops; i++) { 
     device_add_one<<< 1, 254 >>>(A, B); 

     // swap pointers like this does not work 
     swap(A, B); 

     /* This works: 
     float *temp = a; 
     a = b; 
     b = temp; 
     */ 
    } 

    cudaMemcpy(host_array, A, (size * sizeof(float), cudaMemcpyDeviceToHost); 
} 

Answer 1

値渡しを使用しているため、ポインタを交換する関数呼び出しメソッドが機能しません。これは通常のC/C++プログラミングの概念であり、CUDA特有のものではありません。

あなたが機能by valueに（この場合はポインタを含む）変数を渡す場合：

void swap(float *a, float *b) 

C値渡しメカニズムが機能内で使用するために、関数の引数のローカルコピーを作成します体。それらの引数の変更は、呼び出しコンテキストに表示されません。この問題を回避するには、簡単なアプローチは、参照渡しされます（C++で）：ここでは

void swap(float* &a, float* &b) 

働いた例です。

$ cat t393.cu 
#include <stdio.h> 

const int N = 1000; 
float *A, *B; 

__global__ void device_add_one(float *A, float *B) 
{ 
    for (int index = blockIdx.x * blockDim.x + threadIdx.x; 
     index < N; 
     index += blockDim.x * gridDim.x) 
    { 
     // just for the example 
     B[index] = A[index] + 1; 
    } 
} 
void swap(float* &a, float* &b){ 
    float *temp = a; 
    a = b; 
    b = temp; 
} 

void loop(float *host_array, int size, int loops) 
{ 
    cudaMalloc(&A, size * sizeof(float)); 
    cudaMalloc(&B, size * sizeof(float)); 

    cudaMemcpy(A, host_array, (size * sizeof(float)), cudaMemcpyHostToDevice); 

    for (int i = 0; i < loops; i++) { 
     device_add_one<<< 1, 254 >>>(A, B); 

     // swap pointers 
     swap(A, B); 
     //float *temp = A; 
     //A = B; 
     //B = temp; 
    } 

    cudaMemcpy(host_array, A, (size * sizeof(float)), cudaMemcpyDeviceToHost); 
} 

int main(){ 

    float *data = (float *)malloc(N*sizeof(float)); 
    for (int i = 0; i<N; i++) data[i] = i & 3; // fill with 0 1 2 3 0 1 2 3... 
    loop(data, N, 100); 
    for (int i = 0; i<20; i++) printf("%f ", data[i]); 
    printf("\n"); 
    return 0; 
} 
$ nvcc -arch=sm_61 -o t393 t393.cu 
$ cuda-memcheck ./t393 
========= CUDA-MEMCHECK 
100.000000 101.000000 102.000000 103.000000 100.000000 101.000000 102.000000 103.000000 100.000000 101.000000 102.000000 103.000000 100.000000 101.000000 102.000000 103.000000 100.000000 101.000000 102.000000 103.000000 
========= ERROR SUMMARY: 0 errors 
$