2011-04-06 2 views
3

오버랩이있을 수있는 단일 입력 배열에서 블록 단위 공유 메모리 영역을 할당하는 쉬운 방법이 있습니까?CUDA : 테두리가 겹치는 공유 메모리 할당

간단한 예제는 문자열 검색입니다. 보았습니다. 입력 텍스트를 다듬어보고 각 블록의 각 스레드가 text [thread_id]에서 시작하는 패턴을 검색하도록하고 각 블록에 할당 된 데이터를 패턴 길이만큼 겹쳐서 경계를 넘는 일치하는 사례를 아직도 발견. 각 블록에 공유 메모리에 할당 된 총 메모리 크기 즉

은 아마 간단한 뭔가를 놓치고과 CUDA 가이드를 통해 현재 다이빙 생각 해요

(blocksize+patternlength)*sizeof(char) 

이지만, 몇 가지 지침을 부탁드립니다.

업데이트 : 일부 사람들은 내 질문에 대해 오해하고 있거나 (내가 잘못 설명 했음) 의심 스럽습니다.

데이터 세트 QWERTYUIOP이 있다고 가정 해 봅시다. 3 문자 일치를 검색하려고합니다. 각 스레드 블록에 대해 데이터 세트를 임의로 4로 잘라냅니다. QWER TYUI OPxx

이것은 수행하기에 충분히 간단하지만 실제로 3 문자가 일치하면 IOP를 찾는 알고리즘이 실패합니다. 이 경우

가, 내가 원하는 것은 각 블록은 공유 메모리에 가지고입니다 :

QWERTY TYUIOP OPxxxx

즉, 각 블록은 그래서 메모리 국경 문제가 발생하지 않는 블록 크기 + patternlength-1 문자를 할당됩니다 .

더 나은 것을 설명하는 희망.

@jmilloy은 ... 지속되고 있기 때문에 : 대신 나는 짓을하고 싶은 무엇

//VERSION 1: Simple 
__global__ void gpuSearchSimple(char *T, int lenT, char *P, int lenP, int *pFound) 
{ 
    int startIndex = blockDim.x*blockIdx.x + threadIdx.x; 
    int fMatch = 1; 
    for (int i=0; i < lenP; i++) 
    { 
     if (T[startIndex+i] != P[i]) fMatch = 0; 
    } 
    if (fMatch) atomicMin(pFound, startIndex); 
} 
//VERSION 2: Texture 
__global__ void gpuSearchTexture(int lenT, int lenP, int *pFound) 
{ 
    int startIndex = blockDim.x*blockIdx.x + threadIdx.x; 
    int fMatch = 1; 
    for (int i=0; i < lenP; i++) 
    { 
     if (tex1Dfetch(texT,startIndex+i) != tex1Dfetch(texP,i)) fMatch = 0; 
    } 
    if (fMatch) atomicMin(pFound, startIndex); 
} 
//Version 3: Shared 
__global__ void gpuSearchTexSha(int lenT, int lenP, int *pFound) 
{ 
    extern __shared__ char shaP[]; 
    for (int i=0;threadIdx.x+i<lenP; i+=blockDim.x){ 
    shaP[threadIdx.x+i]= tex1Dfetch(texP,threadIdx.x+i); 
    } 
    __syncthreads(); 

    //At this point shaP is populated with the pattern 
    int startIndex = blockDim.x*blockIdx.x + threadIdx.x; 
    // only continue if an earlier instance hasn't already been found 
    int fMatch = 1; 
    for (int i=0; i < lenP; i++) 
    { 
     if (tex1Dfetch(texT,startIndex+i) != shaP[i]) fMatch = 0; 
    } 
    if (fMatch) atomicMin(pFound, startIndex); 
} 

질문의 나머지 부분에서 설명 된 바와 같이, 공유 메모리 덩어리로 텍스트를 넣어하는 것입니다 P, 이후 버전의 텍스쳐 메모리에 텍스트를 유지합니다.

+0

hmmm 이것은 내가 당신이 의미한다고 생각했던 것입니다. 당신이 내 대답에 대해 당신에게 잘못한 것이 무엇인지 설명한다면 도움이 될 것입니다. – jmilloy

+0

또한 두 가지 질문이 있습니다. '블록화'란 무엇입니까? 'text'는 무엇입니까 ('text [thread_id]'에서와 같이)? – jmilloy

답변

1

아니요. 공유 메모리는 블록의 스레드간에 공유되며 할당 된 블록에서만 액세스 할 수 있습니다. 두 개의 다른 블록에서 사용할 수있는 공유 메모리를 가질 수 없습니다.

내가 아는 한 공유 메모리는 실제로 다중 프로세서에 상주하며 스레드는 실행중인 다중 프로세서의 공유 메모리에만 액세스 할 수 있습니다. 그래서 이것은 물리적 한계입니다. (두 블록이 하나의 mp에 상주하는 경우 한 블록의 스레드가 다른 블록에 할당 된 공유 메모리에 예기치 않게 액세스 할 수 있습니다).

전역 메모리에서 공유 메모리로 데이터를 명시 적으로 복사해야한다는 점에 유의하십시오. 문자열의 중첩 영역을 겹치지 않는 공유 메모리에 복사하는 것은 간단합니다.

CUDA 프로그램을 개발하는 데 필요한 대부분의 작업이 필요하다고 생각합니다. 나의 지침은 먼저 공유 메모리를 사용하지 않고 문제를 해결하는 버전부터 시작하는 것입니다. 이를 위해 중복 된 문제를 해결하고 공유 메모리 구현을 쉽게 할 수 있습니다! 대답은 올바른

__global__ void gpuSearchTexSha(int lenT, int lenP, int *pFound) 
{ 
    extern __shared__ char* shared; 

    char* shaP = &shared[0]; 
    char* shaT = &shared[lenP]; 

    //copy pattern into shaP in parallel 
    if(threadIdx.x < lenP) 
     shaP[threadIdx.x] = tex1Dfetch(texP,threadIdx.x); 

    //determine texT start and length for this block 
    blockStartIndex = blockIdx.x * gridDim.x/lenT; 
    lenS = gridDim.x/lenT + lenP - 1; 

    //copy text into shaT in parallel 
    shaT[threadIdx.x] = tex1Dfetch(texT,blockStartIndex + threadIdx.x); 
    if(threadIdx.x < lenP) 
     shaP[blockDim.x + threadIdx.x] = text1Dfetch(texT,blockStartIndex + blockDim.x + threadIdx.x) 

    __syncthreads(); 

    //We have one pattern in shaP for each thread in the block 
    //We have the necessary portion of the text (with overlaps) in shaT 

    int fMatch = 1; 
    for (int i=0; i < lenP; i++) 
    { 
     if (shaT[threadIdx.x+i] != shaP[i]) fMatch = 0; 
    } 
    if (fMatch) atomicMin(pFound, blockStartIndex + threadIdx.x); 
} 

키 노트로 표시되었다


편집이
후 :

    우리는 블록 당 공유 메모리에 패턴의 하나의 사본이 필요
  • - 그들은 모두 사용할 수 있습니다.
  • 블록 당 필요한 공유 메모리는입니다.는
  • 커널 가정 (렌즈는 blocksize + patternlength가 어디) 그 gridDim.x * blockDim.x =
  • 우리가 루프에 대한 위해 (병렬 공유 메모리에 필요를 복사 할 수 있습니다 (버전 1과 동일) 사순절 당신이 경우 충분한 스레드를 가지고)
+0

최근에 편집 한 내용이 더 잘 설명되기를 바랍니다. 내가 원하는 것을 할 수있는 'memcpy'플래그 또는 파생물을 기대했지만, 결국 부분적인 해결책을 찾았습니다. 어쨌든 마크가 있습니다. – Bolster

+1

@ andrew + bolster 내 코드를보고 당신이 생각하는 것을보십시오. 당신은 내가 당신과 다른 방식으로 공유 메모리에 대해 생각하고 있음을 알게 될 것입니다. – jmilloy

1

나는 그 질문이 모두 그렇게 많은 의미가 있는지 확신하지 못한다.

__global__ void kernel() 
{ 
    extern __shared__ int buffer[]; 
    .... 
} 

kernel<<< gridsize, blocksize, buffersize >>>(); 

하지만, 버퍼의 내용은 커널의 시작 부분에 정의되어 있지 : 당신은 동적으로 다음과 같이 런타임에 공유 할당 메모리의 크기를 할 수 있습니다. 패턴 일치가 원하는대로 작동하도록 전역 메모리에서로드 할 겹침 부분을로드하려면 커널에서 스키마를 고안해야합니다.

+0

나는 피하려고 노력하고 있습니다. :)하지만 감사합니다. – Bolster

0

겹치는 공유 메모리가 좋지 않으므로 공유 메모리의 동일한 주소에 액세스 할 때마다 스레드가 동기화해야합니다 (아키텍처> 2.0에서는 완화 됨).

내 마음 속에서 생각해 볼 수있는 가장 간단한 아이디어는 겹칠 텍스트 부분을 복제하는 것입니다.

AAAA BBBB CCCC CCCC DDDD EEEEE : 중복으로 읽기 EEEE

AAAA BBBB CCCC DDDD 대신 정확한 chuncks의 글로벌 메모리에서 읽기의