 
 				// Get kernel attributes
 				cudaFuncAttributes upsweep_kernel_attrs, spine_kernel_attrs;
-				if (retval = B40CPerror(cudaFuncGetAttributes(&upsweep_kernel_attrs, UpsweepReductionKernel<ReductionConfig::Upsweep>),
+				if (retval = B40CPerror(cudaFuncGetAttributes(&upsweep_kernel_attrs, UpsweepReductionKernel<typename ReductionConfig::Upsweep>),
 					"ReductionEnactor cudaFuncGetAttributes upsweep_kernel_attrs failed", __FILE__, __LINE__)) break;
-				if (retval = B40CPerror(cudaFuncGetAttributes(&spine_kernel_attrs, SpineReductionKernel<ReductionConfig::Spine>),
+				if (retval = B40CPerror(cudaFuncGetAttributes(&spine_kernel_attrs, SpineReductionKernel<typename ReductionConfig::Spine>),
 					"ReductionEnactor cudaFuncGetAttributes spine_kernel_attrs failed", __FILE__, __LINE__)) break;
 
 				int max_static_smem = B40C_MAX(upsweep_kernel_attrs.sharedSizeBytes, spine_kernel_attrs.sharedSizeBytes);


			// No need to scan the spine if there's only one CTA in the upsweep grid
			int dynamic_smem = 0;
			UpsweepReductionKernel<ReductionConfig::Upsweep>
					<<<work.grid_size, ReductionConfig::Upsweep::THREADS, dynamic_smem>>>(
				d_src, d_dest, d_work_progress, work, progress_selector);
			if (DEBUG && (retval = B40CPerror(cudaThreadSynchronize(), "ReductionEnactor UpsweepReductionKernel failed ", __FILE__, __LINE__))) break;

			}

			// Upsweep reduction into spine
			UpsweepReductionKernel<typename ReductionConfig::Upsweep>
					<<<grid_size[0], ReductionConfig::Upsweep::THREADS, dynamic_smem[0]>>>(
				d_src, (T*) d_spine, d_work_progress, work, progress_selector);
			if (DEBUG && (retval = B40CPerror(cudaThreadSynchronize(), "ReductionEnactor UpsweepReductionKernel failed ", __FILE__, __LINE__))) break;

			// Spine reduction
			SpineReductionKernel<typename ReductionConfig::Spine>
					<<<grid_size[1], ReductionConfig::Spine::THREADS, dynamic_smem[1]>>>(
				(T*) d_spine, d_dest, spine_elements);
			if (DEBUG && (retval = B40CPerror(cudaThreadSynchronize(), "ReductionEnactor SpineReductionKernel failed ", __FILE__, __LINE__))) break;
