inflearn logo
강의

강의

N
챌린지

챌린지

멘토링

멘토링

N
클립

클립

로드맵

로드맵

지식공유

CUDA 프로그래밍 (5) - C/C++/GPU 병렬 컴퓨팅 - 아토믹 연산 atomic op

30-4 histogram shared - shared memory 사용

[30-4] hist-shared.cu 에서 질문있습니다.

159

정재헌

작성한 질문수 4

0

"hist-shared.cu"의 다음 부분에서 질문있습니다.

if(threadIdx.x < HIST_SIZE) {
    atomicAdd(&(hist[threadIdx.x]), s_hist[threadIdx.x]);
}

threadIdx.x는 0~31로 모두 다르기 때문에 병렬 연산이긴 해도, global memory에 있는 hist 배열의 서로 다른부분으로 write가 이루어진다고 생각했습니다.

 

따라서 서로 겹치는 위치가 없어서 atomic한 연산이 필요없을거라 생각했는데, atomic 연산이 없으면 결과가 이상하게 나오더라구요..

 

혹시 왜 그런지 알 수 있을까요?

 

양질의 강의 제공해주셔서 감사합니다!

c c++ cuda gpu 병렬-처리

답변 2

1

드립커피+한모금더

안녕하세요.

thread block 1개를 기준으로 보면, 말씀하신 내용이 맞습니다.

다만, 현재 예제에서는 thread block 을 여러 개 사용하고, 이것들이 동시에 실행되고 있습니다.

즉, thread block 여러 개에서, theadIdx.x 는 0 ~ 31인 thread 들이, 각 thread block 에서 동시에 메모리 영역에 write 하려고 하므로, 결국 atomicAdd 가 아니면 답이 틀리게 됩니다.

감사합니다.

0

정재헌

이해했습니다 감사합니다!

최신 엔비디아 CUDA 아키텍처에서의 결정적 변경 사항

0

10

1

자문자답- 맞는지 틀린지 확인부탁드립니다.

0

11

0

55강 파이썬에만있는 연산자들

0

11

1

55강의 파이썬에서만 있는 연산자들

0

11

1

[문의] 강의 코드 제공여부

0

24

2

코딩살구클럽 가입 문의

0

38

2

코딩 살구 클럽 컴파일 에러

0

22

1

추천 문제

0

22

2

코딩살구클럽 승인

0

31

1

이 강의를 다 들으면 어떤 강의를 들어야 하나요?

0

42

2

선생님 13.3을 써도 큰 차이가 없나요?

0

33

2

코살구 1주차 1940번 문제 조건과 프라이빗 테스트 불일치 문의

0

35

2

메모리 동적할당시 메모리창 빨간 글씨

0

27

2

문제를 고민하는 시간 관련

0

28

2

29강 5:00

0

33

2

코딩살구클럽

0

43

2

코딩살구클럽 문의

0

47

2

코딩살구클럽 승인

0

38

2

띄어쓰기

0

38

2

공부 순서

1

47

2

DP 경우의 수 설명이 이해가 되지 않습니다.

0

36

2

선생님 gpu->cpu 속도 개선에 대해서 질문드려요

0

618

1

32-2의 warp의 early terminate에 대해 질문드립니다.

0

373

1

28-4 shuffle shared - shared memory 강의에서 질문이 있습니다.

0

434

1