 	int *deviceTADShapeInfoZ = reinterpret_cast<int *>(extraPointers[12]);
 	int *deviceTADOffsetsZ = reinterpret_cast<int *>(extraPointers[13]);
 
-
-
 	if (debug && verbose)
 		printf("F3 opNum:[%i]\n", opNum);
 
-	//dim3 launchDims = getOptimalLaunchParameters<float>(&extraPointers[0], funcAttributes[12], deviceProperties[getDeviceId(extraPointers[2])]);
-
-
 	dim3 launchDims = getReduceLaunchParams(getDeviceId(extraPointers[2]), hostXShapeInfo, hostTADShapeInfo, funcAttributes[12], 1, sizeof(float), 0);
 
+	// this macro builds bunch of IF/ELSE selectors for kernel launch
     DISPATCH_SIMPLE(broadcastSimple, float, PARAMS(x, xShapeInfo, y, yShapeInfo, result, resultShapeInfo, dimension, dimensionLength, deviceTADShapeInfo, deviceTADOffsets, deviceTADShapeInfoZ, deviceTADOffsetsZ), OPS_A(BROADCAST_OPS))
 
 /*

		float16 *extraParams){
	if (debug && verbose)
		printf("H1 opNum:[%i]\n", opNum);

//	dim3 launchDims = getOptimalLaunchParameters<float>(&extraPointers[0], funcAttributes[13], deviceProperties[getDeviceId(extraPointers[2])]);

	cudaStream_t *stream = reinterpret_cast<cudaStream_t *>(&extraPointers[1]);


	if (debug && verbose)
		printf("H3 opNum:[%i]\n", opNum);

	dim3 launchDims = getReduceLaunchParams(getDeviceId(extraPointers[2]), hostXShapeInfo, hostTADShapeInfo, funcAttributes[12], 1, sizeof(float16), 0);

	// this macro builds bunch of IF/ELSE selectors for kernel launch
    DISPATCH_SIMPLE(broadcastSimple, float16, PARAMS(x, xShapeInfo, y, yShapeInfo, result, resultShapeInfo, dimension, dimensionLength, deviceTADShapeInfo, deviceTADOffsets, deviceTADShapeInfoZ, deviceTADOffsetsZ), OPS_A(BROADCAST_OPS))

	if (debug)
		checkCudaErrors(cudaStreamSynchronize(*stream));
