inflearn logo
강의

강의

N
챌린지

챌린지

멘토링

멘토링

N
클립

클립

로드맵

로드맵

지식공유

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

30-4 histogram shared - shared memory 사용

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

185

정재헌

작성한 질문수 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

정재헌

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

25-2 분기커버리지

0

10

0

25-2 이론 재탕

0

11

1

25-1 이론 재탕

0

11

1

서브넷 마스크

0

19

1

모의고사 1회 계산문제 2개 맞았는데

0

31

2

업캐스팅 문제 예시 4번 부모 호출

0

39

1

모의고사 자료 해설이 너무 불친절하네요.

0

75

2

부모의 생성자 출력 관련

0

36

2

24-3 이론파트 재탕문제

0

38

0

누락코드

0

37

1

다차원 배열 포인터 질문있습니다

0

53

3

코딩 살구클럽 가입 문의

0

41

2

24-2 이론 파트 재탕

0

48

1

모의고사 5회

0

42

1

12분 10초 func(values, 0, 1)을 넣을 때

0

40

1

모의고사 1회 4번 문제와 관련한 질문이 있습니다.

0

43

2

23- 3 재탕문제

0

63

1

23- 1 재탕문제

0

46

0

파이썬 추가 코드

0

58

1

배열 포인터 출력 표시

0

54

2

포인터 증감 또는 값 증감

0

55

2

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

0

646

1

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

0

395

1

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

0

460

1