 							checkCudaErrors(cudaStreamSynchronize(*stream));
 
 						// at this point, all IMax indexes are gathered, and we execute
-						fillDimensionalIsMaxFloat<<<128, 64, funcAttributes[36].sharedSizeBytes, *stream>>>(special, hostYShapeInfo, result, resultShapeInfo, tadMaxShapeInfo, dimension, 1, tadMaxOffsets );
+						fillDimensionalIsMaxFloat<<<512, 32, funcAttributes[36].sharedSizeBytes, *stream>>>(special, hostYShapeInfo, result, resultShapeInfo, tadMaxShapeInfo, dimension, 1, tadMaxOffsets );
 
 
 						checkCudaErrors(cudaStreamSynchronize(*stream));

}

void NativeOps::setGridLimit(int gridSize) {
	if (gridSize > 1024)
		gridSize = 1024;
	if (gridSize < 1)
		gridSize = 1;
	blockLimit = gridSize;

	if (verbose && launchDims.x == 1)
		printf("AF2 opNum:[%i]\n", opNum);

	indexReduceFloat<<<launchDims.x, launchDims.y,launchDims.z, *stream>>>(
			opNum,
			x,
			xShapeInfo, shape::rank(hostXShapeInfo),
