 				get_hermitian100_tt_fp16<<<batch_size, 64, SCAN_BATCH * f/2*sizeof(float2)>>>
 					(batch_offset, (half2*) xx, cscColIndex, cscRowIndex, lambda, n, f, XT);
 				#else
-				get_hermitian100<<<batch_size, 64, SCAN_BATCH * f/2*sizeof(float2)>>>
+				get_hermitian100<<<batch_size, 55, SCAN_BATCH * f/2*sizeof(float2)>>>
 					(batch_offset, (float2*)xx, cscColIndex, cscRowIndex, lambda, n, f, XT);
 				#endif
 			}

			two layers: warp divergence unless we split at 32
			require 32 >= SCAN_BATCH
			*/

			//int index = threadIdx.x - (threadIdx.x/32)*32;	//0 to 31;
			int index = threadIdx.x % 32;
			if(threadIdx.x < 2*32 && index < SCAN_BATCH){
				int anchor = start + iter*SCAN_BATCH + index;
				if(anchor < end){
					//IMPORTANT: for loop has constant and identical start and end
					int col = csrColIndex[anchor];
					if(threadIdx.x < 32){
						for (int k = 0; k < 50; k += 2){
							theta.x = __ldg(&thetaT[ F * col + k]);
							theta.y = __ldg(&thetaT[ F * col + k+1]);
							thetaTemp[index * F/2 + k/2] = theta;
						}
					}
					else {
						for (int k = 0; k < 50; k += 2){
							theta.x = __ldg(&thetaT[ F * col + k + 50]);
							theta.y = __ldg(&thetaT[ F * col + k + 51]);
							thetaTemp[index * F/2 + k/2 + 25] = theta;
						}
					}
				}
				//must be the last iteration; no need to check
				//else
				//	memset(&thetaTemp[index*F/2], 0, F*sizeof(float));		
			}
			
/*			//issue: not coalesced access to csrColIndex
			if(threadIdx.x < F && threadIdx.x%2 == 0){
				for(int k = 0; k< SCAN_BATCH; k++){

}

__global__ void
__launch_bounds__(55)
get_hermitian100(const int batch_offset, float2* tt,
		const int* csrRowIndex, const int* csrColIndex, const float lambda, const int m, const int F,
		const float* __restrict__ thetaT) {
