원자 함수는 실제로 CUDA에서 변수를 휘발성으로 만듭니다?

Question

다른 스레드가 해당 변수를 계속 읽는 동안 전역 변수를 업데이트하기 위해 매우 간단한 코드를 스레드 0에 요청했습니다. 그러나 다른 스레드가 실제로 값을 얻지 못함을 발견했습니다.

코드는 매우 간단합니다. 아무도 나에게 그것을 제안하는 방법을 줄 수 있을까? 덕분에 많은

__global__ void addKernel(int *c) 
{ 
int i = threadIdx.x; 
int j = 0; 
if (i == 0) 
{ 
    while(*c < 2000){ 
     int temp = *c; 
     printf("*c = %d\n",*c); 
     atomicCAS(c,temp, temp+1); 
    }  
}else{ 
    while(*c < 1000) 
    { 
     j++; 
    } 
} 

}

Answer 1

내가 비유하고 싶습니다 : 원자 작업이 뮤텍스 것을 잠시 상상이 : 잘 정의 된 프로그램이 될 수 있도록, 두 스레드가 액세스하는 공유 자원은 이어야합니다.은 배타적으로 자원에 액세스하기 위해 뮤텍스를 사용하는 데 동의합니다. 스레드 중 하나가 먼저 뮤텍스를 보유하지 않고 자원에 액세스하면 결과는 정의되지 않습니다.

atomics에서도 마찬가지입니다. 메모리의 특정 위치를 원자 변수로 처리하기로 결정하면 해당 위치에 액세스하는 모든 스레드가 동의해야하며 프로그램에서 의미를 갖도록 처리해야합니다. 만이 원자 적재 및 상점을 통해 조작해야하며 원자 적 및 원자 적 조작의 조합이 아닙니다. 즉

이 :

atomicCAS(c,temp, temp+1); 

은 원자로드 비교 매장을 포함합니다. 결과 명령어는 로드c까지 전역 메모리로 이동하여 비교를 수행하고 새 값으로 전역 메모리로 끝까지 이동합니다.

하지만이 :

while(*c < 2000) 

어떤 방법으로 원자되지 않습니다. 컴파일러 (및 하드웨어)는 c이 다른 스레드에 의해 수정되었을 수도 있다는 것을 알지 못합니다. 따라서 전역 메모리로 이동하는 대신 사용 가능한 가장 빠른 캐시에서 읽습니다. 아마도 컴파일러는 변수를 레지스터에 넣을 것입니다. 왜냐하면 은 현재 스레드에서을 수정하는 다른 사람을 보지 못하기 때문입니다.

while (atomicLoad(c) < 2000) 

하지만 내 지식의 최선을

글을 쓰는 시점에서 CUDA에서 그런 구조가 없다 :

은 당신이 원하는 것은 (상상) 같은 것입니다.

이와 관련하여 volatile 한정자 은입니다. 컴파일러가 변수를 최적화하지 말고 "외부 소스에서 수정 가능"하다고 컴파일러에 알립니다. 이로드가 모든 캐시를 우회하는지는 모르지만 변수의 모든 읽기에 대해로드를 트리거합니다. 실제로는 효과가있을 수 있지만, 이론적으로 나는 당신이 그것에 의존해야한다고 생각하지 않습니다. 게다가, 이것은 변수에 대한 최적화 (예 : 상수 전파 또는보다 나은 성능을 위해 변수를 레지스터로 승격)를 불가능하게합니다.

while(atomicAdd(c, 0) < 2000) 

이것은 글로벌 메모리에서 부하를한다는 것을 원자 명령을 방출한다, 따라서 가장 최근의 표시되어야합니다

는 다음과 같은 해킹을 시도 할 수 있습니다

는 (나는 그것을 시도하지 않은) c의 값. 그러나, 그것은 또한 (이 경우 쓸모없는) 원자 저장소를 소개합니다.