CUDA: медленная тяга :: уменьшить после вызова тяги :: for_each_n - PullRequest
Новая
       48

CUDA: медленная тяга :: уменьшить после вызова тяги :: for_each_n

0 голосов
/ 04 июля 2018

Я пытаюсь взять сумму чисел, используя thrust с GK107 [GeForce GTX 650]. Я смущен, наблюдая, что время выполнения для thrust::reduce значительно увеличивается сразу после инициализации device_vector<curandState> в памяти.

Ниже приведен пример кода: 

#include <iostream>
#include <stack>
#include <ctime>

#include <thrust/device_vector.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/transform.h>
#include <thrust/for_each.h>
#include <curand.h>
#include <curand_kernel.h>

struct tic_toc{
    std::stack<clock_t> tictoc_stack;
    inline void tic() { tictoc_stack.push(clock());}
    inline void toc() {
        std::cout << "Time elapsed: "
            << ((double)(clock() - tictoc_stack.top())) / CLOCKS_PER_SEC << "s"
            << std::endl;
        tictoc_stack.pop();
    }
};

struct curand_setup{
    using init_tuple = thrust::tuple<int, curandState &>;
    const unsigned long long seed;
    curand_setup(unsigned long long _seed) : seed(_seed) {}
    __device__ void operator()(init_tuple t){
        curandState s;
        int id = thrust::get<0>(t);
        curand_init(seed, id, 0, &s);
        thrust::get<1>(t) = s;
    }
};

int main(int argc, char** argv){
    int N = 1<<18;
    std::cout << "N " << N << std::endl;
    tic_toc tt;

    thrust::device_vector<float> val(N,1);

    tt.tic();
    float mean=thrust::reduce(val.begin(),val.end(),0.f,thrust::plus<float>())/N;
    tt.toc();

    thrust::device_vector<curandState> rand_state(N);
    auto rand_init_it = thrust::make_zip_iterator(
            thrust::make_tuple(thrust::counting_iterator<int>(0),rand_state.begin()));
    thrust::for_each_n(rand_init_it, N, curand_setup(0));

    tt.tic();
    mean=thrust::reduce(val.begin(),val.end(),0.f,thrust::plus<float>())/N;
    tt.toc();

    tt.tic();
    mean=thrust::reduce(val.begin(),val.end(),0.f,thrust::plus<float>())/N;
    tt.toc();

    return 0;
}

и вывод: 

Time elapsed: 0.000594s
Time elapsed: 5.60026s
Time elapsed: 0.001098s

Ситуация не изменилась, когда я написал свое собственное ядро ​​для суммирования или скопировал данные в thrust::host_vector и уменьшил их.

Почему thrust::reduce такой медленный сразу после инициализации thrust::device_vector<curandState>, и есть ли способ избежать этой проблемы? Буду признателен за помощь.

Моя система Linux Mint 18.3 с ядром 4.15.0-23-generic.

вывод nvcc --version: nvcc: NVIDIA (R) Cuda compiler driver Copyright (c) 2005-2015 NVIDIA Corporation Built on Tue_Aug_11_14:27:32_CDT_2015 Cuda compilation tools, release 7.5, V7.5.17

1 Ответ

0 голосов
/ 04 июля 2018

Почему thrust::reduce такой медленный сразу после инициализации thrust::device_vector<curandState>

Это не так. Источником вашей путаницы является неправильное измерение времени.

Как правило, вызовы API Thrust, которые работают на устройстве, являются асинхронными на хосте. Единственным исключением являются вызовы, которые возвращают значение (и thrust::reduce является одним из них). В результате средний вызов в вашем коде измеряет не только время выполнения thrust::reduce, но и предыдущий вызов thrust::for_each_n, и это тот предшествующий вызов, который намного медленнее.

Вы можете подтвердить это для себя двумя способами. Если вы измените свой код тяги следующим образом: 

tt.tic();
float mean=thrust::reduce(val.begin(),val.end(),0.f,thrust::plus<float>())/N;
tt.toc();

thrust::device_vector<curandState> rand_state(N);
auto rand_init_it = thrust::make_zip_iterator(
        thrust::make_tuple(thrust::counting_iterator<int>(0),rand_state.begin()));
thrust::for_each_n(rand_init_it, N, curand_setup(0));
cudaDeviceSynchronize(); // wait until for_each is complete

tt.tic();
mean=thrust::reduce(val.begin(),val.end(),0.f,thrust::plus<float>())/N;
tt.toc();

tt.tic();
mean=thrust::reduce(val.begin(),val.end(),0.f,thrust::plus<float>())/N;
tt.toc();

Вы должны получить что-то вроде этого: 

$ nvcc -arch=sm_52 -std=c++11 -o slow_thrust slow_thrust.cu 
$ ./slow_thrust 
N 262144
Time elapsed: 0.000471s
Time elapsed: 0.000621s
Time elapsed: 0.000448s

т.е. когда вы используете cudaDeviceSynchronize() для захвата времени выполнения предыдущего вызова, все вызовы сокращения имеют примерно одинаковое время выполнения. В качестве альтернативы вы можете использовать инструмент профилирования в исходном коде, например: 

$ nvprof --print-gpu-trace ./slow_thrust
N 262144
==7870== NVPROF is profiling process 7870, command: ./slow_thrust
Time elapsed: 0.000521s
Time elapsed: 0.06983s
Time elapsed: 0.000538s
==7870== Profiling application: ./slow_thrust
==7870== Profiling result:
   Start  Duration            Grid Size      Block Size     Regs*    SSMem*    DSMem*      Size  Throughput  SrcMemType  DstMemType           Device   Context    Stream  Name
214.30ms  7.6800us            (512 1 1)       (256 1 1)         8        0B        0B         -           -           -           -  GeForce GTX 970         1         7  void thrust::cuda_cub::core::_kernel_agent<thrust::cuda_cub::__parallel_for::ParallelForAgent<thrust::cuda_cub::__uninitialized_fill::functor<thrust::device_ptr<float>, float>, unsigned long>, thrust::cuda_cub::__uninitialized_fill::functor<thrust::device_ptr<float>, float>, unsigned long>(thrust::device_ptr<float>, float) [109]
214.56ms  5.8550us             (52 1 1)       (256 1 1)        29       44B        0B         -           -           -           -  GeForce GTX 970         1         7  void thrust::cuda_cub::cub::DeviceReduceKernel<thrust::cuda_cub::cub::DeviceReducePolicy<float, int, thrust::plus<float>>::Policy600, thrust::detail::normal_iterator<thrust::device_ptr<float>>, float*, int, thrust::plus<float>>(int, float, thrust::plus<float>, thrust::cuda_cub::cub::GridEvenShare<float>, thrust::cuda_cub::cub::DeviceReducePolicy<float, int, thrust::plus<float>>::Policy600) [128]
214.58ms  2.7200us              (1 1 1)       (256 1 1)        27       44B        0B         -           -           -           -  GeForce GTX 970         1         7  void thrust::cuda_cub::cub::DeviceReduceSingleTileKernel<thrust::cuda_cub::cub::DeviceReducePolicy<float, int, thrust::plus<float>>::Policy600, float*, thrust::detail::normal_iterator<thrust::pointer<float, thrust::cuda_cub::tag, thrust::use_default, thrust::use_default>>, int, thrust::plus<float>, float>(int, float, thrust::plus<float>, thrust::cuda_cub::cub::DeviceReducePolicy<float, int, thrust::plus<float>>::Policy600, float*) [136]
214.60ms  1.1840us                    -               -         -         -         -        4B  3.2219MB/s      Device    Pageable  GeForce GTX 970         1         7  [CUDA memcpy DtoH]
214.98ms  221.27us            (512 1 1)       (256 1 1)        20        0B        0B         -           -           -           -  GeForce GTX 970         1         7  void thrust::cuda_cub::core::_kernel_agent<thrust::cuda_cub::__parallel_for::ParallelForAgent<thrust::cuda_cub::__uninitialized_fill::functor<thrust::device_ptr<curandStateXORWOW>, curandStateXORWOW>, unsigned long>, thrust::cuda_cub::__uninitialized_fill::functor<thrust::device_ptr<curandStateXORWOW>, curandStateXORWOW>, unsigned long>(thrust::device_ptr<curandStateXORWOW>, curandStateXORWOW) [151]
219.51ms  69.492ms            (512 1 1)       (256 1 1)       108        0B        0B         -           -           -           -  GeForce GTX 970         1         7  void thrust::cuda_cub::core::_kernel_agent<thrust::cuda_cub::__parallel_for::ParallelForAgent<thrust::cuda_cub::for_each_f<thrust::zip_iterator<thrust::tuple<thrust::counting_iterator<int, thrust::use_default, thrust::use_default, thrust::use_default>, thrust::detail::normal_iterator<thrust::device_ptr<curandStateXORWOW>>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, thrust::detail::wrapped_function<curand_setup, void>>, int>, thrust::cuda_cub::for_each_f<thrust::zip_iterator<thrust::tuple<thrust::counting_iterator<int, thrust::use_default, thrust::use_default, thrust::use_default>, thrust::detail::normal_iterator<thrust::device_ptr<curandStateXORWOW>>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, thrust::detail::wrapped_function<curand_setup, void>>, int>(thrust::use_default, thrust::use_default) [160]
289.00ms  9.5360us             (52 1 1)       (256 1 1)        29       44B        0B         -           -           -           -  GeForce GTX 970         1         7  void thrust::cuda_cub::cub::DeviceReduceKernel<thrust::cuda_cub::cub::DeviceReducePolicy<float, int, thrust::plus<float>>::Policy600, thrust::detail::normal_iterator<thrust::device_ptr<float>>, float*, int, thrust::plus<float>>(int, float, thrust::plus<float>, thrust::cuda_cub::cub::GridEvenShare<float>, thrust::cuda_cub::cub::DeviceReducePolicy<float, int, thrust::plus<float>>::Policy600) [179]
289.01ms  3.4880us              (1 1 1)       (256 1 1)        27       44B        0B         -           -           -           -  GeForce GTX 970         1         7  void thrust::cuda_cub::cub::DeviceReduceSingleTileKernel<thrust::cuda_cub::cub::DeviceReducePolicy<float, int, thrust::plus<float>>::Policy600, float*, thrust::detail::normal_iterator<thrust::pointer<float, thrust::cuda_cub::tag, thrust::use_default, thrust::use_default>>, int, thrust::plus<float>, float>(int, float, thrust::plus<float>, thrust::cuda_cub::cub::DeviceReducePolicy<float, int, thrust::plus<float>>::Policy600, float*) [187]
289.07ms  1.3120us                    -               -         -         -         -        4B  2.9075MB/s      Device    Pageable  GeForce GTX 970         1         7  [CUDA memcpy DtoH]
289.66ms  9.9520us             (52 1 1)       (256 1 1)        29       44B        0B         -           -           -           -  GeForce GTX 970         1         7  void thrust::cuda_cub::cub::DeviceReduceKernel<thrust::cuda_cub::cub::DeviceReducePolicy<float, int, thrust::plus<float>>::Policy600, thrust::detail::normal_iterator<thrust::device_ptr<float>>, float*, int, thrust::plus<float>>(int, float, thrust::plus<float>, thrust::cuda_cub::cub::GridEvenShare<float>, thrust::cuda_cub::cub::DeviceReducePolicy<float, int, thrust::plus<float>>::Policy600) [211]
289.68ms  3.3280us              (1 1 1)       (256 1 1)        27       44B        0B         -           -           -           -  GeForce GTX 970         1         7  void thrust::cuda_cub::cub::DeviceReduceSingleTileKernel<thrust::cuda_cub::cub::DeviceReducePolicy<float, int, thrust::plus<float>>::Policy600, float*, thrust::detail::normal_iterator<thrust::pointer<float, thrust::cuda_cub::tag, thrust::use_default, thrust::use_default>>, int, thrust::plus<float>, float>(int, float, thrust::plus<float>, thrust::cuda_cub::cub::DeviceReducePolicy<float, int, thrust::plus<float>>::Policy600, float*) [219]
289.69ms  1.3120us                    -               -         -         -         -        4B  2.9075MB/s      Device    Pageable  GeForce GTX 970         1         7  [CUDA memcpy DtoH]

Здесь вы можете видеть, что три вызова, составляющие операцию сокращения, занимают в совокупности 8-13 микросекунд каждый, тогда как для for_each_n требуется 69 миллисекунд для завершения.

...