 	// Get SM version of compiled kernel assembly
 	//
 	cudaFuncGetAttributes(&_spine_scan_kernel_attrs, SrtsScanSpine<void>);
-	_kernel_sm_version = _spine_scan_kernel_attrs.binaryVersion * 10;
+	_kernel_ptx_version = _spine_scan_kernel_attrs.ptxVersion * 10;
 	
 
 	//

	//
	
	// 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;
	
	SrtsScanSpine<void><<<_grid_size, RADIXSORT_SPINE_THREADS, dynamic_smem>>>(
		converted_storage.d_spine,

			} 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 {
