qudarap/tests/crovella_t66.cu

0001 /**
0002 name=crovella_t66 ; nvcc $name.cu -o /tmp/$name && /tmp/$name
0003
0004 https://stackoverflow.com/questions/64464892/how-to-use-thrustcopy-if-using-pointers
0005
0006 **/
0007
0008 #include <thrust/copy.h>
0009 #include <iostream>
0010 #include <iomanip>
0011 #include <thrust/device_ptr.h>
0012
0013 struct is_not_zero
0014 {
0015     __host__ __device__
0016     bool operator()( double x)
0017     {
0018         return (x != 0);
0019     }
0020 };
0021
0022
0023 int main()
0024 {
0025     const int ds = 20 ;
0026
0027     float src[ds];
0028     for (int i = 0; i < ds; i++) src[i] = i % 3 == 0 ? 0.f : 100.5f + float(i)  ;
0029
0030     float* d_src ;
0031     cudaMalloc(&d_src, ds*sizeof(float));
0032     cudaMemcpy(d_src, src, ds*sizeof(float), cudaMemcpyHostToDevice);
0033     thrust::device_ptr<float> t_src(d_src);
0034
0035     is_not_zero predicate ;
0036     size_t num_sel = thrust::count_if(t_src, t_src+ds, predicate );
0037     std::cout << " num_sel " << num_sel << std::endl ;
0038
0039     float* dst = new float[num_sel] ;
0040     for (int i = 0; i < num_sel ; i++) dst[i] = 0.f ;
0041
0042     float* d_dst ;
0043     cudaMalloc(&d_dst,        num_sel*sizeof(float));
0044     cudaMemset(d_dst, 0,      num_sel*sizeof(float));
0045     thrust::device_ptr<float> t_dst(d_dst);   // thrust::device_pointer_cast(d_dst);
0046
0047     thrust::copy_if(t_src, t_src+ds, t_dst, predicate );
0048
0049     cudaMemcpy(dst, d_dst, num_sel*sizeof(float), cudaMemcpyDeviceToHost);
0050
0051     for(int i=0;i<ds;i++)
0052         std::cout
0053             << " src " << std::setw(10) << ( i < ds ? src[i] : -1.f )
0054             << " dst " << std::setw(10) << ( i < num_sel ? dst[i] : -1.f )
0055             << std::endl
0056             ;
0057
0058 }