MOLOKINI

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

Graphic Processing Unit 구조 GPU라는 놈이 nVidia가 표준도 아닌것이 가장 많이 사용되고 있기 때문에, nVidia의 CUDA만 본다. 내 연구실 컴퓨터 GPU도 GTX280이다. ㅇㅋ CPU와 GPU의 구조이다. 딱봐도 느껴지는것은 GPU는 병렬구조를 갖고 있다는 것 그리하야 CPU에서 하지 못하는 병렬 계산을 GPU에서는 해낼 수 있다. GPU 전체가 그래픽카드를 의미하는 말은 아니고 그래픽 카드에서 계산을 담당하는 부분이 바로 GPU인 것이다. CUDA는 이미지 및 동영상, 그래픽 처리에 아주 월등한 성능을 보여주고 최근 다양한 용도로 사용되고 있다. 기본 아키텍쳐가 CPU 변수 -> GPU 처리 -> 다시 CPU가 처리된 내용을 처리 물론 복잡한 행렬식이나 벡터계산을 병렬..

컴퓨터 그래픽에서 수학공식 가운데 스플라인 함수를 이용하여 표현하는 매끄러운 곡선스플라인 곡선은 적은 수의 제어점으로 매끄러운 곡선을 만들기 위한 방법의 하나이며, 함수의 차수가 높을 수록 선이 매끄럽게 되고 보통 3차 스플라인이 많이 사용된다.Bezier 곡선과 같이 널리 사용되고 있는데, 스플라인 곡선의 데이터로 Bezier 곡선으로의 변환이 가능한 것이 특징이다.스플라인 곡선은 모든 제어점을 지나가는 것과 지나가지 않는 것이 있으며, 후자는 컴퓨터 그래픽 등으로 곡면을 만드는데에 사용되고 있다.자동차나 비행기의 표면과 같은 유선형의 곡선을 설계할 때 많이 사용된다. 정의는 여기까지 이 스플라인 곡선은, 1차, 2차, 3차 곡선이 있다.차분하게 1차부터 1차 스플라인n개의 점이 주어졌을 때, n-1개..

ANOVA Test ANalysis Of VArience로, 분산분석이라한다. - 샘플링한 조사 집단간의 평균 차이가 우연한 결과인지 아닌지의 여부를 측정하는 기법 - 두 집단 이상의 평균의 차이 검증 - 여러 집단의 변이 또는 분산 정도의 비교를 통한 평균 차이 검증 - 집단 간 변이가 집단 내의 변이보다 상대적으로 차이가 있는가 하는 것에 대한 검증 - 두 모집단 이상의 평균에 대한 추리

테일러 전개 1. 테일러 전개 요약 극한을 매우매우 쉽게 만들어주는 위대한 정리심화미적에서의 극한은 대부분 x -> 0의 극한을 다루는데, 바로 그런 조건에서 초월함수의 근사값으로 대수식을 대체하는 것이다. 2. 테일러 급수그럼 먼저 테일러 급수를 알아야한다. 함수엔 두가지 종류가 있다. 대수함수(무리, 유리함수)와 초월함수 - 대수함수 : 다항, 분수함수,, 흔히 배우는것들 - 초월함수 : 지수함수, 삼각함수 등 대수함수는 근을 구하는게 어렵지 않지만,초월함수는 말그대로 근을 구해내는게 만만치 않다.그래서 초월함수를 다항함수로 근사할 수 있다면 얼마나 좋겠냐 바로 테일러 급수를 이용하면 그것이 가능하다. 테일러 급수는 멱급수로부터 파생되는데 여기서 a_{k}는 미정계수이며 윗 끝인 n을 조절함에 따라 ..