1

それぞれ8000要素を持つ6つのベクトルの平均と標準偏差を見つけるコードを書いています。CUDAを使ってそれを実行し、動作を高速化できないかと考えていました。CUDAを使用して平均を見つける方法を考えることができましたが、CUDAを使用して標準偏差を計算する方法を理解できません。ここで誰か助けてくれませんか？

4

5 に答える 5

6

これは、平均と標準を含む、1 回のパスで多数の要約統計を計算するThrust の例です。偏差。

#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/transform_reduce.h>
#include <thrust/functional.h>
#include <thrust/extrema.h>
#include <cmath>
#include <limits>

// This example computes several statistical properties of a data
// series in a single reduction.  The algorithm is described in detail here:
// http://en.wikipedia.org/wiki/Algorithms_for_calculating_variance#Parallel_algorithm
//
// Thanks to Joseph Rhoads for contributing this example


// structure used to accumulate the moments and other 
// statistical properties encountered so far.
template <typename T>
struct summary_stats_data
{
    T n;
    T min;
    T max;
    T mean;
    T M2;
    T M3;
    T M4;

    // initialize to the identity element
    void initialize()
    {
      n = mean = M2 = M3 = M4 = 0;
      min = std::numeric_limits<T>::max();
      max = std::numeric_limits<T>::min();
    }

    T variance()   { return M2 / (n - 1); }
    T variance_n() { return M2 / n; }
    T skewness()   { return std::sqrt(n) * M3 / std::pow(M2, (T) 1.5); }
    T kurtosis()   { return n * M4 / (M2 * M2); }
};

// stats_unary_op is a functor that takes in a value x and
// returns a variace_data whose mean value is initialized to x.
template <typename T>
struct summary_stats_unary_op
{
    __host__ __device__
    summary_stats_data<T> operator()(const T& x) const
    {
         summary_stats_data<T> result;
         result.n    = 1;
         result.min  = x;
         result.max  = x;
         result.mean = x;
         result.M2   = 0;
         result.M3   = 0;
         result.M4   = 0;

         return result;
    }
};

// summary_stats_binary_op is a functor that accepts two summary_stats_data 
// structs and returns a new summary_stats_data which are an
// approximation to the summary_stats for 
// all values that have been agregated so far
template <typename T>
struct summary_stats_binary_op 
    : public thrust::binary_function<const summary_stats_data<T>&, 
                                     const summary_stats_data<T>&,
                                           summary_stats_data<T> >
{
    __host__ __device__
    summary_stats_data<T> operator()(const summary_stats_data<T>& x, const summary_stats_data <T>& y) const
    {
        summary_stats_data<T> result;

        // precompute some common subexpressions
        T n  = x.n + y.n;
        T n2 = n  * n;
        T n3 = n2 * n;

        T delta  = y.mean - x.mean;
        T delta2 = delta  * delta;
        T delta3 = delta2 * delta;
        T delta4 = delta3 * delta;

        //Basic number of samples (n), min, and max
        result.n   = n;
        result.min = thrust::min(x.min, y.min);
        result.max = thrust::max(x.max, y.max);

        result.mean = x.mean + delta * y.n / n;

        result.M2  = x.M2 + y.M2;
        result.M2 += delta2 * x.n * y.n / n;

        result.M3  = x.M3 + y.M3;
        result.M3 += delta3 * x.n * y.n * (x.n - y.n) / n2; 
        result.M3 += (T) 3.0 * delta * (x.n * y.M2 - y.n * x.M2) / n;

        result.M4  = x.M4 + y.M4;
        result.M4 += delta4 * x.n * y.n * (x.n * x.n - x.n * y.n + y.n * y.n) / n3;
        result.M4 += (T) 6.0 * delta2 * (x.n * x.n * y.M2 + y.n * y.n * x.M2) / n2;
        result.M4 += (T) 4.0 * delta * (x.n * y.M3 - y.n * x.M3) / n;

        return result;
    }
};

template <typename Iterator>
void print_range(const std::string& name, Iterator first, Iterator last)
{
    typedef typename std::iterator_traits<Iterator>::value_type T;

    std::cout << name << ": ";
    thrust::copy(first, last, std::ostream_iterator<T>(std::cout, " "));  
    std::cout << "\n";
}


int main(void)
{
    typedef float T;

    // initialize host array
    T h_x[] = {4, 7, 13, 16};

    // transfer to device
    thrust::device_vector<T> d_x(h_x, h_x + sizeof(h_x) / sizeof(T));

    // setup arguments
    summary_stats_unary_op<T>  unary_op;
    summary_stats_binary_op<T> binary_op;
    summary_stats_data<T>      init;

    init.initialize();

    // compute summary statistics
    summary_stats_data<T> result = thrust::transform_reduce(d_x.begin(), d_x.end(), unary_op, init, binary_op);

    std::cout <<"******Summary Statistics Example*****"<<std::endl;
    print_range("The data", d_x.begin(), d_x.end());

    std::cout <<"Count              : "<< result.n << std::endl;
    std::cout <<"Minimum            : "<< result.min <<std::endl;
    std::cout <<"Maximum            : "<< result.max <<std::endl;
    std::cout <<"Mean               : "<< result.mean << std::endl;
    std::cout <<"Variance           : "<< result.variance() << std::endl;
    std::cout <<"Standard Deviation : "<< std::sqrt(result.variance_n()) << std::endl;
    std::cout <<"Skewness           : "<< result.skewness() << std::endl;
    std::cout <<"Kurtosis           : "<< result.kurtosis() << std::endl;

    return 0;
}
于 2012-09-12T06:14:03.467 に答える
2

これは私の専門分野ではありませんが、標準偏差を計算するための単一パスの反復アルゴリズムがあり、これを削減に変換できる場合があります。特に、クヌース、TAOCP、vol. 2. 1 つの欠点は、すべてのステップで除算が必要なことですが、これは必要なメモリ アクセスとのバランスが取れている可能性があります。アルゴリズムの有用なオンライン リファレンスは次のようです。

http://www.johndcook.com/standard_deviation.html

于 2012-09-12T05:09:22.813 に答える
0

データマイニングのためにCUDAでこの問題を解決しました。私はライブラリを使用していません。しかし、それは私に良い結果をもたらしました。問題は、128*100 万サンプルの標準偏差と平均を見つけることです。これが私がしたことです。

  1. 私のデバイスには 16KB の共有メモリがあります。そして、フロートを使用しています。したがって、共有メモリは 4,000 要素に対応できます。私のデバイスのブロックあたりの最大スレッドは 512 です。したがって、8 つのブロックを持つことができます。16KB を 8 ブロックに分割すると、2,000KB になります (つまり、1 スレッドに対して 1 float)。通常、これは一致しません。より良いデバイスがある場合は、この計算をもう一度行う必要があります。

  2. 標準偏差を求めるために、各ブロックには 512 個の要素があります。シングルスレッドを使用して平方（要素平均）を見つけることができます。

  3. 次の課題は、これを追加して、これらの要素の合計を見つけることです。平均を見つけるために使用したのと同じ方法を試してください。512エレメント用。結果をグローバル メモリにコピーします。

  4. 繰り返します。結果の平方根を求めます。

PS: それに応じて、グローバル メモリの呼び出しが最小限になるように計画してください。平均と標準偏差はメモリからデータを頻繁に呼び出します。

于 2012-09-12T14:48:23.040 に答える
0

Thrustを使用してベクトルの平均と分散を求めるソリューションは次のとおりです。

template <typename T> struct square {
    __host__ __device__ T operator()(const T& x) const {
        return x * x;
    }
};

template <typename T> void mean_and_var(T a, int n, double* p_mean, double* p_var) {
    double sum = thrust::reduce(a, &a[n], 0.0, thrust::plus<double>());
    double sum_square = thrust::transform_reduce(
        a,
        &a[n],
        square<double>(),
        0.0,
        thrust::plus<double>()
    );
    double mean = sum / n;
    *p_mean = mean;
    *p_var = (sum_square / n) - mean*mean;
}

CPU と GPU を比較した自己完結型のソース ファイルのソリューションを次に示します (質問に正確に答えるにはmean_and_var、長さ 8000 の 6 つの異なるベクトルで 6 回呼び出す必要がありますが、これで要点がわかります）：

#include "stdio.h"
#include <thrust/reduce.h>
#include <thrust/device_vector.h>

#define PROFILING_INIT                                              \
    cudaEvent_t start, stop;                                        \
    float elapsedTime;

#define PROFILING_START                                             \
    cudaEventCreate(&start);                                        \
    cudaEventCreate(&stop);                                         \
    cudaEventRecord(start, 0);

#define PROFILING_STOP                                              \
    cudaEventRecord(stop, 0);                                       \
    cudaEventSynchronize(stop);                                     \
    cudaEventElapsedTime(&elapsedTime, start, stop);                \
    printf("Time elapsed:  %.3g ms\n", elapsedTime);

#define N (6*8000)
// #define N (6*8000*10)
// #define N (6*8000*100)
double a[N];

template <typename T> struct square {
    __host__ __device__ T operator()(const T& x) const {
        return x * x;
    }
};

void mean_and_var_cpu(double* a, int n, double* p_mean, double* p_var) {
    double sum = 0, sum_square = 0, mean;
    for (int i = 0; i < n; i++) {
        sum += a[i];
        sum_square += (a[i] * a[i]);
    }
    mean = sum / n;
    *p_mean = mean;
    *p_var = (sum_square / n) - mean*mean;
}

template <typename T> void mean_and_var(T a, int n, double* p_mean, double* p_var) {
    double sum = thrust::reduce(a, &a[n], 0.0, thrust::plus<double>());
    double sum_square = thrust::transform_reduce(a, &a[n], square<double>(), 0.0, thrust::plus<double>());
    double mean = sum / n;
    *p_mean = mean;
    *p_var = (sum_square / n) - mean*mean;
}

int main() {
    for (int i = 0; i < N; i++) {
        a[i] = i;
    }

    double mean, var;

    PROFILING_INIT;

    printf("With thrust:\n");
    PROFILING_START;
    mean_and_var<double*>(a, N, &mean, &var);
    PROFILING_STOP;
    printf("Mean = %f, var = %f\n", mean, var);

    printf("With thrust, using device memory:\n");
    thrust::device_vector<double> a_dev(N);
    thrust::copy(a, &a[N], a_dev.begin());
    PROFILING_START;
    mean_and_var<thrust::device_ptr<double>>(&a_dev[0], N, &mean, &var);
    PROFILING_STOP;
    printf("Mean = %f, var = %f\n", mean, var);

    printf("On CPU:\n");
    PROFILING_START;
    mean_and_var_cpu(a, N, &mean, &var);
    PROFILING_STOP;
    printf("Mean = %f, var = %f\n", mean, var);

}

コードの品質 (私は主に C プログラマーで、Thrust の目的で C++ に適応しようとしています) とコメントの欠如についてお詫びします。

statisticsすべての場合において、答えは Pythonモジュールを使用して計算された平均と母分散に一致します。

import statistics
x = 8000*6
print(statistics.mean(list(range(x))))
# print(statistics.variance(list(range(x))))
print(statistics.pvariance(list(range(x))))

以下は、Jetson Nano 開発ボードで実行されたプロファイリング情報です。

N 推力にかかった時間 (ms) デバイス メモリを使用して Thrust にかかった時間 (ミリ秒) 単純な CPU の実装にかかった時間 (ミリ秒)
6*8000 11.2 2.38 1.11
6*8000*10 108 7.62 11.7
6*8000*100 1.12e+03 41.7 109

プロファイリングからの結論:

  • Thrust を使用する場合は、必ずデバイス メモリを使用してください。
  • 入力サイズが十分に大きい場合、GPU は CPU よりも高速です。
  • 小さい入力サイズの場合、プラットフォームによっては、GPU よりも CPU を使用した方が速い場合があります (小さい入力サイズで高速化するために私の Thrust コードを改善する理由/方法を誰かが教えてくれれば、それは素晴らしいことです! )
于 2022-03-02T18:12:50.207 に答える