CUDA Thrustライブラリの解決方法 - for_each同期エラー？

Question

私はCUDAの推力ライブラリを使用してCUDAで単純な動的ベクトルを修正しようとしています。しかし、エラーがいくつかの同期プロセスに関連していることを示す画面上に "launch_closure_by_value"エラーが表示されます。

このエラーのため、単純な1Dダイナミックアレイの変更はできません。

エラーの原因となっているコードセグメントは、次のとおりです。

void 
setIndexedGridInfo(float* a, float*b) 
{ 

    thrust::device_ptr<float> d_oldData(a); 
    thrust::device_ptr<float> d_newData(b); 

    float c = 0.0; 

    thrust::for_each(
     thrust::make_zip_iterator(thrust::make_tuple(d_oldData,d_newData)), 
     thrust::make_zip_iterator(thrust::make_tuple(d_oldData+8,d_newData+8)), 
     grid_functor(c)); 
} 

grid_functorが_kernel.cu

で定義されている：System.cu

float* a= (float*)(malloc(8*sizeof(float))); 
a[0]= 0; a[1]= 1; a[2]= 2; a[3]= 3; a[4]= 4; a[5]= 5; a[6]= 6; a[7]= 7; 
float* b = (float*)(malloc(8*sizeof(float))); 
setIndexedGridInfo(a,b); 

System.cuでコードセグメントに定義されている私はsetIndexedGridを呼び出す.cppファイルから

struct grid_functor 
{ 
    float a; 

    __host__ __device__ 
    grid_functor(float grid_Info) : a(grid_Info) {} 

    template <typename Tuple> 
    __device__ 
    void operator()(Tuple t) 
    { 
     volatile float data = thrust::get<0>(t); 
     float pos = data + 0.1; 
     thrust::get<1>(t) = pos; 
    } 

}; 

私は出力ウィンドウ（私はVisual Studioを使用しています）：

Particles.exeで0x000007fefdc7cacdで

初回例外： のMicrosoft C++の例外：メモリ位置でcudaError_enum 0x0029eb60 ...初回例外で0x000007fefdc7cacd でsmokeParticles.exe：マイクロソフトC++の例外： 推力::システムメモリ位置0x0029ecf0で:: SYSTEM_ERROR .. Particles.exeで0x000007fefdc7cacdで未処理の例外 ：マイクロソフトC++ 例外：メモリ位置 0x0029ecf0で推力::システム:: SYSTEM_ERROR ..

引き起こしている何

問題？

Answer 1

デバイスメモリにポインタが必要な関数でホストメモリポインタを使用しようとしています。

float* a= (float*)(malloc(8*sizeof(float))); 
a[0]= 0; a[1]= 1; a[2]= 2; a[3]= 3; a[4]= 4; a[5]= 5; a[6]= 6; a[7]= 7; 
float* b = (float*)(malloc(8*sizeof(float))); 
setIndexedGridInfo(a,b); 

..... 

thrust::device_ptr<float> d_oldData(a); 
thrust::device_ptr<float> d_newData(b); 

thrust::device_ptrは「ラッピング」という推力がそれを使用することができますので、CUDAのAPIに割り当てられたデバイスメモリポインタを対象としています。このコードは、問題です。ホストポインターを直接デバイスポインターとして扱おうとしています。それは違法です。あなたはこのようなあなたのsetIndexedGridInfo機能を変更することができます：

void setIndexedGridInfo(float* a, float*b, const int n) 
{ 

    thrust::device_vector<float> d_oldData(a,a+n); 
    thrust::device_vector<float> d_newData(b,b+n); 

    float c = 0.0; 

    thrust::for_each(
     thrust::make_zip_iterator(thrust::make_tuple(d_oldData.begin(),d_newData.begin())), 
     thrust::make_zip_iterator(thrust::make_tuple(d_oldData.end(),d_newData.end())), 
     grid_functor(c)); 
} 

device_vectorコンストラクタは、デバイスのメモリを割り当てた後、デバイスにホストメモリの内容をコピーします。それはあなたが見ているエラーを修正するはずですが、私はあなたがfor_eachイテレータで何をしようとしているのか、あなたがwrttienを持っているファンクタが正しいかどうかはわかりません。

---

編集：

#include <cstdlib> 
#include <cstdio> 
#include <thrust/device_vector.h> 
#include <thrust/for_each.h> 
#include <thrust/copy.h> 

struct grid_functor 
{ 
    float a; 

    __host__ __device__ 
    grid_functor(float grid_Info) : a(grid_Info) {} 

    template <typename Tuple> 
    __device__ 
    void operator()(Tuple t) 
    { 
     volatile float data = thrust::get<0>(t); 
     float pos = data + 0.1f; 
     thrust::get<1>(t) = pos; 
    } 

}; 

void setIndexedGridInfo(float* a, float*b, const int n) 
{ 

    thrust::device_vector<float> d_oldData(a,a+n); 
    thrust::device_vector<float> d_newData(b,b+n); 

    float c = 0.0; 

    thrust::for_each(
     thrust::make_zip_iterator(thrust::make_tuple(d_oldData.begin(),d_newData.begin())), 
     thrust::make_zip_iterator(thrust::make_tuple(d_oldData.end(),d_newData.end())), 
     grid_functor(c)); 

    thrust::copy(d_newData.begin(), d_newData.end(), b); 
} 

int main(void) 
{ 
    const int n = 8; 
    float* a= (float*)(malloc(n*sizeof(float))); 
    a[0]= 0; a[1]= 1; a[2]= 2; a[3]= 3; a[4]= 4; a[5]= 5; a[6]= 6; a[7]= 7; 
    float* b = (float*)(malloc(n*sizeof(float))); 
    setIndexedGridInfo(a,b,n); 

    for(int i=0; i<n; i++) { 
     fprintf(stdout, "%d (%f,%f)\n", i, a[i], b[i]); 
    } 

    return 0; 
} 

私はOS 10.6.8ホスト上でこのコードをコンパイルして実行することができます：ここで

はあなたのコードの完全な、コンパイル、実行可能なバージョンがありますCUDA 4。1このように：

$ nvcc -Xptxas="-v" -arch=sm_12 -g -G thrustforeach.cu 
./thrustforeach.cu(18): Warning: Cannot tell what pointer points to, assuming global memory space 
./thrustforeach.cu(20): Warning: Cannot tell what pointer points to, assuming global memory space 
./thrustforeach.cu(18): Warning: Cannot tell what pointer points to, assuming global memory space 
./thrustforeach.cu(20): Warning: Cannot tell what pointer points to, assuming global memory space 
ptxas info : Compiling entry function '_ZN6thrust6detail7backend4cuda6detail23launch_closure_by_valueINS2_18for_each_n_closureINS_12zip_iteratorINS_5tupleINS0_15normal_iteratorINS_10device_ptrIfEEEESB_NS_9null_typeESC_SC_SC_SC_SC_SC_SC_EEEEi12grid_functorEEEEvT_' for 'sm_12' 
ptxas info : Used 14 registers, 160+0 bytes lmem, 16+16 bytes smem, 4 bytes cmem[1] 
ptxas info : Compiling entry function '_ZN6thrust6detail7backend4cuda6detail23launch_closure_by_valueINS2_18for_each_n_closureINS_12zip_iteratorINS_5tupleINS0_15normal_iteratorINS_10device_ptrIfEEEESB_NS_9null_typeESC_SC_SC_SC_SC_SC_SC_EEEEj12grid_functorEEEEvT_' for 'sm_12' 
ptxas info : Used 14 registers, 160+0 bytes lmem, 16+16 bytes smem, 4 bytes cmem[1] 

$ ./a.out 
0 (0.000000,0.100000) 
1 (1.000000,1.100000) 
2 (2.000000,2.100000) 
3 (3.000000,3.100000) 
4 (4.000000,4.100000) 
5 (5.000000,5.100000) 
6 (6.000000,6.100000) 
7 (7.000000,7.100000)