CUDA 공유 메모리 할당에 대해 상수 선언 위치

Question

1024 행렬에 대한 적합도 함수를 실행 중입니다. 각 행렬은 자체 블록을 가지며 같은 크기입니다. 각 블록에는 n*n 개의 스레드 (행렬의 차원)가 있으며 n*n 공유 메모리가 있어야 손쉽게 총계 축소를 수행 할 수 있습니다. 그러나 모든 행렬에 대한 n 차원은 런타임 전에 가변적입니다 (즉, 수동으로 변경할 수 있지만 항상 2의 거듭 제곱이므로 합계는 간단 함). 여기서 공유 메모리는 상수를 사용하여 할당되어야하지만 호스트로부터 커널에 전달할 값이 필요합니다. n은 커널에 전달하기 위해 CPU에서 볼 수 있도록 (커널 내에서) 공유 메모리의 크기를 선언하는 데 사용할 수 있도록 치수를 어디서 선언해야합니까?

내 코드는 다음과 같이 구성되어있다 : main.cu에서

내가 커널 전화 : kernel.cu에 다음

const int num_states = 1024 
const int dimension = 4 

fitness <<< num_states, dimension * dimension >>> (device_array_of_states, dimension, num_states, device_fitness_return); 

과 내가 가진 :

__global__ void fitness(
    int *numbers, 
    int dimension, 
    int num_states, 
    int *fitness_return) { 
    __shared__ int fitness[16]; <<-- needs to be dimension * dimension 
    //code 
} 

numbers 1024 행렬을 나타내는 배열입니다 , dimension은 행 및 열 길이이고, num_states은 1024입니다. fitness_return 각 행렬의 적합도 값을 보유하는 길이가 1024 인 배열. 커널에서 공유 메모리는 dimension의 제곱으로 하드 코딩되어 있습니다 (이 경우 dimension은 4입니다).

어디에서 어떻게 dimension을 선언하여 공유 메모리를 할당하고 커널을 호출하는 데 사용할 수 있습니까?이 방법으로 한 번에 dimension 만 업데이트하면됩니까? 당신의 도움을 주셔서 감사합니다.

Answer 1

할당 된 공유 메모리의 양은 모든 블록에서 동일합니다. 각 블록에서 실제로 공유 메모리의 양을 늘릴 수도 있지만 여전히 사용 가능합니다. 또한 공유 메모리의 양은 관계없이 제한적이므로 n * n 요소는 maximum amount of space (일반적으로 48KiB)을 초과 할 수 없습니다. float 형 요소 (각 4 바이트)는 n < 340 정도를 의미합니다.

이제 공유 메모리를 할당하는 두 가지 방법이 있습니다 : 정적 및 동적. 이러한 경우

__shared__ int fitness[16]; 

을, 크기가 장치 측 코드 컴파일에서 컴파일 시간 (알려진해야합니다

정적 할당은 작동하지 않을 것입니다 예를 들어,로 준 것입니다 시간) - 이것은 당신에게 해당하지 않습니다.

동적 공유 메모리 할당, 당신이하지 할 커널 코드의 크기를 지정 - leave it empty 당신을하고 extern 앞에 추가 :

extern __shared__ int fitness[]; 

커널을 시작할 때 대신, 당신은 양을 지정을, 다른 블록의 스레드는 그것이 무엇인지를 반드시 알 필요는 없습니다.

그러나 귀하의 경우 스레드 은이 무엇인지 알 필요가 있습니다. 음, 그냥 커널 인수로 전달하십시오.그래서,

__global__ void fitness(
    int *numbers, 
    int dimension, 
    int num_states, 
    int *fitness_return, 
    unsigned short fitness_matrix_order /* that's your n*/) 
{ 
    extern __shared__ int fitness[]; 
    /* ... etc ... */ 
} 

엔비디아의 Parallel-for-all 블로그는 특히 정적 및 동적 공유 메모리 할당을 커버 공유 메모리를 사용을 더욱 심도 깊게 소개와 a nice post 있습니다.