 	//
 	
 	// GF100 and GT200 get the same smem allocation for every kernel launch (pad the reduction/top-level-scan kernels)
-	dynamic_smem = (_kernel_sm_version >= 130) ? scan_scatter_attrs.sharedSizeBytes - _spine_scan_kernel_attrs.sharedSizeBytes : 0;
+	dynamic_smem = (_kernel_ptx_version >= 130) ? scan_scatter_attrs.sharedSizeBytes - _spine_scan_kernel_attrs.sharedSizeBytes : 0;
 	
 	SrtsScanSpine<void><<<_grid_size, RADIXSORT_SPINE_THREADS, dynamic_smem>>>(
 		converted_storage.d_spine,

	
	if (RADIXSORT_DEBUG) {
		
		printf("Bottom-level reduction & scan kernels:\n\tgrid_size: %d, \n\tthreads: %d, \n\tcycle_elements: %d, \n\tnum_big_blocks: %d, \n\tbig_block_elements: %d, \n\tnormal_block_elements: %d\n\textra_elements_last_block: %d\n\n",
			_grid_size, RADIXSORT_THREADS, _cycle_elements, _work_decomposition.num_big_blocks, _work_decomposition.big_block_elements, _work_decomposition.normal_block_elements, _work_decomposition.extra_elements_last_block);
		printf("Top-level spine scan:\n\tgrid_size: %d, \n\tthreads: %d, \n\tspine_block_elements: %d\n\n", 

			} else if (_device_sm_version < 200) {
				
				// GT200 (has some kind of TLB or icache drama)
				unsigned int orig_max_grid_size = _device_props.multiProcessorCount * RADIXSORT_SCAN_SCATTER_CTA_OCCUPANCY(_kernel_ptx_version);
				if (_keys_only) { 
					orig_max_grid_size *= (_num_elements + (1024 * 1024 * 96) - 1) / (1024 * 1024 * 96);
				} else {
