From 0ca14d740fcf76579d17f5b573f3f003f04592bc Mon Sep 17 00:00:00 2001 From: Andrzej Janik Date: Mon, 13 Dec 2021 22:25:26 +0100 Subject: Better reporting of unrecognized tokens --- hip_runtime-sys/build.rs | 18 ++- hip_runtime-sys/lib/amdhip64.def | 308 +++++++++++++++++++++++++++++++++++++++ hip_runtime-sys/lib/amdhip64.lib | Bin 0 -> 71480 bytes ptx/src/lib.rs | 112 ++++++++++++++ ptx/src/ptx.lalrpop | 11 +- zluda_dump/src/trace.rs | 39 +++-- 6 files changed, 461 insertions(+), 27 deletions(-) create mode 100644 hip_runtime-sys/lib/amdhip64.def create mode 100644 hip_runtime-sys/lib/amdhip64.lib diff --git a/hip_runtime-sys/build.rs b/hip_runtime-sys/build.rs index 9c3d9b4..3bc1250 100644 --- a/hip_runtime-sys/build.rs +++ b/hip_runtime-sys/build.rs @@ -1,9 +1,21 @@ use std::env::VarError; +use std::{env, path::PathBuf}; fn main() -> Result<(), VarError> { println!("cargo:rustc-link-lib=dylib=amdhip64"); - //println!("cargo:rustc-link-search=native=/opt/rocm/lib/"); - println!("cargo:rustc-link-search=native=/home/ubuntu/hipamd/build/lib"); - println!("cargo:rustc-link-search=native=/home/vosen/hipamd/build/lib"); + if cfg!(windows) { + let env = env::var("CARGO_CFG_TARGET_ENV")?; + if env == "msvc" { + let mut path = PathBuf::from(env::var("CARGO_MANIFEST_DIR")?); + path.push("lib"); + println!("cargo:rustc-link-search=native={}", path.display()); + } else { + println!("cargo:rustc-link-search=native=C:\\Windows\\System32"); + }; + } else { + //println!("cargo:rustc-link-search=native=/opt/rocm/lib/"); + println!("cargo:rustc-link-search=native=/home/ubuntu/hipamd/build/lib"); + println!("cargo:rustc-link-search=native=/home/vosen/hipamd/build/lib"); + } Ok(()) } diff --git a/hip_runtime-sys/lib/amdhip64.def b/hip_runtime-sys/lib/amdhip64.def new file mode 100644 index 0000000..515d5c6 --- /dev/null +++ b/hip_runtime-sys/lib/amdhip64.def @@ -0,0 +1,308 @@ +LIBRARY AMDHIP64 +EXPORTS +__hipPopCallConfiguration +__hipPushCallConfiguration +__hipRegisterFatBinary +__hipRegisterFunction +__hipRegisterManagedVar +__hipRegisterSurface +__hipRegisterTexture +__hipRegisterVar +__hipUnregisterFatBinary +hipApiName +hipArray3DCreate +hipArrayCreate +hipArrayDestroy +hipBindTexture +hipBindTexture2D +hipBindTextureToArray +hipBindTextureToMipmappedArray +hipChooseDevice +hipConfigureCall +hipCreateChannelDesc +hipCreateSurfaceObject +hipCreateTextureObject +hipCtxCreate +hipCtxDestroy +hipCtxDisablePeerAccess +hipCtxEnablePeerAccess +hipCtxGetApiVersion +hipCtxGetCacheConfig +hipCtxGetCurrent +hipCtxGetDevice +hipCtxGetFlags +hipCtxGetSharedMemConfig +hipCtxPopCurrent +hipCtxPushCurrent +hipCtxSetCacheConfig +hipCtxSetCurrent +hipCtxSetSharedMemConfig +hipCtxSynchronize +hipDestroyExternalMemory +hipDestroyExternalSemaphore +hipDestroySurfaceObject +hipDestroyTextureObject +hipDeviceCanAccessPeer +hipDeviceComputeCapability +hipDeviceDisablePeerAccess +hipDeviceEnablePeerAccess +hipDeviceGet +hipDeviceGetAttribute +hipDeviceGetByPCIBusId +hipDeviceGetCacheConfig +hipDeviceGetLimit +hipDeviceGetName +hipDeviceGetP2PAttribute +hipDeviceGetPCIBusId +hipDeviceGetSharedMemConfig +hipDeviceGetStreamPriorityRange +hipDevicePrimaryCtxGetState +hipDevicePrimaryCtxRelease +hipDevicePrimaryCtxReset +hipDevicePrimaryCtxRetain +hipDevicePrimaryCtxSetFlags +hipDeviceReset +hipDeviceSetCacheConfig +hipDeviceSetSharedMemConfig +hipDeviceSynchronize +hipDeviceTotalMem +hipDriverGetVersion +hipDrvMemcpy2DUnaligned +hipDrvMemcpy3D +hipDrvMemcpy3DAsync +hipEnableActivityCallback +hipEventCreate +hipEventCreateWithFlags +hipEventDestroy +hipEventElapsedTime +hipEventQuery +hipEventRecord +hipEventSynchronize +hipExtGetLinkTypeAndHopCount +hipExtLaunchKernel +hipExtLaunchMultiKernelMultiDevice +hipExtMallocWithFlags +hipExtModuleLaunchKernel +hipExtStreamCreateWithCUMask +hipExtStreamGetCUMask +hipExternalMemoryGetMappedBuffer +hipFree +hipFreeArray +hipFreeHost +hipFreeMipmappedArray +hipFuncGetAttribute +hipFuncGetAttributes +hipFuncSetAttribute +hipFuncSetCacheConfig +hipFuncSetSharedMemConfig +hipGLGetDevices +hipGetChannelDesc +hipGetCmdName +hipGetDevice +hipGetDeviceCount +hipGetDeviceFlags +hipGetDeviceProperties +hipGetErrorName +hipGetErrorString +hipGetLastError +hipGetMipmappedArrayLevel +hipGetSymbolAddress +hipGetSymbolSize +hipGetTextureAlignmentOffset +hipGetTextureObjectResourceDesc +hipGetTextureObjectResourceViewDesc +hipGetTextureObjectTextureDesc +hipGetTextureReference +hipGraphAddDependencies +hipGraphAddEmptyNode +hipGraphAddKernelNode +hipGraphAddMemcpyNode +hipGraphAddMemcpyNode1D +hipGraphAddMemsetNode +hipGraphCreate +hipGraphDestroy +hipGraphExecDestroy +hipGraphExecKernelNodeSetParams +hipGraphGetNodes +hipGraphGetRootNodes +hipGraphInstantiate +hipGraphKernelNodeGetParams +hipGraphKernelNodeSetParams +hipGraphLaunch +hipGraphMemcpyNodeGetParams +hipGraphMemcpyNodeSetParams +hipGraphMemsetNodeGetParams +hipGraphMemsetNodeSetParams +hipGraphicsGLRegisterBuffer +hipGraphicsMapResources +hipGraphicsResourceGetMappedPointer +hipGraphicsUnmapResources +hipGraphicsUnregisterResource +hipHccModuleLaunchKernel +hipHostAlloc +hipHostFree +hipHostGetDevicePointer +hipHostGetFlags +hipHostMalloc +hipHostRegister +hipHostUnregister +hipImportExternalMemory +hipImportExternalSemaphore +hipInit +hipInitActivityCallback +hipIpcCloseMemHandle +hipIpcGetEventHandle +hipIpcGetMemHandle +hipIpcOpenEventHandle +hipIpcOpenMemHandle +hipKernelNameRef +hipLaunchByPtr +hipLaunchCooperativeKernel +hipLaunchCooperativeKernelMultiDevice +hipLaunchKernel +hipMalloc +hipMalloc3D +hipMalloc3DArray +hipMallocArray +hipMallocHost +hipMallocManaged +hipMallocMipmappedArray +hipMallocPitch +hipMemAdvise +hipMemAllocHost +hipMemAllocPitch +hipMemGetAddressRange +hipMemGetInfo +hipMemPrefetchAsync +hipMemPtrGetInfo +hipMemRangeGetAttribute +hipMemRangeGetAttributes +hipMemcpy +hipMemcpy2D +hipMemcpy2DAsync +hipMemcpy2DFromArray +hipMemcpy2DFromArrayAsync +hipMemcpy2DToArray +hipMemcpy2DToArrayAsync +hipMemcpy3D +hipMemcpy3DAsync +hipMemcpyAsync +hipMemcpyAtoH +hipMemcpyDtoD +hipMemcpyDtoDAsync +hipMemcpyDtoH +hipMemcpyDtoHAsync +hipMemcpyFromArray +hipMemcpyFromSymbol +hipMemcpyFromSymbolAsync +hipMemcpyHtoA +hipMemcpyHtoD +hipMemcpyHtoDAsync +hipMemcpyParam2D +hipMemcpyParam2DAsync +hipMemcpyPeer +hipMemcpyPeerAsync +hipMemcpyToArray +hipMemcpyToSymbol +hipMemcpyToSymbolAsync +hipMemcpyWithStream +hipMemset +hipMemset2D +hipMemset2DAsync +hipMemset3D +hipMemset3DAsync +hipMemsetAsync +hipMemsetD16 +hipMemsetD16Async +hipMemsetD32 +hipMemsetD32Async +hipMemsetD8 +hipMemsetD8Async +hipMipmappedArrayCreate +hipMipmappedArrayDestroy +hipMipmappedArrayGetLevel +hipModuleGetFunction +hipModuleGetGlobal +hipModuleGetTexRef +hipModuleLaunchKernel +hipModuleLaunchKernelExt +hipModuleLoad +hipModuleLoadData +hipModuleLoadDataEx +hipModuleOccupancyMaxActiveBlocksPerMultiprocessor +hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags +hipModuleOccupancyMaxPotentialBlockSize +hipModuleOccupancyMaxPotentialBlockSizeWithFlags +hipModuleUnload +hipOccupancyMaxActiveBlocksPerMultiprocessor +hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags +hipOccupancyMaxPotentialBlockSize +hipPeekAtLastError +hipPointerGetAttributes +hipProfilerStart +hipProfilerStop +hipRegisterActivityCallback +hipRegisterApiCallback +hipRemoveActivityCallback +hipRemoveApiCallback +hipRuntimeGetVersion +hipSetDevice +hipSetDeviceFlags +hipSetupArgument +hipSignalExternalSemaphoresAsync +hipStreamAddCallback +hipStreamAttachMemAsync +hipStreamBeginCapture +hipStreamCreate +hipStreamCreateWithFlags +hipStreamCreateWithPriority +hipStreamDestroy +hipStreamEndCapture +hipStreamGetFlags +hipStreamGetPriority +hipStreamIsCapturing +hipStreamQuery +hipStreamSynchronize +hipStreamWaitEvent +hipTexObjectCreate +hipTexObjectDestroy +hipTexObjectGetResourceDesc +hipTexObjectGetResourceViewDesc +hipTexObjectGetTextureDesc +hipTexRefGetAddress +hipTexRefGetAddressMode +hipTexRefGetArray +hipTexRefGetBorderColor +hipTexRefGetFilterMode +hipTexRefGetFlags +hipTexRefGetFormat +hipTexRefGetMaxAnisotropy +hipTexRefGetMipmapFilterMode +hipTexRefGetMipmapLevelBias +hipTexRefGetMipmapLevelClamp +hipTexRefGetMipmappedArray +hipTexRefSetAddress +hipTexRefSetAddress2D +hipTexRefSetAddressMode +hipTexRefSetArray +hipTexRefSetBorderColor +hipTexRefSetFilterMode +hipTexRefSetFlags +hipTexRefSetFormat +hipTexRefSetMaxAnisotropy +hipTexRefSetMipmapFilterMode +hipTexRefSetMipmapLevelBias +hipTexRefSetMipmapLevelClamp +hipTexRefSetMipmappedArray +hipUnbindTexture +hipWaitExternalSemaphoresAsync +hiprtcAddNameExpression +hiprtcCompileProgram +hiprtcCreateProgram +hiprtcDestroyProgram +hiprtcGetCode +hiprtcGetCodeSize +hiprtcGetErrorString +hiprtcGetLoweredName +hiprtcGetProgramLog +hiprtcGetProgramLogSize diff --git a/hip_runtime-sys/lib/amdhip64.lib b/hip_runtime-sys/lib/amdhip64.lib new file mode 100644 index 0000000..1d7df96 Binary files /dev/null and b/hip_runtime-sys/lib/amdhip64.lib differ diff --git a/ptx/src/lib.rs b/ptx/src/lib.rs index 4ade4e8..db9fc23 100644 --- a/ptx/src/lib.rs +++ b/ptx/src/lib.rs @@ -28,6 +28,8 @@ pub mod ast; mod test; mod translate; +use std::fmt; + pub use crate::ptx::ModuleParser; pub use lalrpop_util::lexer::Token; pub use lalrpop_util::ParseError; @@ -36,6 +38,86 @@ pub use translate::to_spirv_module; pub use translate::KernelInfo; pub use translate::TranslateError; +pub trait ModuleParserExt { + fn parse_checked<'input>( + txt: &'input str, + ) -> Result, Vec, ast::PtxError>>>; + + // Returned AST might be malformed. Some users, like logger, want to look at + // malformed AST to record information - list of kernels or such + fn parse_unchecked<'input>( + txt: &'input str, + ) -> ( + ast::Module<'input>, + Vec, ast::PtxError>>, + ); +} + +impl ModuleParserExt for ModuleParser { + fn parse_checked<'input>( + txt: &'input str, + ) -> Result, Vec, ast::PtxError>>> { + let mut errors = Vec::new(); + let maybe_ast = ptx::ModuleParser::new().parse(&mut errors, txt); + match (&*errors, maybe_ast) { + (&[], Ok(ast)) => Ok(ast), + (_, Err(unrecoverable)) => { + errors.push(unrecoverable); + Err(errors) + } + (_, Ok(_)) => Err(errors), + } + } + + fn parse_unchecked<'input>( + txt: &'input str, + ) -> ( + ast::Module<'input>, + Vec, ast::PtxError>>, + ) { + let mut errors = Vec::new(); + let maybe_ast = ptx::ModuleParser::new().parse(&mut errors, txt); + let ast = match maybe_ast { + Ok(ast) => ast, + Err(unrecoverable_err) => { + errors.push(unrecoverable_err); + ast::Module { + version: (0, 0), + directives: Vec::new(), + } + } + }; + (ast, errors) + } +} + +pub struct DisplayParseError<'a, Loc, Tok, Err>(pub &'a str, pub &'a ParseError); + +impl<'a, Loc, Tok, Err> fmt::Display for DisplayParseError<'a, Loc, Tok, Err> +where + Loc: fmt::Display + Into + Copy, + Tok: fmt::Display, + Err: fmt::Display, +{ + fn fmt(&self, f: &mut fmt::Formatter) -> fmt::Result { + match self.1 { + ParseError::UnrecognizedToken { + token: (start, token, end), + .. + } => { + let full_instruction = + unsafe { self.0.get_unchecked((*start).into()..(*end).into()) }; + write!( + f, + "`{}` unrecognized token `{}` found at {}:{}", + full_instruction, token, start, end + ) + } + _ => self.fmt(f), + } + } +} + pub(crate) fn without_none(x: Vec>) -> Vec { x.into_iter().filter_map(|x| x).collect() } @@ -53,3 +135,33 @@ pub(crate) fn vector_index<'input>( }), } } + +#[cfg(test)] +mod tests { + use crate::{DisplayParseError, ModuleParser, ModuleParserExt}; + + #[test] + fn error_report_unknown_instructions() { + let module = r#" + .version 6.5 + .target sm_30 + .address_size 64 + + .visible .entry add( + .param .u64 input, + ) + { + .reg .u64 x; + does_not_exist.u64 x, x; + ret; + }"#; + let errors = match ModuleParser::parse_checked(module) { + Err(e) => e, + Ok(_) => panic!(), + }; + assert_eq!(errors.len(), 1); + let reporter = DisplayParseError(module, &errors[0]); + let build_log_string = format!("{}", reporter); + assert!(build_log_string.contains("does_not_exist")); + } +} diff --git a/ptx/src/ptx.lalrpop b/ptx/src/ptx.lalrpop index 181a727..f6fd9cd 100644 --- a/ptx/src/ptx.lalrpop +++ b/ptx/src/ptx.lalrpop @@ -526,8 +526,15 @@ Statement: Option>> = { ";" => Some(ast::Statement::Instruction(p, i)), PragmaStatement => None, "{" "}" => Some(ast::Statement::Block(without_none(s))), - ! ";" => { - let (err, _) = (<>); + @L ! ";" @R => { + let (start, mut err, _, end) = (<>); + // TODO: report this error more generally, perhaps in user error? + err.error = match err.error { + ParseError::UnrecognizedToken { token: (_, token, _), expected } => { + ParseError::UnrecognizedToken { token: (start, token, end), expected } + } + e => e + }; errors.push(err.error); None } diff --git a/zluda_dump/src/trace.rs b/zluda_dump/src/trace.rs index 3bdf807..eac6bbd 100644 --- a/zluda_dump/src/trace.rs +++ b/zluda_dump/src/trace.rs @@ -1,4 +1,5 @@ use ptx::{ast::PtxError, Token}; +use ptx::{DisplayParseError, ModuleParserExt}; use crate::{cuda::CUmodule, dark_api, log, Settings}; use std::{ @@ -170,24 +171,18 @@ impl StateTracker { submodule_index: Option, module_text: &str, ) { - let mut errors = Vec::new(); - let ast = ptx::ModuleParser::new().parse(&mut errors, module_text); - let ast = match (&*errors, ast) { - (&[], Ok(ast)) => ast, - (err_vec, res) => { - fn_logger.log(log::LogEntry::ModuleParsingError( - DumpWriter::get_file_name(module_index, version, submodule_index, "log"), - )); - fn_logger.log_io_error(self.writer.save_module_error_log( - module_index, - version, - submodule_index, - err_vec, - res.err(), - )); - return; - } - }; + let (ast, errors) = ptx::ModuleParser::parse_unchecked(module_text); + if !errors.is_empty() { + fn_logger.log(log::LogEntry::ModuleParsingError( + DumpWriter::get_file_name(module_index, version, submodule_index, "log"), + )); + fn_logger.log_io_error(self.writer.save_module_error_log( + module_index, + version, + submodule_index, + &*errors, + )); + } } pub(crate) fn module_exists(&self, hmod: CUmodule) -> bool { @@ -238,8 +233,7 @@ impl DumpWriter { module_index: usize, version: Option, submodule_index: Option, - recoverable: &[ptx::ParseError, PtxError>], - unrecoverable: Option, PtxError>>, + errors: &[ptx::ParseError, PtxError>], ) -> io::Result<()> { let mut log_file = match &self.dump_dir { None => return Ok(()), @@ -252,8 +246,9 @@ impl DumpWriter { "log", )); let mut file = File::create(log_file)?; - for error in unrecoverable.iter().chain(recoverable.iter()) { - writeln!(file, "{}", error)?; + for error in errors { + let pretty_print_error = DisplayParseError("", error); + writeln!(file, "{}", pretty_print_error)?; } Ok(()) } -- cgit v1.2.3