 	_passes								= passes;
 	_num_elements 						= num_elements;
 	_keys_only 							= IsKeysOnly<V>();
-	_cycle_elements 					= RADIXSORT_CYCLE_ELEMENTS(_kernel_sm_version , ConvertedKeyType, V);
+	_cycle_elements 					= RADIXSORT_CYCLE_ELEMENTS(_kernel_ptx_version , ConvertedKeyType, V);
 	_grid_size 							= GridSize(max_grid_size);
 	_swizzle_pointers_for_odd_passes	= swizzle_pointers_for_odd_passes;
 	

	//
	
	// 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,

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