나는 몇 가지 비교를하고 두 객체가 충돌하는지 아닌지를 결정하는 커널을 가지고있다. 충돌 객체의 ID를 출력 버퍼에 저장하려고합니다. 출력 버퍼에 갭이 있기를 원하지 않습니다. 각 충돌을 출력 버퍼의 고유 인덱스에 기록하려고합니다.CUDA : atomicAdd가 스레드를 직렬화하는 데 너무 많은 시간이 걸린다.
그래서 공유 메모리 (로컬 합계)와 전역 메모리 (글로벌 합계)에 원자 변수를 만들었습니다. 아래 코드는 충돌이 발견 될 때 공유 변수의 증가를 보여줍니다. 나는 현재 전역 메모리에서 원자 변수를 증가시키는 것에 문제가 없다.
__global__ void mykernel(..., unsigned int *gColCnt) {
...
__shared__ unsigned int sColCnt;
__shared__ unsigned int sIndex;
if (threadIdx.x == 0) {
sColCnt = 0;
}
__syncthreads();
unsigned int index = 0;
if (colliding)
index = atomicAdd(&sColCnt, 1); //!!Time Consuming!!
__syncthreads();
if (threadIdx.x == 0)
sIndex = atomicAdd(gColCnt, sColCnt);
__syncthreads();
if (sColCnt + sIndex > outputSize) { //output buffer is not enough
//printf("Exceeds outputsize: %d + %d > %d\n", sColCnt, sIndex, outputSize);
return;
}
if (colliding) {
output[sIndex + index] = make_uint2(startId, toId);
}
}
내 문제는 많은 스레드가 원자 변수를 증가 시키려고 할 때 직렬화된다는 것입니다. 접두어 - 합계와 같은 것을 작성하기 전에이 작업을 효율적으로 수행 할 수있는 방법이 있는지 묻고 싶습니다.
이 한 줄로 인해 커널의 경과 시간이 13msec에서 44msec로 증가합니다.
예문 코드를 찾았지만 엔비디아의 토론 게시판이 다운 되었기 때문에 참조 된 링크가 실패합니다. https://stackoverflow.com/a/3836944/596547
편집 : 나는 위의 너무 내 코드의 끝을 추가했습니다. 사실 나는 계층 구조를 가지고 있습니다. 모든 코드 행의 영향을 보려면 모든 객체가 서로 충돌하는 장면, 극단적 인 경우 및 거의 충돌이없는 또 다른 극단적 인 경우를 설정합니다.
마지막으로 전역 변수 (gColCnt)에 공유 원자 변수를 추가하여 충돌 횟수를 알리고 올바른 인덱스 값을 찾습니다. 나는 어떤 식 으로든 여기에 atomicAdd를 사용해야한다고 생각한다.
'atomicAdd'는 정의에 의해 직렬화, 그래서 당신은 충돌이 드문 드문 될 것으로 예상 할 때 당신은 단지에 의존해야한다. 아마도 원자 구조를 계층 적으로 사용하도록 계산을 재구성 할 수 있습니다. 먼저, 각 스레드 블록에서'__shared__' 변수에 누적됩니다. 후 처리 (예 : 위의 3 번째'__syncthreads '이후)에서는 각 블록의 충돌을 전역 메모리의 단일 변수에 누적 할 수 있습니다. –
사실 나는 계층 구조를 가지고 있습니다. 그러나 동일한 블록의 스레드는 모든 객체가 서로 충돌하는 첫 번째 극단적 인 경우에도 __shared__ 변수에 대해 atomicAdd에서 직렬화됩니다. – phoad
www.cuvilib.com/Reduction.pdf M. Harris의 자습서를 발견했습니다. 나는 그것을 이용하려고 노력할 것이다. – phoad