Cuda Block Thread 설정.

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

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

 

그림 참조 : https://phychai.wordpress.com/2015/05/26/cuda-programming-thread-block-grid/

//=====================================================================================
//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); //==>>느려짐