fix bugs in computeGD using prediction

Author: zeyiwen

Device/DevicePredictorHelper.cu

@@ -43,11 +43,7 @@ __global__ void PredMultiTarget(real *pdTargetValue, int numofDenseIns, const Tr
 
 	int pid = 0; //node id
 	const TreeNode *curNode = pAllTreeNode + pid;
-	if(curNode->nodeId != 0)
-	{
-		printf("id of root node is %d should be 0\n", curNode->nodeId);
-		return;
-	}
+	CONCHECKER(curNode->nodeId == 0);
 	int counter = 0;
 	while(curNode->featureId != -1)//!curNode->isLeaf()
 	{
@@ -69,11 +65,7 @@ __global__ void PredMultiTarget(real *pdTargetValue, int numofDenseIns, const Tr
 		curNode = pAllTreeNode + pid;
 
 		counter++;
-		if(counter > maxDepth)//for skipping from deadlock
-		{
-			printf("%s has bugs; fid=%d\n", __PRETTY_FUNCTION__, fid);
-			break;
-		}
+		CONCHECKER(counter <= maxDepth);
 	}
 	pdTargetValue[targetId] += pAllTreeNode[pid].predValue;
 }

Device/FindSplit/ComputeGD.cu

@@ -72,7 +72,7 @@ void DeviceSplitter::ComputeGD(vector<RegTree> &vTree, vector<vector<KeyValue> >
 	KernelConf conf;
 	//start prediction
 //###################  for my future experiments
-bool bOptimisePred = true;
+bool bOptimisePred = false;
 	if(nNumofTree > 0 && numofUsedFea >0 && bOptimisePred == false)//numofUsedFea > 0 means the tree has more than one node.
 	{
 		uint startPos = 0;
@@ -85,14 +85,11 @@ bool bOptimisePred = true;
 		int *pNumofFea = manager.m_pDNumofFea + startInsId;
 		int numofInsToFill = nNumofIns;
 
-		dim3 dimGridThreadForEachIns;
-		conf.ComputeBlock(numofInsToFill, dimGridThreadForEachIns);
-		int sharedMemSizeEachIns = 1;
 		//memset dense instances
 		real *pTempDense = manager.m_pdDenseInsEachBag + bagId * bagManager.m_maxNumUsedFeaATree * bagManager.m_numIns;
 		checkCudaErrors(cudaMemset(pTempDense, -1, sizeof(real) * bagManager.m_maxNumUsedFeaATree * bagManager.m_numIns));
 		GETERROR("before FillMultiDense");
-		FillMultiDense<<<dimGridThreadForEachIns, sharedMemSizeEachIns, 0, (*(cudaStream_t*)pStream)>>>(
+		FillMultiDense<<<numofInsToFill, 1, 0, (*(cudaStream_t*)pStream)>>>(
 											  pDevInsValue, pInsStartPos, pDevFeaId, pNumofFea,
 											  pTempDense,
 										  	  manager.m_pSortedUsedFeaIdBag + bagId * bagManager.m_maxNumUsedFeaATree,
@@ -106,10 +103,7 @@ bool bOptimisePred = true;
 	{
 		assert(pLastTree != NULL);
 		if(bOptimisePred == false){
-			dim3 dimGridThreadForEachIns;
-			conf.ComputeBlock(nNumofIns, dimGridThreadForEachIns);
-			int sharedMemSizeEachIns = 1;
-			PredMultiTarget<<<dimGridThreadForEachIns, sharedMemSizeEachIns, 0, (*(cudaStream_t*)pStream)>>>(
+			PredMultiTarget<<<nNumofIns, 1, 0, (*(cudaStream_t*)pStream)>>>(
 												  bagManager.m_pTargetValueEachBag + bagId * bagManager.m_numIns,
 												  nNumofIns, pLastTree,
 												  manager.m_pdDenseInsEachBag + bagId * bagManager.m_maxNumUsedFeaATree * bagManager.m_numIns,

Device/FindSplit/FindFeaCsr.cu

@@ -31,7 +31,6 @@ void DeviceSplitter::FeaFinderAllNode2(void *pStream, int bagId)
 	BagManager bagManager;
 	BagCsrManager csrManager(bagManager.m_numFea, bagManager.m_maxNumSplittable, bagManager.m_numFeaValue);
 	int numofSNode = bagManager.m_curNumofSplitableEachBag_h[bagId];
-	printf("numof sn=%d\n", numofSNode);
 
 	IndexComputer indexComp;
 	indexComp.AllocMem(bagManager.m_numFea, numofSNode, bagManager.m_maxNumSplittable);
@@ -259,13 +258,6 @@ void DeviceSplitter::FeaFinderAllNode2(void *pStream, int bagId)
 	cudaStreamSynchronize((*(cudaStream_t*)pStream));
 	GETERROR("after ComputeGainDense");
 
-	//change the gain of the first feature value to 0
-//	int blockSizeFirstGain;
-//	dim3 dimNumofBlockFirstGain;
-//	conf.ConfKernel(numSeg, blockSizeFirstGain, dimNumofBlockFirstGain);
-//	FirstFeaGain<<<dimNumofBlockFirstGain, blockSizeFirstGain, 0, (*(cudaStream_t*)pStream)>>>(
-//			csrManager.pEachCsrFeaStartPos, numSeg, pGain_d, csrManager.curNumCsr);
-
 	//	cout << "searching" << endl;
 	cudaDeviceSynchronize();
 	clock_t start_search = clock();
@@ -294,7 +286,6 @@ void DeviceSplitter::FeaFinderAllNode2(void *pStream, int bagId)
 					  	  	  	  	  	 bagManager.m_pRChildStatEachBag + bagId * bagManager.m_maxNumSplittable,
 					  	  	  	  	  	 bagManager.m_pLChildStatEachBag + bagId * bagManager.m_maxNumSplittable);
 	cudaStreamSynchronize((*(cudaStream_t*)pStream));
-//	printf("completed splitting node.....................................\n");
 	checkCudaErrors(cudaFree(pMaxGain_d));
 	checkCudaErrors(cudaFree(pMaxGainKey_d));
 }

Device/FindSplit/FindFeaKernel.h

@@ -82,11 +82,6 @@ __global__ void ComputeGainDense(const nodeStat *pSNodeStat, const int *pid2SNPo
     	}
     	else
     		pGainOnEachFeaValue[gTid] = all2Left;
-//    	if(7773118 == gTid)
-//    		printf("gain=%f, totalHess=%f, totalMissHess=%f, last+1=%f, last-1=%f, last=%f, last-2=%f, last=%u\n",
-//    				pGainOnEachFeaValue[gTid], totalHess, totalMissingHess, pHessPrefixSumOnEachFeaValue[lastFvaluePos+1],
-//    				pHessPrefixSumOnEachFeaValue[lastFvaluePos],
-//    				pHessPrefixSumOnEachFeaValue[lastFvaluePos-1], pHessPrefixSumOnEachFeaValue[lastFvaluePos-2], lastFvaluePos);
 	}
 	else{
 		//if the previous fea value is the same as the current fea value, gain is 0 for the current fea value.
@@ -172,10 +167,6 @@ __global__ void FindSplitInfo(const uint *pEachFeaStartPosEachNode, const T *pEa
 	uint lastFvaluePos = segStartPos + segLen - 1;
 	if(key == 0 || (key > 0 && pnKey[key] != pnKey[key - 1])){//first element of the feature
         const real gap = fabs(pDenseFeaValue[key]) + DeviceSplitter::rt_eps;
-        //printf("############## %u v.s. %u; %u\n", bestFeaId, bestFeaId, pPrefixSumHess[key]);
-//        printf("missing %f all to one node: fid=%u, pnKey[%u]=%u != pnKey[%u]=%u, segLen=%u, parentHess=%f, startPos=%u..........................\n",
-//        		pPrefixSumHess[key], bestFeaId, key, pnKey[key], key-1, pnKey[key-1], pEachFeaLenEachNode[segId],
-//        		snNodeStat[snPos].sum_hess, pEachFeaStartPosEachNode[segId]);
 		if(pDefault2Right[key] == true){//all non-missing to left
 			pBestSplitPoint[snPos].m_bDefault2Right = true;
 			pBestSplitPoint[snPos].m_fSplitValue = pDenseFeaValue[lastFvaluePos] + gap;

SharedUtility/initCuda.cu

@@ -63,8 +63,8 @@ int GetMaxMemDevice(int count){
 
 bool InitCUDA(CUcontext &context, char gpuType = 'T')
 {
-//   UseDevice(0, context);
-//	return true;
+   UseDevice(1, context);
+	return true;
 
     int count;
 

auto-test.py

@@ -2,7 +2,7 @@
 fname="stderr.test"
 dataset=["abalone", "cadata", "covetype", "e2006", "higgs", "car ins", 
 		"log1p", "new20", "realsim", "susy", "year pre", "3d"]
-result=["2.26841", "72842", "0.365866", "0.3716", "0.446322", "38.9215",
+result=["2.26841", "72842", "0.348459", "0.371375", "0.446233", "38.9214",
 		"0.37719", "0.784022", "0.705245", "0.376718", "9.34178", "3.42847e+07"]
 new_result=[]
 

gbdt_main.cu

@@ -101,14 +101,18 @@ int main(int argc, char *argv[])
 	printf("max numof used fea is %d\n", maxNumofUsedFeature);
 
 	//decide if want to use csr
-	if(numFeaValue > pow(2, 28) || (numFeaValue > pow(2, 20) && numFea < numIns / 1000) || numFea / 10 < numIns){
+	if(numFeaValue > pow(2, 28) || (numFeaValue > pow(2, 20) && numFea < numIns / 1000) || numFea / 10 < numIns)
 		CsrCompressor::bUseCsr = true;
-cerr << "use CSR" << endl;
-	}
-	else{
+	else
 		CsrCompressor::bUseCsr = false;
-cerr << "use non-CSR" << endl;
-	}
+
+//CsrCompressor::bUseCsr = false;
+
+	if(CsrCompressor::bUseCsr == true)
+		cerr << "use CSR" << endl;
+	else
+		cerr << "use non-CSR" << endl;
+
 
 	//fill the bags
 	BagManager bagManager;