2014-10-11 8 views
0

그래서 나는 OpenCL을 위해이 사용자 정의 RNG 라이브러리의 사용을 만들려고 노력 해요 :가장 좋은 방법은

//! Represents the state of a particular generator 
typedef struct{ uint x; uint c; } mwc64x_state_t; 
: http://cas.ee.ic.ac.uk/people/dt10/research/rngs-gpu-mwc64x.html

도서관은 상태 구조체를 정의

그리고 임의의 UINT를 생성하기 위해 다음과 같은 기능으로 상태로 전달합니다 있도록 상태를 업데이트

uint MWC64X_NextUint(mwc64x_state_t *s) 

내가 통과 t를 다시 함수에 넣으면 시퀀스의 다음 "임의"숫자가 생성됩니다.

내가 만드는 프로젝트에 대해 다른 작업 그룹/항목뿐만 아니라 여러 장치에서 동시에 난수를 생성 할 수 있어야하며이를 설계하는 가장 좋은 방법을 찾는 데 문제가 있습니다. 마찬가지로 장치/명령 대기열 당 1 개의 mwc64x_state_t 객체를 만들어 그 상태를 전역 변수로 전달해야합니까? 또는 한 번에 모든 장치에 대해 하나의 상태 개체를 만들 수 있습니까? 아니면 전역 변수로 전달하지 않고 각 커널 함수 내에서 로컬로 새 상태를 선언합니까?

라이브러리는 또한이 기능을 함께 제공 : 아마도 여러에 "스트림"는 RNG를 분할하기로하지만 내 커널이 포함되어

void MWC64X_SeedStreams(mwc64x_state_t *s, ulong baseOffset, ulong perStreamOffset) 

그것이 믿을 수 없을만큼 느린 있습니다. 예를 들어 다음과 같은 간단한 작업을 수행 할 경우 :

__kernel void myKernel() 
{ 
    mwc64x_state_t rng; 
    MWC64X_SeedStreams(&rng, 0, 10000); 
} 

그러면 커널 호출 속도가 40 배 정도 느려집니다.

라이브러리에는 예제 코드로 제공되는 일부 소스 코드가 있지만 예제 코드는 제한적이며 도움이되지 않습니다.

누구나 openCL의 RNG에 익숙하거나이 특정 라이브러리를 사용한 적이 있다면 귀하의 조언에 진심으로 감사드립니다.

답변

3

MWC64X_SeedStreams 기능은 실제로 적어도 MWC64X_NextUint 통화에 비해 에서, 상대적으로 느리지 만,이 에서 사용할 수 많은 서브 스트림으로 큰 글로벌 스트림을 분할 을 시도 대부분의 병렬 된 RNG의 사실이다 평행. 가정에서는 NextUint를 커널 내에서 번 (예 : 100 개 이상)을 여러 번 호출 할 것이지만 SeedStream은 맨 위에 표시됩니다.

(mwc64x/test/estimate_pi.cpp 및 mwc64x/test/test_mwc64x.cl) 라이브러리와 함께 제공되는 EstimatePi 예제의 주석 버전입니다.

__kernel void EstimatePi(ulong n, ulong baseOffset, __global ulong *acc) 
{ 
    // One RNG state per work-item 
    mwc64x_state_t rng; 

    // This calculates the number of samples that each work-item uses 
    ulong samplesPerStream=n/get_global_size(0); 

    // And then skip each work-item to their part of the stream, which 
    // will from stream offset: 
    // baseOffset+2*samplesPerStream*get_global_id(0) 
    // up to (but not including): 
    // baseOffset+2*samplesPerStream*(get_global_id(0)+1) 
    // 
    MWC64X_SeedStreams(&rng, baseOffset, 2*samplesPerStream); 


    // Now use the numbers 
    uint count=0; 
    for(uint i=0;i<samplesPerStream;i++){ 
     ulong x=MWC64X_NextUint(&rng); 
     ulong y=MWC64X_NextUint(&rng); 
     ulong x2=x*x; 
     ulong y2=y*y; 
     if(x2+y2 >= x2) 
      count++; 
    } 
    acc[get_global_id(0)] = count; 
} 

그래서 의도는 n이 클 수 및 samplesPerStream 주위 백 이상 유지되도록 작업 항목의 성장 수가 증가해야한다는 것입니다.

당신이 다음 필요 스트림 분할에 계층 구조의 다른 수준을 추가하기 위해 여러 장치에 여러 개의 커널을 원하는 경우에, 그래서 예를 들어, 당신은있는 경우 :

  • K : 장치의 수 (아마도에를 병렬 컴퓨터)
  • W : 디바이스 당 번호 작업 항목
  • C : 호출 번호 작업 항목 당 NextUint하는

Y o = K W 모든 C 전체작업 항목에 대해 NextUint를 호출합니다. 당신의 장치로 식별하는 경우 K = 0 .. (K-1), 다음 각 커널 내에서 당신이 할 것 :

MWC64X_SeedStreams(&rng, W*C*k, C); 

그런 다음 스트림 내에서 인덱스는 다음과 같습니다

[0    .. N) : Parts of stream used across all devices 
[k*(W*C)  .. (k+1)*(W*C)) : Used within device k 
[k*(W*C)+(i*C) .. (k*W*C)+(i+1)*C) : Used by work-item i in device k. 

에게 그것을 각 커널이 C 샘플보다 적은 수를 사용한다면 괜찮습니다. 필요하다면 을 과소 평가할 수 있습니다.

(저는 도서관의 저자입니다).

+0

정보 및 설명을 제공해 주셔서 감사합니다. 불행히도 내 커널 작업 항목은 NextUint()에 대해 각각 약 12 ​​개 정도의 호출 만하면되므로 SeedStreams() 호출의 큰 오버 헤드를 보지 못하는 것처럼 보입니다. – user1855952