2012-09-12 4 views
1

나는 각각 8000 개의 요소가있는 6 개의 벡터의 평균 및 표준 편차를 찾기위한 코드를 작성하고 있습니다. CUDA를 사용하여 작업을 수행 할 수 있는지 궁금해서 작업 속도를 높였습니다. CUDA를 사용하여 평균을 찾는 방법을 생각할 수는 있지만 CUDA를 사용하여 표준 편차를 계산하는 방법을 이해할 수 없습니다. 아무도 나 좀 도와 줄래?CUDA를 사용한 표준 편차

+0

몇 부동 소수점 연산은 적어도 수백 배 너무 적은 속도로 CUDA를 보람있게 사용하십시오. 많은 양의 데이터가있는 GPU는 단순히 너무 적습니다. 데이터를 GPU로 이동하고 계산을 초기화 할 때의 대기 시간은 지난 10 년 동안 구축 된 모든 호스트 CPU의 스프레드 시트 *에서 계산을 수행하는 데 걸리는 시간보다 커집니다. 그 문제의 크기에서, 당신은 이길 수 없습니다. – talonmies

+0

이것은 아주 작은 부분 집합입니다 ... 실제로 처리 할 2 천만 X 6 요소가 있습니다 – user1439690

+1

왜 그런 질문을하지 않았습니까? – talonmies

답변

0

데이터 마이닝을 위해 CUDA에서이 문제를 해결했습니다. 나는 모든 도서관을 사용하지 않았다. 그러나, 그것은 나에게 좋은 결과를 주었다. 문제는 표준 편차와 128 * 100 만 샘플의 평균을 찾는 것입니다. 이것이 내가 한 일이다.

  1. 내 장치의 공유 메모리는 16KB입니다. 그리고 저는 수레를 사용하고 있습니다. 따라서 공유 메모리는 4,000 개의 요소를 수용 할 수 있습니다. 내 장치에 대한 블록 당 최대 스레드는 512입니다. 따라서 8 블록을 가질 수 있습니다. 16KB를 8 개의 블록으로 나누면 2,000KB가됩니다 (즉, 1 개의 스레드에 대해 1 개의 플로트). 일반적으로 일치하지 않습니다. 더 나은 장치를 가지고 있다면이 수학을 다시해야합니다.

  2. 표준 편차를 찾기 위해 각 블록에는 512 개의 요소가 있습니다. 단일 스레드를 사용하여 사각형 (요소 평균)을 찾을 수 있습니다.

  3. 다음 과제는이를 추가하여 이러한 요소의 합계를 찾는 것입니다. 평균을 찾는 데 사용한 것과 같은 방법을 시도하십시오. 512 요소. 결과를 전역 메모리로 복사하십시오.

  4. 반복. 결과의 제곱근을 찾습니다.

추신 : 글로벌 메모리 호출이 최소가되도록 계획하십시오. 평균 및 표준 편차는 데이터를 메모리에서 자주 호출합니다.

+0

솔루션을 주셔서 감사합니다 ... 나는 또한 내 글로벌 메모리 호출을 최소화하려고 할 것입니다 ... – user1439690

2

이것은 내 전문 분야 밖에 있지만 표준 편차를 계산하기위한 단일 패스 반복 알고리즘이 있습니다.이 알고리즘은 축소로 변환 될 수 있습니다. 특히 Knuth, TAOCP, vol. 1에서 설명한대로 Welford의 알고리즘을 생각하고 있습니다. 2. 하나의 단점은 모든 단계에서 나눗셈이 필요하지만, 이것은 필요한 메모리 액세스와 잘 균형을 이룰 것입니다. 알고리즘에 대한 가능한 온라인 참조로 나타납니다 평균과 표준을 포함하여 단일 패스에서 요약 통계의 수를 계산

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

6

Here is a Thrust example합니다. 일탈.

#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; 
} 
0

늦은 대답은,하지만 난 내 코드에 thrust::transform_reduce를 사용하여이 문제를 해결 한 (GTX는 1,070에서 100,000 수레 테스트) : 48000 개 번호에

#include <thrust/transform_reduce.h> 
#include <thrust/device_vector.h> 
#include <thrust/functional.h> 

#include <functional> 
#include <cmath> 

/* 
* @struct varianceshifteop 
* @brief a unary function that shifts input data 
* by their mean and computes the squares of them 
*/ 
struct varianceshifteop 
    : std::unary_function<float, float> 
{ 
    varianceshifteop(float m) 
     : mean(m) 
    { /* no-op */ } 

    const float mean; 

    __device__ float operator()(float data) const 
    { 
     return ::pow(data - mean, 2.0f); 
    } 
}; 

int main(int argc, char** argv) 
{ 
    thrust::device_vector<float> data{ ... }; 

    // sum elements and divide by the number of elements 
    float mean = thrust::reduce(
     data.cbegin(), 
     data.cend(), 
     0.0f, 
     thrust::plus<float>())/data.size(); 

    // shift elements by mean, square, and add them 
    float variance = thrust::transform_reduce(
      data.cbegin(), 
      data.cend(), 
      varianceshifteop(mean), 
      0.0f, 
      thrust::plus<float>())/(data.size() - 1); 

    // standard dev is just a sqrt away 
    float stdv = std::sqrtf(variance); 

    return 0; 
}