 	__syncthreads();
 	
 	// global thread index
-	const unsigned int t   = blockDim.x * blockIdx.x + threadIdx.x;  
-
+	//const unsigned int t   = blockDim.x * blockIdx.x + threadIdx.x;  
+
+	const unsigned int th_group    = (blockDim.x * blockIdx.x + threadIdx.x)/THREAD_PER_ROW;  
 	//thread group number, 
-	int th_group = t/THREAD_PER_ROW;
-	int th_mod = t%THREAD_PER_ROW;
+	//int th_group = t/THREAD_PER_ROW;
+	//int th_mod = t%THREAD_PER_ROW;
+	//const unsigned int th_mod    = (blockDim.x * blockIdx.x + threadIdx.x)%THREAD_PER_ROW;  
 	//determines the class membership, (first cls1_N threads belongs to first class),0 - first class, 1- second class
 	int th_cls = th_group/cls1_N_aligned;
 	

				//wrong index
				//arIdx = (i*PREFETCH_SIZE+j )*align+rowSliceStart+threadIdx.x;
				//corrected version
				arIdx = (i*PREFETCH_SIZE+j )*align+rowSliceStart+sliceOffset*THREAD_PER_ROW+threadIdx.x%THREAD_PER_ROW;
				preColls[j]=colIdx[arIdx];
				preVals[j]=vals[arIdx];
			}

//		{
//					
//			//index within a subset of two considered class
//			int row = th_cls_offset+th_cls*cls_count[0];
//			
//			results[row]=blockIdx.x;
//			results[row+num_rows]=sliceStartNr;
//			
//			// results[row]=row;
//			// results[row+shClsSum]=row+shClsSum;
//			
//		}
//	}
