공유 메모리 할당

공유 메모리 할당

상수 매개 변수를 사용하여 공유 메모리를 할당하려고 하는데 오류가 발생했습니다. 커널은 다음과 같습니다.

__global__ void Kernel(const int count)
{
    __shared__ int a[count];
}

그리고 다음과 같은 오류가 발생합니다.

오류: 식의 값은 상수여야 합니다.

카운트는 상수입니다!이 오류가 발생하는 이유는 무엇입니까?어떻게 하면 이 상황을 피할 수 있을까요?

CUDA는 동적 공유 메모리 할당을 지원합니다.커널을 다음과 같이 정의하는 경우:

__global__ void Kernel(const int count)
{
    extern __shared__ int a[];
}

그런 다음 커널 실행의 세 번째 인수로 필요한 바이트 수를 전달합니다.

Kernel<<< gridDim, blockDim, a_size >>>(count)

실행 시 크기를 조정할 수 있습니다.런타임은 블록당 동적으로 선언된 단일 할당만 지원합니다.추가로 필요한 경우 해당 단일 할당 내에서 오프셋에 대한 포인터를 사용해야 합니다.또한 공유 메모리가 32비트 워드를 사용하는 포인터를 사용할 때는 공유 메모리 할당 유형에 관계없이 모든 할당이 32비트 워드로 정렬되어야 합니다.

const"읽기 전용"을 의미하는 것이 아니라 "읽기 전용"을 의미합니다.

상수 식은 컴파일러가 컴파일할 때 값을 알 수 있는 식입니다.

1:값으로 합니다("1"과 하지 않음: " 상값사와여하공용메선유언리모않음옵지을하동일션수와▁option선()언▁with않▁shared음▁one").const)

__global__ void Kernel(int count_a, int count_b)
{
    __shared__ int a[100];
    __shared__ int b[4];
}

옵션 2: 커널 시작 구성에서 공유 메모리를 동적으로 선언합니다.

__global__ void Kernel(int count_a, int count_b)
{
    extern __shared__ int *shared;
    int *a = &shared[0]; //a is manually set at the beginning of shared
    int *b = &shared[count_a]; //b is manually set at the end of a
}

sharedMemory = count_a*size(int) + size_b*size(int);
Kernel <<<numBlocks, threadsPerBlock, sharedMemory>>> (count_a, count_b);

참고: 동적으로 공유된 메모리에 대한 포인터는 모두 동일한 주소로 지정됩니다.공유 메모리에서 두 개의 어레이를 수동으로 설정하는 방법을 설명하기 위해 두 개의 공유 메모리 어레이를 사용합니다.

"CUDAC 프로그래밍 가이드"에서:

실행 구성은 다음 양식의 식을 삽입하여 지정합니다.

<<<Dg, Db, Ns, S>>>

위치:

• Dg는 dim3 유형이며 그리드의 치수와 크기를 지정합니다...
• Db는 dim3 유형이며 각 블록의 치수와 크기를 지정합니다...
• Ns는 size_t 유형이며 정적으로 할당된 메모리 외에 이 호출에 대해 블록당 동적으로 할당되는 공유 메모리의 바이트 수를 지정합니다.이 동적으로 할당된 메모리는 __shared__에 언급된 외부 배열로 선언된 변수에 의해 사용됩니다. Ns는 기본값이 0인 선택적 인수입니다.
• S는 cudaStream_t 유형이며 연결된 스트림을 지정합니다...

따라서 동적 매개 변수 Ns를 사용하여 사용자는 커널에 공유 변수가 아무리 많더라도 하나의 커널 함수가 사용할 수 있는 공유 메모리의 총 크기를 지정할 수 있습니다.

이와 같이 공유 변수를 선언할 수 없습니다.

__shared__ int a[count];

어레이의 최대 크기에 대해 충분히 확신한다면 다음과 같이 직접 선언할 수 있습니다.

__shared__ int a[100];

하지만 이 경우에는 프로그램에 블록이 몇 개 있는지 걱정해야 합니다. 공유 메모리를 블록에 고정하고 완전히 활용되지 않으면 글로벌 메모리(높은 지연 시간)를 사용하여 컨텍스트 전환을 수행하게 되므로 성능이 저하됩니다.

이 문제에 대한 좋은 해결책이 있습니다.

extern __shared__ int a[];

그리고 메모리에서 커널을 호출하는 동안 메모리를 할당합니다.

Kernel<<< gridDim, blockDim, a_size >>>(count)

그러나 커널에서 할당하는 것보다 블록에서 더 많은 메모리를 사용하는 경우 예기치 않은 결과를 얻을 수 있기 때문에 여기에서도 문제가 발생해야 합니다.

언급URL : https://stackoverflow.com/questions/5531247/allocating-shared-memory