 				cudacall(cudaFree(XT_fp16));
 				#elif defined(CUMF_XX_FP16)
 				get_hermitian100_tt_fp16<<<batch_size, 64, SCAN_BATCH * f/2*sizeof(float2)>>>
-					(batch_offset, (half2*) xx, cscColIndex, cscRowIndex, lambda, n, f, XT);
+					(batch_offset, (half2*) xx, cscColIndex, cscRowIndex, lambda, n, f, (float2*)XT);
 				#else
 				get_hermitian100<<<batch_size, 64, SCAN_BATCH * f/2*sizeof(float2)>>>
 					(batch_offset, (float2*)xx, cscColIndex, cscRowIndex, lambda, n, f, (float2*)XT);

__launch_bounds__(64, 6)
get_hermitian100_tt_fp16(const int batch_offset, half2* tt,
		const int* csrRowIndex, const int* csrColIndex, const float lambda, const int m, const int F,
		const float* __restrict__ thetaT) {
	extern __shared__ float2 thetaTemp[];
	int row = blockIdx.x + batch_offset;
	if (row < m) {

		}
		//iteration: copy gmem-->smem; aggregate smem-->register
		for (int iter = 0; iter < iterations; iter ++){
			//copy texture --> smem, and sync
			/*
			This is the fastest implementation
			thetaT is NOT coalesced loaded but cached by L1 and L2
			faster than coalesced version (see the next paragraph commented out) 
			because it concurrently load multiple thetaT columns
			two threads per theta column, e.g., threads 0 & 1 for theta[0], threads 2 & 3 for theta[1]
			require: blockDim.x (64) >= 2*SCAN_BATCH
			*/
///* 
			if(threadIdx.x < 2*SCAN_BATCH){
				int anchor = start + iter*SCAN_BATCH + threadIdx.x/2;
				if(anchor < end){
					int col = csrColIndex[anchor];
					//IMPORTANT: for loop has constant and identical start and end
					for (int k = 0; k < 50; k += 2)
						//thetaTemp[threadIdx.x*F/4 + k/2] =__ldg(&thetaT[ F/2 * col + threadIdx.x%2*F/4 + k/2]);
						thetaTemp[threadIdx.x*F/4 + k/2] = thetaT[ F/2 * col + threadIdx.x%2*F/4 + k/2];
				}
			}
//*/
			__syncthreads();

			//tile: 10*10
			if(threadIdx.x < 55){
				if(iter < iterations - 1){
					for(int k = 0; k < SCAN_BATCH; k++)
						accumulate_in_registers();
				}
				else{
					for(int k = 0; k < end - start - iter*SCAN_BATCH; k++)
						accumulate_in_registers();
				}
				
			}
		}
		//end of iteration in copying from smem and aggregating in register
