MSW_Creator

Nvidia Cuda GPU연산을 위한 기본 Core수를 알자.

https://www.studio1productions.com/Articles/NVidia-GPU-Chart.htm

Cuda Thread, Grid, Block 설정 시에 참고하세요.

Grid->Block->Thread와 각 Thread에 위치 정보와 Data 전달 개념을 잡아가야 합니다. 

//=====================================================================================
//Cuda102.cuh
	typedef struct _XYPos_
	{
		int x;
		int y;

		int blockIdx_x;
		int blockDim_x;
		int threadIdx_x;

		int blockIdx_y;
		int blockDim_y;
		int threadIdx_y;
	}XYPos;
	class CudaProc
	{
	public:
		CudaProc(void);
		virtual ~CudaProc(void);
	//----------------------------------------------------------
	public:
		int ProcBlockThread(int nXCnt, int nYCnt, XYPos *pXYList);
	};



//=====================================================================================
//Cuda102.cu
__global__ 
void XYPosCheck(int nXCnt, int nYCnt, XYPos *pXYList)
{
	//1) 2차원 배열 접근 형태의 tidX와 tidY를 구한다.
	int tidX = blockIdx.x * blockDim.x + threadIdx.x;
	int tidY = blockIdx.y * blockDim.y + threadIdx.y;

	//2) tidX와 tidY는 16의 배수형태로 Thread가 할당되어 주어진 Memory보다 많게 된다.
	//  Count X와 Count Y의 범위 내에서만 메모리에 접근하여 Data를 체워준다.
	if (tidX < nXCnt && tidY < nYCnt)
	{
		pXYList[tidX + (nXCnt * tidY)].x = tidX;
		pXYList[tidX + (nXCnt * tidY)].y = tidY;

		pXYList[tidX + (nXCnt * tidY)].blockIdx_x = blockIdx.x;
		pXYList[tidX + (nXCnt * tidY)].blockDim_x = blockDim.x;
		pXYList[tidX + (nXCnt * tidY)].threadIdx_x = threadIdx.x;

		pXYList[tidX + (nXCnt * tidY)].blockIdx_y= blockIdx.y;
		pXYList[tidX + (nXCnt * tidY)].blockDim_y= blockDim.y;
		pXYList[tidX + (nXCnt * tidY)].threadIdx_y = threadIdx.y;
	
		float fAsphericZ_mm = K_Aspheric(12.3, (float)tidX, (float)tidY);
	}
}
__host__
int CudaProc::ProcBlockThread(int nXCnt, int nYCnt, XYPos *pXYList)
{
	int nCnt_X = nXCnt;
	int nCnt_Y = nYCnt;

	//1)처리를 위한 data 생성.
	int nThreadCnt = (nCnt_X*nCnt_Y);
	int nRetMemsize = nThreadCnt * sizeof(XYPos);

	XYPos  *C_pRet_XY = nullptr;
	cudaMalloc((void**)&C_pRet_XY, nRetMemsize);

	//2) Cudat Thread, Block 할당.

	dim3 threads(16, 16);

	int nGridCntX = nCnt_X / threads.x + (nCnt_X % threads.x > 0 ? 1 : 0);
	int nGridCntY = nCnt_Y / threads.y + (nCnt_Y % threads.y > 0 ? 1 : 0);
	dim3 grid(nGridCntX, nGridCntY);
	//a) 기본 Thread Count는 16*16의 크기로 256개를 잡는다.
	//b) grid는 입력된 X Data의 X성분을 16개의 X성분 Thread로 나누어 준다.
	//c) grid는 입력된 Y Data의 Y성분을 16개의 Y성분 Thread로 나누어 준다.
	//d) 나눈 나머지가 발생하면 Grid항목 1(thread 16개)을 추가 한다.
	//(XYPosCheck()안에서 tidX와 tidY의 법위가 넘는것은 Thread수의 배수로 증가 해서이다.)

	//3) Data 처리
	XYPosCheck << < grid, threads >> > (nXCnt, nYCnt, C_pRet_XY);

	//4) 결과 Data 전달.
	cudaMemcpy((void*)pXYList, (void*)C_pRet_XY, nRetMemsize, cudaMemcpyDeviceToHost);

	//5) 생성 메모리 처리.
	cudaFree(C_pRet_XY);
	return 1;
}
//=====================================================================================
//CudaTestDlg.cpp
void CCudaTestDlg::OnBnClickedBtCudaTest()
{
	CudaProc HWAccGPU;
	int nXCnt = 210;
	int nYCnt = 260;

	XYPos *pXYList= new XYPos[nXCnt * nYCnt];
	//Cuda Object
	HWAccGPU.ProcBlockThread(nXCnt, nYCnt, pXYList);

	for (int iy = 0; iy < nXCnt * nYCnt; iy++)
	{
		TRACE(L"==.....>>>>>>%d) X:%d, Y:%d\r\n", iy, pXYList[iy].x, pXYList[iy].y);
		TRACE(L"==.....            X) %d, %d, %d\r\n", 
        		pXYList[iy].blockIdx_x, pXYList[iy].blockDim_x, pXYList[iy].threadIdx_x);
		TRACE(L"==.....            Y) %d, %d, %d\r\n", 
        		pXYList[iy].blockIdx_y, pXYList[iy].blockDim_y, pXYList[iy].threadIdx_y);
	}

	delete[] pXYList;
}
//=====================================================================================
//결과

==.....>>>>>>0) X:0, Y:0
==.....            X) 0, 16, 0
==.....            Y) 0, 16, 0
==.....>>>>>>1) X:1, Y:0
==.....            X) 0, 16, 1
==.....            Y) 0, 16, 0
==.....>>>>>>2) X:2, Y:0
==.....            X) 0, 16, 2
==.....            Y) 0, 16, 0
==.....>>>>>>3) X:3, Y:0
==.....            X) 0, 16, 3
==.....            Y) 0, 16, 0
==.....>>>>>>4) X:4, Y:0
==.....            X) 0, 16, 4
==.....            Y) 0, 16, 0
==.....>>>>>>5) X:5, Y:0
==.....            X) 0, 16, 5
==.....            Y) 0, 16, 0
==.....>>>>>>6) X:6, Y:0
==.....            X) 0, 16, 6
==.....            Y) 0, 16, 0


...

== .....>>>>>>54594) X:204, Y : 259
== .....               X) 12, 16, 12
== .....               Y) 16, 16, 3
== .....>>>>>>54595) X:205, Y : 259
== .....               X) 12, 16, 13
== .....               Y) 16, 16, 3
== .....>>>>>>54596) X:206, Y : 259
== .....               X) 12, 16, 14
== .....              Y) 16, 16, 3
== .....>>>>>>54597) X:207, Y : 259
== .....              X) 12, 16, 15
== .....              Y) 16, 16, 3
== .....>>>>>>54598) X:208, Y : 259
== .....              X) 13, 16, 0
== .....              Y) 16, 16, 3
== .....>>>>>>54599) X:209, Y : 259
== .....              X) 13, 16, 1
== .....              Y) 16, 16, 3

1) 입력 Data는 54600개이고, X 210, Y 260개 Data라고 보면 된다.

2) XYPosCheck() 함수 안에서 tidX, tidY의  접근 인자들을 정상적으로 Memory에 표기하여 돌려보내준다.

3) X와 Y Data크기가 tidX, tidY와 같은 크기로 설정하여 적용하면 의미 없는 Thread는 없어지지만 원하는 속도가 나오기 힘들다.(2의 배수로 Threa Block이 할당된다.)

Thread Block Test
dim3 threads(64, 64); //==>>오동작
dim3 threads(32, 32); //==>>느려짐
dim3 threads(16, 16); //==>>보통
dim3 threads(8, 8); //==>>미미하게 빨라짐
dim3 threads(4, 4); //==>>느려짐

uafxcwd.lib  error LNK2005

Link 추가 연결

1차 Test

대상 : 644*484 Color이미지의 Red와 Green색상 변경.

코드 : CPU용 이중 for문으로 구성, GPU Cuda 병렬 처리 방식으로 구성.

속도 : CPU 3 msec, GPU 11 msec CPU가 빠르다.(debug, release 같은 현상)

       -> Memory 생성 복사가 추가 되었음으로 CPU 쪽에도 추가 ( 속도 변화 없음)

결론 : Data양이 적을때에는 CPU가 3배 이상 빠르다.(GPU로 메모리 복사가 생기지 않음으로)

     - CPU : 적은 Data, 복잡한 연산에 적합.

     - GPU : 다량의 Data, 단순(CPU에 비해) 연산에 적합.

2차 Test

대상 : 상동

코드 : 각 처리위치 반복 수행 제일 안쪽에서 같은 동작 1000회 반복.

속도 :  debug Mode : CPU 2438 msec, GPU 2829 msec

         Release Mode: CPU 1180msec, GPU 3 msec(393배 빠름)

*cudaDeviceSynchronize();로 GPU 동작 완료 체크.

추가 : 이미지 크기 9999*7514 크기 Test

   1회교환

        Debug : CPU 930 msec, GPU  : 2598 msec

        Release : CPU 139 msec, GPU : 358 msec

   100회 교환

        Release : CPU 32288 msec, GPU : 378(최초 511) msec

   1000회 교환

        Release : CPU 측정 불가, GPU : 352 msec

1) Cuda 10.2 Download

2) "cuda_10.2.89_441.22_win 10.exe"설치(2.5 GByte)

3) VS2017 설정 (Consol Test Project 생성 후)

4) Cuda Lib 추가.(Debug & Release : cuda.lib; cudart.lib; cublas.lib; cufft.lib)

5) Test.cuh, Test.cu 파일 추가.

6) 기본 Test 코드 작성

7) 컴파일 및 실행