2014-09-22 7 views
1

GPU 공유 메모리에 큰 배열을로드하고 싶습니다. 내가 벨로우즈처럼 사용할 때 :GPU 공유 메모리에 큰 배열을 효율적으로로드하는 방법은 무엇입니까?

1 : int index = threadidx.x;

2 : 공유 unsigned char x [1000];

3 : x [i] = array [i];

그런 다음 1000 개의 스레드와 하나의 블록으로 커널 코드를 호출하면 모든 스레드에 대해 메모리 액세스가 발생합니까? 이 배열을 하나의 메모리 액세스로로드하여 공유 메모리에로드 할 수 있습니까?

모든 의견을 크게 기뻐할 것입니다.

답변

2

아니요 단일 액세스로 수행 할 수 없습니다.

공유 메모리로드와 병렬로 스레드를 사용하면 앞에서 설명한 것처럼 가장 빠른 방법입니다. 공유 메모리는 만 CUDA 커널의 스레드가 수행하는 메모리 연산에 의해로드 될 수 있습니다. 공유 메모리를로드하는 API 함수는 없습니다.

스레드 블록의 스레드 수보다 큰 배열이있는 경우 here과 같은 루핑 방식을 사용할 수 있습니다.

+0

이렇게하면 공유 메모리에 배열 요소를 동시에로드 할 수 있습니까? – kiki

+0

"평행" "한 번에"활동 만 워프 레벨에 있습니다. 각 워프 (32 스레드)는 전역 메모리 ('array [i]')에 대한 액세스를 단일 액세스/트랜잭션으로 결합하고 공유 메모리 ('x [i]')에 대한 액세스를 단일 액세스로 결합합니다. 다중 워프가 병렬로 실행될 수도 있지만, 해당 활동은 단일 메모리 트랜잭션으로 결합되지 않습니다. –

+0

일반적으로 스레드 당 최소 4 바이트를로드하여로드/저장 효율성을 향상시킬 수 있습니다. 기본 데이터 구조가 바이트 지향적 일지라도 적절한 포인터 연산을 사용하여 기본 포인터 캐스팅 (예 :'char *'에서'int *')으로 수행 할 수 있습니다. –

0

일반적으로 패러다임은 블록에 대해 for 루프를 사용합니다. 공유 메모리는 블록간에 만 공유되므로 공유 메모리가 블록 수보다 클 경우 멋진 루프가 필요할 것입니다.

커널 시작 부분에 다음과 같이하십시오.

//loadbuff is a T*, T is whatever type you want 
//sharedmemsize is some (compile time) constant 
__shared__ T sharedmem[sharedmemsize]; 
int index = threadidx.x; 
int blocksize = blockDim.x; 

for (int i = index; i < sharedmemsize; i+=blocksize) 
{ 
    sharedmem[i]=loadbuff[i]; 
} 
__syncthreads(); 
+0

'blocksize [i]'는 코드에서 이해가되지 않습니다. 스레드 블록과 데이터 배열 모두에 대해 임의의 크기를 처리하는 접근법은 [여기] (http://stackoverflow.com/questions/24419822/efficiently-initializing-shared-memory-array-in-cuda/24419969#24419969) –

+0

@ RobertCrovella, 미안 해요, 오타였습니다. – IdeaHat

+0

@RobertCrovella 또한 링크의 메소드는 내 메소드와 동등하지만, C 구문 만 사용합니다.당신이 그런 종류의 일을하고 있다면 괜찮습니다. – IdeaHat