워프는 원자 적 연산과 어떻게 작동합니까?

Question

워프의 스레드는 물리적으로 병렬로 실행되므로, 스레드 X 중 하나가 원자 연산을 시작하면 다른 연산은 무엇을할까요? 기다림? 스레드 X가 원자 큐에 푸시되고 액세스 (뮤텍스)를 받고 메모리로 보호되는 동안 모든 스레드가 대기하게됩니다.이 뮤텍스로 보호되고 실제 뮤텍스가 보호됩니다.

일부 작업에 다른 스레드를 사용할 수있는 방법이 있습니까? 예를 들어, 일부 메모리를 읽는 것과 같은 방법이 있습니까? 원자 연산이 대기 시간을 숨길 수 있습니까? 내 말은, 15 개의 유휴 스레드는 잘 안되네. 원자는 정말로 느립니다. 그렇죠? 어떻게 가속화 할 수 있습니까? 어떤 패턴으로 작업 할 수 있습니까?

공유 메모리가있는 원자 적 연산이 은행이나 전체 메모리에 대해 잠겨 있습니까? 은 (mutexs없이) 예를 들어, __shared__ float smem[256];

• Thread1이 다른 은행과 함께 작동 Thread2이 atomicAdd(smem + 1, 1);

그 스레드를 실행 atomicAdd(smem, 1);

을 실행하지만 일반적으로 공유 메모리에있다. 그들은 parralel을 실행합니까, 아니면 큐에 대기 할 것입니까? Thread1과 Thread2가 분리 된 warps 또는 일반적인 war1 인 경우이 예제와 어떤 차이가 있습니까?

Answer 1

10 개의 질문과 비슷한 것으로 기록됩니다. 대답하기가 어렵습니다. 질문 당 하나의 질문을하는 것이 좋습니다.

일반적으로 워프의 모든 스레드는 동일한 명령어 스트림을 실행합니다. 그래서 두 가지 경우 우리가 고려할 수있다 : 조건문없이

1. (예를 들어 ... 다음 ... 다른 경우) 발생이 경우, 모든 스레드가 실행되는 동일한 명령은 원자 명령 할 수 . 그런 다음 32 개 스레드 모두가 같은 위치에있을 필요는 없지만 원자를 실행합니다. 이 atomics는 모두 SM에 의해 처리되고 어느 정도는 직렬화됩니다 (동일한 위치를 업데이트하는 경우 완전히 직렬화됩니다).
2. (조건이) 예를 들어 if (!threadIdx.x) AtomicAdd(*data, 1);이라고 가정하면 스레드 0은 원자를 실행하고 은 그렇지 않습니다. 우리가 다른 사람에게 할 일을 줄 수있는 것처럼 보일지도 모르지만 다른 것, 그러나 lockstep warp 실행은 이것을 허용하지 않습니다. if (true) 경로를 사용하는 모든 스레드가 함께 실행되고 if (false) 경로를 실행하는 모든 스레드가 함께 실행되지만 true 및 false 경로가 직렬화되도록 워프 실행이 직렬화됩니다. 다시 워프에서 서로 다른 명령어를 실행하는 스레드를 다르게 지정할 수는 없습니다 .

그물은 워프 내에서 하나의 스레드가 원자 단위로 처리 할 수 ​​없지만 다른 스레드는 동시에 다른 처리를 수행 할 수 없습니다.

많은 다른 질문은 메모리 트랜잭션이 시작된 명령 사이클이 끝날 때 완료 될 것으로 예상됩니다. 전역 메모리와 공유 메모리를 사용하여 이전 쓰기 트랜잭션을 다른 스레드에서 볼 수 있도록 코드에서 특별한 단계를 수행해야합니다 (트랜잭션 완료의 증거라고 할 수 있음).) 이것을 수행하는 일반적인 방법 중 하나는 __syncthreads() 또는 __threadfence()과 같은 장벽 명령어를 사용하는 것입니다.하지만 이러한 장벽 명령어가 없으면 스레드는 쓰기가 완료 될 때까지 대기하지 않습니다. A (읽기에 의존적 인 연산)은 스레드를 정지시킬 수 있습니다. 쓰기는 일반적으로 스레드를 멈추게 할 수 없습니다. 그 중 하나가 원자 작업을 시작하면

는

때문에, 다른 무엇을 할 것인가 :

지금 당신의 질문에 대해 볼 수 있습니다? 기다림?

아니요, 그들은 기다리지 않습니다. 원자 연산은 SM에서 원자 단위를 처리하는 기능 단위로 전달되고 모든 스레드는 함께 잠금 단계로 진행됩니다. 원자는 일반적으로 읽기를 의미하기 때문에 읽기가 워프를 멈출 수 있습니다. 그러나 스레드는 원자 연산이 완료 될 때까지 기다리지 않습니다 (즉, 쓰기). 그러나이 위치의 후속 읽기 은 원자 (쓰기)가 완료되기를 기다리는 워프를 다시 정지시킬 수 있습니다. 전역 메모리를 업데이트하는 것이 보장되는 전역 원자의 경우 원래의 SM (활성화 된 경우) 및 해당 위치에 항목으로 포함 된 L2의 L1을 무효화합니다.

일부 작업에 다른 스레드를 사용할 수있는 방법이 있습니까? 예를 들어, 원자 적 연산은 대기 시간을 숨길 수 있습니까?

사실 처음에는 언급 한 이유가 있습니다.

원자는 정말로 느립니까? 어떻게 가속화 할 수 있습니까? 어떤 패턴으로 작업 할 수 있습니까?

예, 아토 훨씬 더 천천히들이 활동을 지배하는 경우 프로그램 실행을 할 수 있습니다 (예 : 순진 감소 또는 순진 히스 토 그래밍 등을.) 일반적으로 말하기, 원자 작업을 가속화 할 수있는 방법을 사용하거나 사용하지 않을 것입니다 드물게 프로그램 활동을 지배하지 않는 방식으로 예를 들어 순진한 감소는 모든 요소를 ​​전체 합계에 추가하기 위해 원자를 사용합니다. 똑똑한 병렬 축소는 스레드 블럭에서 수행되는 작업에 전혀 아토믹을 사용하지 않습니다. 스레드 블록 감소가 끝나면 스레드 블록 부분 합을 전역 합계로 업데이트하는 데 단일 원자를 사용할 수 있습니다. 즉, 임의로 많은 수의 요소를 32 원자 추가 또는 그 이하의 순서로 빠르게 병렬로 줄일 수 있음을 의미합니다. 그것보다는 오히려 2.

공유 메모리 단일 커널 호출 할 수있는 병렬 감소를 가능하게한다는 점을 제외 아토의이 스페어 링의 ​​사용은 기본적으로, 전체 프로그램 실행에 띄지 않을 것이다 : 그들은 parralel을 실행하거나 않음을 그들은 대기열에 올거야?

대기열에 있습니다. 그 이유는 공유 메모리에서 원자 연산을 처리 할 수있는 기능 유닛의 수가 제한되어 있기 때문에 단일 사이클에서 워프의 모든 요청을 처리하기에 충분하지 않기 때문입니다.

원자력 작업의 처리량과 관련된 질문에 답하는 것을 피했습니다. AFAIK 문서에이 데이터가 잘 지정되어 있지 않기 때문입니다. 충분한 동시 또는 거의 동시적인 원자 연산을 실행하면 원자 기능 유닛이 가득 찬 대기열로 인해 원자 명령에 일부 워프가 실속 할 수 있습니다. 나는 이것이 사실이라는 것을 모르고 그것에 관해서는 대답 할 수 없다.