remove redundant code

Author: zeyiwen

Device/CSR/CsrSplit.cu

@@ -79,9 +79,6 @@ __global__ void fillFvalue(const real *csrFvalue, uint numCsr, const uint *preRo
 		atomicAdd(newCsrFeaLen + (firstPartId + 1) * numFea + feaId, 1);
 		newCsrLen[basePos + numCsrCurPart] = csrLen;
 	}
-//	if(0.369250 == temp){
-//		printf("csr len of 0.369250 is %d and %d; old len=%d, csrId=%d\n", newCsrLen[basePos], newCsrLen[basePos + numCsrCurPart], oldCsrLen[csrId], csrId);
-//	}
 }
 
 __global__ void newCsrLenFvalue(const int *preFvalueInsId, int numFeaValue, const int *pInsId2Nid, int maxNid,
@@ -101,10 +98,7 @@ __global__ void newCsrLenFvalue(const int *preFvalueInsId, int numFeaValue, cons
 		RangeBinarySearch(gTid, eachCsrStart, numCsr, csrId);
 	CONCHECKER(csrId < numCsr);
 	__syncthreads();
-	//first csrId
-//	if(gTid >= numFeaValue && tid == 0){
-//		printf("oh shitt####################################################\n");
-//	}
+
 	if(tid == 0)
 		firstCsrId = csrId;
 	__syncthreads();
@@ -115,30 +109,16 @@ __global__ void newCsrLenFvalue(const int *preFvalueInsId, int numFeaValue, cons
 		int insId = preFvalueInsId[gTid];//insId is not -1, as preFvalueInsId is dense.
 		int pid = pInsId2Nid[insId] - maxNid - 1;//mapping to new node
 
-//		if(csrId == 2809992 && csrFvalue[csrId] == 0.369250){
-//			printf("gpu pid=%d, insId=%d, csrfvalue=%f, left_fv=%f, right_fv=%f, gTid=%d, blkid=%d\n",
-//					pid, insId, csrFvalue[csrId], csrFvalue[csrId - 1], csrFvalue[csrId + 1], gTid, blockIdx.y * gridDim.x + blockIdx.x);
-//		}
-
 		csrId2Pid[csrId] = pid < 0 ? LARGE_1B_UCHAR : pid;
 		if(pid >= 0 && pid % 2 == 0){//not leaf node and it's first part.
 			uint counterOffset = csrId - firstCsrId;
-//			if(305030 == csrId || 478798 == csrId){
-//				printf("first csr id=%d, csrId=%d\n", firstCsrId, csrId);
-//			}
+
 			__threadfence();
 			uint orgValue = atomicAdd(csrCounter + counterOffset, 1);
-//			if(csrFvalue[csrId] == 0.369250){
-//				printf("gpu pid=%d, insId=%d, csrfvalue=%f, firstCsrId=%d, cntOffset=%d, orgValue=%d, cnt=%d\n",
-//						pid, insId, csrFvalue[csrId], firstCsrId, counterOffset, orgValue, csrCounter[counterOffset]);
-//			}
 		}
 	}
 
 	__syncthreads();
-//	if(blockIdx.y * gridDim.x + blockIdx.x == 2216293 && tid == 95){
-//		printf("###################################################################### cnt=%d\n", csrCounter[tid]);
-//	}
 	//compute len of each csr
 	if(csrCounter[tid] > 0){
 		uint numCsrCurPart;
@@ -149,9 +129,6 @@ __global__ void newCsrLenFvalue(const int *preFvalueInsId, int numFeaValue, cons
 					  numCsrPrePartsAhead, posInPart, numCsrCurPart, feaId);
 
 		atomicAdd(csrNewLen + numCsrPrePartsAhead * 2 + posInPart, csrCounter[tid]);
-//		if(blockIdx.y * gridDim.x + blockIdx.x == 2216293 && tid == 95){
-//			printf("csrNewLen=%d, numcsrPrePartAhead=%d, posInPart=%d, cnt=%d\n", csrNewLen[numCsrPrePartsAhead * 2 + posInPart], numCsrPrePartsAhead, posInPart, csrCounter[tid]);
-//		}
 	}
 }
 

Device/FindSplit/FindFeaCsr.cu

@@ -150,8 +150,6 @@ void AfterCompression(GBDTGPUMemManager &manager, BagCsrManager &csrManager, Bag
 	int numSeg = bagManager.m_numFea * numofSNode;
 	//construct keys for exclusive scan
 	checkCudaErrors(cudaMemset(csrManager.getMutableCsrKey(), -1, sizeof(uint) * csrManager.curNumCsr));
-//	checkCudaErrors(cudaMemset(pCSRMultableKey, -1, sizeof(uint) * csrManager.curNumCsr));
-	printf("done constructing key... number of segments is %d\n", numSeg);
 
 	//set keys by GPU
 	uint maxSegLen = 0;
@@ -339,10 +337,6 @@ cudaDeviceSynchronize();
 		int blockSizeFillFvalue;
 		dim3 dimNumBlockToFillFvalue;
 		conf.ConfKernel(csrManager.curNumCsr, blockSizeFillFvalue, dimNumBlockToFillFvalue);
-//fid hess sum
-uint *hess_cnt_d;
-checkCudaErrors(cudaMalloc((void**)&hess_cnt_d, sizeof(uint) * bagManager.m_numFea));
-checkCudaErrors(cudaMemset(hess_cnt_d, 0, sizeof(uint) * bagManager.m_numFea));
 		fillFvalue<<<dimNumBlockToFillFvalue, blockSizeFillFvalue>>>(csrManager.getCsrFvalue(), csrManager.curNumCsr, csrManager.pEachCsrFeaStartPos,
 				   bagManager.m_pPreNumSN_h[0], bagManager.m_numFea, csrManager.getCsrKey(), pOldCsrLen_d, pCsrId2Pid,
 //				   bagManager.m_pPreNumSN_h[bagId], bagManager.m_numFea, pCSRKey, pOldCsrLen_d, pCsrId2Pid,
@@ -426,7 +420,6 @@ checkCudaErrors(cudaFree(pCsrMarker));
 								  csrManager.pEachNodeSizeInCsr, csrManager.pEachCsrNodeStartPos, csrManager.getMutableCsrFvalue(), csrManager.getMutableCsrLen());
 	}
 	//need to compute for every new tree
-	printf("reserve memory\n");
 	if(indexComp.histogram_d.reservedSize < csrManager.curNumCsr * (2 + 2 * sizeof(real)/sizeof(uint))){//make sure enough memory for reuse
 		printf("reallocate memory for histogram (sn=%u): %u v.s. %u.......\n", numofSNode, indexComp.histogram_d.reservedSize,
 				csrManager.curNumCsr * (2 + 2 * sizeof(real)/sizeof(uint)));
@@ -448,7 +441,6 @@ checkCudaErrors(cudaFree(pCsrMarker));
 		pCsrStartPos_d = (uint*)indexComp.partitionMarker.addr;
 	printf("comp gd and hess\n");
 
-
 	ComputeGDHess<<<dimNumofBlockForGD, blockSizeForGD, sharedMemSizeForGD>>>(csrManager.getCsrLen(), pCsrStartPos_d,
 			bagManager.m_pInsGradEachBag,
 			bagManager.m_pInsHessEachBag,
@@ -457,7 +449,6 @@ checkCudaErrors(cudaFree(pCsrMarker));
 	GETERROR("after ComputeGD");
 	clock_t csr_len_end = clock();
 	total_csr_len_t += (csr_len_end - csr_len_t);
-	printf("done comp gd and hess\n");
 }
 
 void DeviceSplitter::FeaFinderAllNode2(void *pStream, int bagId)