 	unsigned int 		_device_sm_version;
 	
 	// Information about our kernel assembly
-	unsigned int 		_kernel_sm_version;
+	unsigned int 		_kernel_ptx_version;
 	cudaFuncAttributes 	_spine_scan_kernel_attrs;
 	
 protected:

			} 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_sm_version);
				if (_keys_only) { 
					orig_max_grid_size *= (_num_elements + (1024 * 1024 * 96) - 1) / (1024 * 1024 * 96);
				} else {

			FlushKernel<void><<<_grid_size, RADIXSORT_THREADS, scan_scatter_attrs.sharedSizeBytes>>>();

	// GF100 and GT200 get the same smem allocation for every kernel launch (pad the reduction/top-level-scan kernels)
	dynamic_smem = (_kernel_ptx_version >= 130) ? scan_scatter_attrs.sharedSizeBytes - reduce_kernel_attrs.sharedSizeBytes : 0;

	RakingReduction<ConvertedKeyType, V, PASS, RADIX_BITS, BIT, PreprocessFunctor> <<<_grid_size, threads, dynamic_smem>>>(
		converted_storage.d_from_alt_storage,
