CUDA의 차이 방정식으로 설명 된 지수 이동 평균 필터 구현

Question

현재 CUDA를 사용하여 지수 이동 평균 필터을 나타내는 차등 방정식을 계산하려고합니다. alpha 및 beta이 상수는 상기 차분 방정식이 필터에 대한 시스템 응답을 얻기 위해 조작 할 수있는 방법

alpha = (2.0/(1 + Period)) 
beta = 1 - alpha 

로 정의되는 경우, 필터는 다음과 같은 차분 방정식

y[n] = y[n-1] * beta + alpha * x[n] 

하여 설명? GPU에이 필터를 구현하는 효율적인 방법은 무엇입니까?

나는 GTX CUDA 추력 프리미티브를 사용하여 다음 아래 표시와 나는 위의 차이 방정식을 조작하는 제안 것 (570)

Answer 1

작업입니다. -

차이 식 조작 차이 식 간단한 대수으로

명시 적으로 FORM, 당신은 다음 찾을 수 있습니다

y[n]/beta^n = y[0] + alpha * x[1]/beta + alpha * x[2]/beta^2 + ... 

: 따라서

y[1] = beta * y[0] + alpha * x[1] 

y[2] = beta^2 * y[0] + alpha * beta * x[1] + alpha * x[2] 

y[3] = beta^3 * y[0] + alpha * beta^2 * x[1] + alpha * beta * x[2] + alpha * x[3] 

을 명시 적 형태는 다음과 같다

쿠다 쓰루풋 구현

당신은 다음 단계로 위의 명시적인 형태로 구현할 수 있습니다

1. 이 d_input[0] = 1.를 제외하고 alpha에 입력 순서 d_input 초기화를;
2. 벡터 d_1_over_beta_to_the_n을 1, 1/beta, 1/beta^2, 1/beta^3, ...과 동일하게 정의하십시오.
3. 요소 단위로 곱하기 d_inputd_1_over_beta_to_the_n;
4. y[n]/beta^n의 시퀀스를 얻기 위해 inclusive_scan을 수행하십시오.
5. 위의 시퀀스를 1, 1/beta, 1/beta^2, 1/beta^3, ...으로 나눕니다.

편집은

위의 방법 선형 시변 (LTV) 시스템을 권장 할 수있다. Linear Time-Invariant (LTI) 시스템의 경우 Paul이 언급 한 FFT 접근 방식을 권장 할 수 있습니다. FIR filter in CUDA에 CUDA Thrust 및 cuFFT를 사용하여 이러한 접근 방식의 예를 제공합니다.

다른 접근 방식에 대한 전체 코드

#include <thrust/sequence.h> 

#include <thrust/device_vector.h> 
#include <thrust/host_vector.h> 

int main(void) 
{ 
    int N = 20; 

    // --- Filter parameters 
    double alpha = 2.7; 
    double beta  = -0.3; 

    // --- Defining and initializing the input vector on the device 
    thrust::device_vector<double> d_input(N,alpha * 1.); 
    d_input[0] = d_input[0]/alpha; 

    // --- Defining the output vector on the device 
    thrust::device_vector<double> d_output(d_input); 

    // --- Defining the {1/beta^n} sequence 
    thrust::device_vector<double> d_1_over_beta(N,1./beta); 
    thrust::device_vector<double> d_1_over_beta_to_the_n(N,1./beta); 
    thrust::device_vector<double> d_n(N); 
    thrust::sequence(d_n.begin(), d_n.end()); 
    thrust::inclusive_scan(d_1_over_beta.begin(), d_1_over_beta.end(), d_1_over_beta_to_the_n.begin(), thrust::multiplies<double>()); 
    thrust::transform(d_1_over_beta_to_the_n.begin(), d_1_over_beta_to_the_n.end(), d_input.begin(), d_input.begin(), thrust::multiplies<double>());  
    thrust::inclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), thrust::plus<double>()); 
    thrust::transform(d_output.begin(), d_output.end(), d_1_over_beta_to_the_n.begin(), d_output.begin(), thrust::divides<double>()); 

    for (int i=0; i<N; i++) { 
     double val = d_output[i]; 
     printf("Device vector element number %i equal to %f\n",i,val); 
    } 

    // --- Defining and initializing the input vector on the host 
    thrust::host_vector<double> h_input(N,1.); 

    // --- Defining the output vector on the host 
    thrust::host_vector<double> h_output(h_input); 

    h_output[0] = h_input[0]; 
    for(int i=1; i<N; i++) 
    { 
     h_output[i] = h_input[i] * alpha + beta * h_output[i-1]; 
    } 

    for (int i=0; i<N; i++) { 
     double val = h_output[i]; 
     printf("Host vector element number %i equal to %f\n",i,val); 
    } 

    for (int i=0; i<N; i++) { 
     double val = h_output[i] - d_output[i]; 
     printf("Difference between host and device vector element number %i equal to %f\n",i,val); 
    } 

    getchar(); 
} 

Answer 2

, 당신은 지수 이동 평균 창을 절단 한 다음 신호와 창 지수 사이의 회선을 수행하여 필터링 된 신호를 계산할 수 있습니다.convolution은 자유 CUDA FFT 라이브러리 (cuFFT)를 사용하여 계산할 수 있습니다. 왜냐하면 알고있는 것처럼 컨볼 루션은 푸리에 도메인에서 두 신호의 점진적 곱셈으로 표현할 수 있기 때문입니다 (이것은 Convolution Theorem이라는 적절한 이름입니다, 이것은 O (n log (n))의 복잡성으로 실행됩니다. 이러한 유형의 접근 방식은 CUDA 커널 코드를 최소화하고 GeForce 570에서도 매우 빠르게 실행됩니다. 특히, 모든 계산을 단일 (부동) 정밀도로 수행 할 수 있다면 그렇습니다.