aboutsummaryrefslogtreecommitdiffhomepage
diff options
context:
space:
mode:
Diffstat
-rw-r--r--hip_runtime-sys/build.rs18
-rw-r--r--hip_runtime-sys/lib/amdhip64.def308
-rw-r--r--hip_runtime-sys/lib/amdhip64.libbin0 -> 71480 bytes
-rw-r--r--ptx/src/lib.rs112
-rw-r--r--ptx/src/ptx.lalrpop11
-rw-r--r--zluda_dump/src/trace.rs39
6 files changed, 461 insertions, 27 deletions
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
--- /dev/null
+++ b/hip_runtime-sys/lib/amdhip64.lib
Binary files 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<ast::Module<'input>, Vec<ParseError<usize, Token<'input>, 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<ParseError<usize, Token<'input>, ast::PtxError>>,
+ );
+}
+
+impl ModuleParserExt for ModuleParser {
+ fn parse_checked<'input>(
+ txt: &'input str,
+ ) -> Result<ast::Module<'input>, Vec<ParseError<usize, Token<'input>, 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<ParseError<usize, Token<'input>, 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<Loc, Tok, Err>);
+
+impl<'a, Loc, Tok, Err> fmt::Display for DisplayParseError<'a, Loc, Tok, Err>
+where
+ Loc: fmt::Display + Into<usize> + 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<T>(x: Vec<Option<T>>) -> Vec<T> {
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<ast::Statement<ast::ParsedArgParams<'input>>> = {
<p:PredAt?> <i:Instruction> ";" => Some(ast::Statement::Instruction(p, i)),
PragmaStatement => None,
"{" <s:Statement*> "}" => 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<usize>,
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<usize>,
submodule_index: Option<usize>,
- recoverable: &[ptx::ParseError<usize, Token<'input>, PtxError>],
- unrecoverable: Option<ptx::ParseError<usize, Token<'input>, PtxError>>,
+ errors: &[ptx::ParseError<usize, Token<'input>, 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(())
}
generated by cgit v1.2.3 (git 2.39.1) at 2024-12-25 04:49:59 +0000