MOLOKINI

행렬과 벡터의 곱셈 프로그램으로 최적화 기법을 적어보았다 그리고 메모리 사용에 있어 다양한 경우에 효과적인 메모리 사용 방법을 선택할 수 있는 지표를 제시했다 메모리 읽기와 쓰기의 참조방식에 있어 프로그램의 유형은 아래와 같이 분류할 수 있다. - 다수의 트랜잭션으로 전역메모리 읽기, 쓰기를 하는 프로그램 * 메모리 접근이 빈번할 때에는 공유 메모리를 사용하여 읽기, 쓰기를 하는것이 좋다. * NVIDIA GPU Computing Development Kit의 'transposeNew' 프로그램 참고 - 소수의 트랜잭션으로 전역메모리 읽기, 쓰기를 하는 프로그램 * 공유 메모리를 사용하자, 그런데 전역메모리 참조횟수를 줄여 이득을 얻지 못하는 경우에는 이득을 얻지 못한다. - 소수의 트랜잭션으로 전역메모..

알고리즘, 블록, 메모리, 분기문 중 분기문이다. 행렬과 벡터의 곱셈 후 블록 크기의 공유 메모리에 저장되어 있는 값을 합하는 과정에서 각각의 스레드 인덱스를 제어하여 필요한 값만을 합할 수 있다. 하지만, if문을 이용해 스레드를 제어하는 것 보다는 사용하게 되지 않는 값의 합까지 동일하게 병렬적으로 계산하도록 함으로 성능이 더 향상되어 if문을 제거하였다. 하프워프 단위 이하 스레드 공유메모리 접근시 읽기 및 쓰기 충돌로 인한 오류값이 발생하지 때문에, 동기화를 위한 __syncthreads() 함수 호출은 필요없다 그래서 알고리즘, 블록, 메모리, 분기문 최적화를 통해 무려 56배나 빨라진 계산속도를 볼 수 있다. 최적화된 행렬 곱셈 - 알고리즘, 블록, 메모리, 분기문 최적화 ThreadID = ..

행렬 M * 벡터 V = 결과 W 의 최적화 알고리즘, 블록, 메모리, 분기문 중 알고리즘 파트 봤고 이번에는 블록과 메모리 최적화 부분이다. 블록 최적화 스레드 블록 크기 : 블록 당 최대 스레드 수, 그리드 크기 : SM의 수 X 2 (60) 스레드는 그렇다치는데 그리드는 SM의 수 X 2가 왜 최적화일까? - 활성화 비율을 고려하여 SM당 2개의 블록이 할당되도록 한다. - 여러개의 블록이 할당 될 경우 뱅크충돌이 일어나구 그러면 성능에 악영향을 미친다 메모리 최적화 초기방법 : 커널의 입력데이터 행렬 M과 벡터 V의 인수저장에는 전역메모리 사용 개선방법 : 행렬 M의 행은 변하지만 곱해지는 벡터 V는 계속 같은 내용이 호출되기 때문에 전역메모리보다 빠른 텍스쳐 메모리에 벡터 V를 저장하여 성능향상..

행렬과 벡터의 곱셈을 통해 최적화 기법의 영향을 분석해본다. 알고리즘, 블록 크기, 메모리, 분기문제거 요 네가지 부분으로 알아본다. 알고리즘 파트 행렬 M : NxN 벡터 V : Nx1 결과행렬 W : Nx1 곱하는 행렬의 행과 곱해지는 행렬의 열이 결과행렬이니까 이걸 가지고 최적화 한거랑 안한거랑 성능평가를 할 겁니다 기존의 알고리즘대로 하자면 그리드와 블록은 1차원 하나의 스레드는 결과 벡터 W의 복수의 항에 들어갈 값을 계산한다. 행렬 M의 블록 ID번째 행과 벡터 V의 벡터내적을 통하여 벡터 W의 블록 ID번째 항을 계산한다. - 그니까 행렬의 ID행과 벡터의 ID항을 곱해, 걍 행렬곱 스레드 크기보다 큰 행은, '행번호 / 스레드크기의 나머지'번째의 스레드가 계산을 한다. 최적화된 알고리즘대로..

__syncthreads() : 각 스레드의 작업을 모두 동일하게 맞추어준다 - 때문에 지연 감춤의 효과를 떨어뜨린다. - 빠른속도의 공유 메모리를 사용한다고해서 반드시 유리하지만은 않다는 것을 의미한다. 공유메모리에 전역 메모리 일부를 읽어 데이터를 처리하는 방법을 많이 사용하지만, 공유메모리에 전역 메모리를 복사한 후 __syncthreads()를 사용할 경우 지연 감춤은 일어나지 않는다. 즉, 전역메모리 참조 횟수 대비 적당한 시간의 명령 수행시간이 확보되고, __syncthread()가 그 명령 수행시간 사이를 구분짓지 않는 경우, 지연감춤이 일어나는 전역 메모리만 사용하는 경우가 공유메모리를 함께 사용하는 것 보다 더 나을 수 있다. 결국은 메모리 매니지먼트가 그만큼 GPU 프로그래밍에서는 상당..

상대적인 전역 메모리 참조 시간 전역 메모리 참조는 상당한 수의 사이클을 소비하기 때문에, 프로그램 수행 시 참조 횟수에 따라 성능이 크게 변한다. 같은 스레드 블록, 그리드 구조의 프로그램은 커널 내부에서 더 적은 횟수의 전역 메모리 참조를 하는 것이 더 빠르다. 여기서 전역메모리 참조라 하는것은 지역메모리 참조도 포함한다. 하지만, 전역메모리 참조를 하는데 걸리는 시간은 조건에 따라 크게 달라진다. 같은 전역 메모리 참조라도 커널의 전체 실행시간에 끼치는 영향은 더 클 수도 작을 수도 있는데, 커널에서 전역 메모리 참조를 할 때 발생하는 지연 감춤 때문이다. CUDA의 전역 메모리 지연 감춤 명령은 32개씩 워프단위로 실행된다. 하지만 SM 안의 프로세스 코어와 TPC(텍스쳐 프로세서 클러스터)의 전..

SM당 활성 스레드의 개수는 GPU Compute Capability를 통해 확인한다. nVidia 프로파일러에서 occupancy 항목을 확인한다. GTX280은 1.3이다. Compute Capability는 그냥 그래픽카드가 지원하는 계산능력 버전(?) 뭐 그 정도로 해두자 요게 바로 Compute Capability 버전에 따른 계산능력 스펙 차이 최대 활성 스레드의 수는 1024, 컴파일러 옵션 : -ptxas-options = -v를 통해 커널이 사용하는 레지스터 수와 지역메모리, 공유메모리, 상수메모리 수를 알 수 있다. - 의도와 다르게 레지스터 대신 지역메모리가 사용될 수 있기 때문에 컴파일 결과를 통해 이를 반드시 확인한다. -maxrregcount : 최대 레지스터 수 제한 - 레지스..

활성화 비율 = SM 활성화 워프 수 / SM 활성화 최대 워프 수 여기서 워프를 다시 말하면, SM당 존재하는 32개의 스레드 GTX280은 30개의 SM으로 구성되어 있으니깨, 30개의 워프가 있는것이다 활성화 비율을 높여야 하는 이유 활성화 된 스레드들이 많아야 메모리 참조시 발생하는 지연시간 동안 다른 활성화 된 스레드에 대한 연산을 수행하도록 해서 SM이 쉬는 시간을 줄이게 될 수 있기 때문이다. SM을 놀게하면 GPU의 연산능력을 최대화 시키지 못해 성능 저하를 유발한다. 한 블록은 한 SM에서 처리된다, SM에는 자원 즉, 레지스터, 스레드 수, 공유메모리가 허용하는 최대한의 블록들이 할당되어 활성화 된다. 활성화 된 블록의 수를 늘리면 SM을 유효화(일시키고) 지연 감춤 효과를 유도할 수 ..

무조건이라고 말할 수는 없지만, CPU 프로그래밍(걍 평소 코딩)을 할 때에는 전역변수의 사용은 자제하고 지역변수의 사용을 최대화 시켜야만 했다. 왜냐, 전역변수는 프로그램이 실행되는 동안 계속 메모리를 차지하고 있고, 지역변수는 메서드가 종료되는 순간 소멸되기 때문에 하지만, 우리의 GPU는 그렇지 않다. 메모리의 속도 자체가 전역메모리가 더 빠르기 때문에, 지역메모리보다는 전역메모리를 사용해야한다, 물론 공유메모리나 레지스터를 사용하는 방법이 가장 좋은 방법이다. 그럼, 전에 썼던 것 처럼,, 가장 빠른 메모리인 레지스터만 사용하다보면 레지스터에 트랜잭션들이 밀리고 밀린 트랜잭션들이 병렬처리가 아닌 순차처리되는 상황이 생기기 때문에 성능이 되려 느려진다고 했다. 이렇게 레지스터가 사용량이 많아지면 어..

계층적 메모리 접근 최적화 무슨말이가? 그거슨, GPU에는 여러 메모리가 있는데, 그 메모리들을 빠른 순서대로 사용한다는 의미로 속도순으로 나열하면 레지스터 = 공유메모리 > 상수메모리 > 텍스쳐메모리 > 전역메모리 > 지역메모리 요런 형태로, 레지스터가 가장 빠르고 레지스터와 비슷한 접근속도를 갖는것이 바로 공유메모리 우리는 어떻게 하면 이 레지스터와 공유메모리에서 계산을 할 것인가를 생각해야한다. 그렇지만, 무조건 접근 메모리를 빠른 메모리를 사용한다고해서 빠르기만 한 것도 아니다. 빠른 메모리를 사용했는데도 느린 예를 보여드리겠습니다, 뱅크충돌 : 스레드가 동시에 같은 뱅크영역을 접근하려할 때 발생하는 문제 동시에 같은 영역을 접근하려 했기 때문에, 스레드는 한번에 메모리에 접근하지 못한다 때문에 ..