diff options
| author | Andrzej Janik <[email protected]> | 2024-09-13 01:07:31 +0200 |
|---|---|---|
| committer | GitHub <[email protected]> | 2024-09-13 01:07:31 +0200 |
| commit | 46def3e7e09dbf4d3e7287a72bfecb73e6e429c5 (patch) | |
| tree | 6eebad3f9722ee9127c2640300ae20047d4acd9d /ptx/src | |
| parent | 193eb29be825370449afb1fe2358f6a654aa0986 (diff) | |
| download | ZLUDA-46def3e7e09dbf4d3e7287a72bfecb73e6e429c5.tar.gz ZLUDA-46def3e7e09dbf4d3e7287a72bfecb73e6e429c5.zip | |
Connect new parser to LLVM bitcode backend (#269)
This is very incomplete. Just enough code to emit LLVM bitcode and continue further development
Diffstat (limited to 'ptx/src')
| -rw-r--r-- | ptx/src/pass/emit_llvm.rs | 692 | ||||
| -rw-r--r-- | ptx/src/pass/mod.rs | 30 | ||||
| -rw-r--r-- | ptx/src/test/spirv_run/mod.rs | 304 |
3 files changed, 760 insertions, 266 deletions
diff --git a/ptx/src/pass/emit_llvm.rs b/ptx/src/pass/emit_llvm.rs new file mode 100644 index 0000000..44debba --- /dev/null +++ b/ptx/src/pass/emit_llvm.rs @@ -0,0 +1,692 @@ +// We use Raw LLVM-C bindings here because using inkwell is just not worth it.
+// Specifically the issue is with builder functions. We maintain the mapping
+// between ZLUDA identifiers and LLVM values. When using inkwell, LLVM values
+// are kept as instances `AnyValueEnum`. Now look at the signature of
+// `Builder::build_int_add(...)`:
+// pub fn build_int_add<T: IntMathValue<'ctx>>(&self, lhs: T, rhs: T, name: &str, ) -> Result<T, BuilderError>
+// At this point both lhs and rhs are `AnyValueEnum`. To call
+// `build_int_add(...)` we would have to do something like this:
+// if let (Ok(lhs), Ok(rhs)) = (lhs.as_int(), rhs.as_int()) {
+// builder.build_int_add(lhs, rhs, dst)?;
+// } else if let (Ok(lhs), Ok(rhs)) = (lhs.as_pointer(), rhs.as_pointer()) {
+// builder.build_int_add(lhs, rhs, dst)?;
+// } else if let (Ok(lhs), Ok(rhs)) = (lhs.as_vector(), rhs.as_vector()) {
+// builder.build_int_add(lhs, rhs, dst)?;
+// } else {
+// return Err(error_unrachable());
+// }
+// while with plain LLVM-C it's just:
+// unsafe { LLVMBuildAdd(builder, lhs, rhs, dst) };
+
+use std::convert::{TryFrom, TryInto};
+use std::ffi::CStr;
+use std::ops::Deref;
+use std::ptr;
+
+use super::*;
+use llvm_zluda::analysis::{LLVMVerifierFailureAction, LLVMVerifyModule};
+use llvm_zluda::bit_writer::LLVMWriteBitcodeToMemoryBuffer;
+use llvm_zluda::core::*;
+use llvm_zluda::prelude::*;
+use llvm_zluda::{LLVMCallConv, LLVMZludaBuildAlloca};
+
+const LLVM_UNNAMED: &CStr = c"";
+// https://llvm.org/docs/AMDGPUUsage.html#address-spaces
+const GENERIC_ADDRESS_SPACE: u32 = 0;
+const GLOBAL_ADDRESS_SPACE: u32 = 1;
+const SHARED_ADDRESS_SPACE: u32 = 3;
+const CONSTANT_ADDRESS_SPACE: u32 = 4;
+const PRIVATE_ADDRESS_SPACE: u32 = 5;
+
+struct Context(LLVMContextRef);
+
+impl Context {
+ fn new() -> Self {
+ Self(unsafe { LLVMContextCreate() })
+ }
+
+ fn get(&self) -> LLVMContextRef {
+ self.0
+ }
+}
+
+impl Drop for Context {
+ fn drop(&mut self) {
+ unsafe {
+ LLVMContextDispose(self.0);
+ }
+ }
+}
+
+struct Module(LLVMModuleRef);
+
+impl Module {
+ fn new(ctx: &Context, name: &CStr) -> Self {
+ Self(unsafe { LLVMModuleCreateWithNameInContext(name.as_ptr(), ctx.get()) })
+ }
+
+ fn get(&self) -> LLVMModuleRef {
+ self.0
+ }
+
+ fn verify(&self) -> Result<(), Message> {
+ let mut err = ptr::null_mut();
+ let error = unsafe {
+ LLVMVerifyModule(
+ self.get(),
+ LLVMVerifierFailureAction::LLVMReturnStatusAction,
+ &mut err,
+ )
+ };
+ if error == 1 && err != ptr::null_mut() {
+ Err(Message(unsafe { CStr::from_ptr(err) }))
+ } else {
+ Ok(())
+ }
+ }
+
+ fn write_bitcode_to_memory(&self) -> MemoryBuffer {
+ let memory_buffer = unsafe { LLVMWriteBitcodeToMemoryBuffer(self.get()) };
+ MemoryBuffer(memory_buffer)
+ }
+
+ fn write_to_stderr(&self) {
+ unsafe { LLVMDumpModule(self.get()) };
+ }
+}
+
+impl Drop for Module {
+ fn drop(&mut self) {
+ unsafe {
+ LLVMDisposeModule(self.0);
+ }
+ }
+}
+
+struct Builder(LLVMBuilderRef);
+
+impl Builder {
+ fn new(ctx: &Context) -> Self {
+ Self::new_raw(ctx.get())
+ }
+
+ fn new_raw(ctx: LLVMContextRef) -> Self {
+ Self(unsafe { LLVMCreateBuilderInContext(ctx) })
+ }
+
+ fn get(&self) -> LLVMBuilderRef {
+ self.0
+ }
+}
+
+impl Drop for Builder {
+ fn drop(&mut self) {
+ unsafe {
+ LLVMDisposeBuilder(self.0);
+ }
+ }
+}
+
+struct Message(&'static CStr);
+
+impl Drop for Message {
+ fn drop(&mut self) {
+ unsafe {
+ LLVMDisposeMessage(self.0.as_ptr().cast_mut());
+ }
+ }
+}
+
+impl std::fmt::Debug for Message {
+ fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result {
+ std::fmt::Debug::fmt(&self.0, f)
+ }
+}
+
+pub struct MemoryBuffer(LLVMMemoryBufferRef);
+
+impl Drop for MemoryBuffer {
+ fn drop(&mut self) {
+ unsafe {
+ LLVMDisposeMemoryBuffer(self.0);
+ }
+ }
+}
+
+impl Deref for MemoryBuffer {
+ type Target = [u8];
+
+ fn deref(&self) -> &Self::Target {
+ let data = unsafe { LLVMGetBufferStart(self.0) };
+ let len = unsafe { LLVMGetBufferSize(self.0) };
+ unsafe { std::slice::from_raw_parts(data.cast(), len) }
+ }
+}
+
+pub(super) fn run<'input>(
+ id_defs: &GlobalStringIdResolver<'input>,
+ call_map: MethodsCallMap<'input>,
+ directives: Vec<Directive<'input>>,
+) -> Result<MemoryBuffer, TranslateError> {
+ let context = Context::new();
+ let module = Module::new(&context, LLVM_UNNAMED);
+ let mut emit_ctx = ModuleEmitContext::new(&context, &module, id_defs);
+ for directive in directives {
+ match directive {
+ Directive::Variable(..) => todo!(),
+ Directive::Method(method) => emit_ctx.emit_method(method)?,
+ }
+ }
+ module.write_to_stderr();
+ if let Err(err) = module.verify() {
+ panic!("{:?}", err);
+ }
+ Ok(module.write_bitcode_to_memory())
+}
+
+struct ModuleEmitContext<'a, 'input> {
+ context: LLVMContextRef,
+ module: LLVMModuleRef,
+ builder: Builder,
+ id_defs: &'a GlobalStringIdResolver<'input>,
+ resolver: ResolveIdent,
+}
+
+impl<'a, 'input> ModuleEmitContext<'a, 'input> {
+ fn new(
+ context: &Context,
+ module: &Module,
+ id_defs: &'a GlobalStringIdResolver<'input>,
+ ) -> Self {
+ ModuleEmitContext {
+ context: context.get(),
+ module: module.get(),
+ builder: Builder::new(context),
+ id_defs,
+ resolver: ResolveIdent::new(&id_defs),
+ }
+ }
+
+ fn kernel_call_convention() -> u32 {
+ LLVMCallConv::LLVMAMDGPUKERNELCallConv as u32
+ }
+
+ fn func_call_convention() -> u32 {
+ LLVMCallConv::LLVMCCallConv as u32
+ }
+
+ fn emit_method(&mut self, method: Function<'input>) -> Result<(), TranslateError> {
+ let func_decl = method.func_decl.borrow();
+ let name = method
+ .import_as
+ .as_deref()
+ .unwrap_or_else(|| match func_decl.name {
+ ast::MethodName::Kernel(name) => name,
+ ast::MethodName::Func(id) => self.id_defs.reverse_variables[&id],
+ });
+ let name = CString::new(name).map_err(|_| error_unreachable())?;
+ let fn_type = self.function_type(
+ func_decl.return_arguments.iter().map(|v| &v.v_type),
+ func_decl.input_arguments.iter().map(|v| &v.v_type),
+ );
+ let fn_ = unsafe { LLVMAddFunction(self.module, name.as_ptr(), fn_type) };
+ for (i, param) in func_decl.input_arguments.iter().enumerate() {
+ let value = unsafe { LLVMGetParam(fn_, i as u32) };
+ let name = self.resolver.get_or_add(param.name);
+ unsafe { LLVMSetValueName2(value, name.as_ptr().cast(), name.len()) };
+ self.resolver.register(param.name, value);
+ }
+ let call_conv = if func_decl.name.is_kernel() {
+ Self::kernel_call_convention()
+ } else {
+ Self::func_call_convention()
+ };
+ unsafe { LLVMSetFunctionCallConv(fn_, call_conv) };
+ if let Some(statements) = method.body {
+ let variables_bb =
+ unsafe { LLVMAppendBasicBlockInContext(self.context, fn_, LLVM_UNNAMED.as_ptr()) };
+ let variables_builder = Builder::new_raw(self.context);
+ unsafe { LLVMPositionBuilderAtEnd(variables_builder.get(), variables_bb) };
+ let real_bb =
+ unsafe { LLVMAppendBasicBlockInContext(self.context, fn_, LLVM_UNNAMED.as_ptr()) };
+ unsafe { LLVMPositionBuilderAtEnd(self.builder.get(), real_bb) };
+ let mut method_emitter = MethodEmitContext::new(self, fn_, variables_builder);
+ for statement in statements {
+ method_emitter.emit_statement(statement)?;
+ }
+ unsafe { LLVMBuildBr(method_emitter.variables_builder.get(), real_bb) };
+ }
+ Ok(())
+ }
+
+ fn function_type(
+ &self,
+ return_args: impl ExactSizeIterator<Item = &'a ast::Type>,
+ input_args: impl ExactSizeIterator<Item = &'a ast::Type>,
+ ) -> LLVMTypeRef {
+ if return_args.len() == 0 {
+ let mut input_args = input_args
+ .map(|type_| match type_ {
+ ast::Type::Scalar(scalar) => match scalar {
+ ast::ScalarType::Pred => {
+ unsafe { LLVMInt1TypeInContext(self.context) }
+ }
+ ast::ScalarType::S8 | ast::ScalarType::B8 | ast::ScalarType::U8 => {
+ unsafe { LLVMInt8TypeInContext(self.context) }
+ }
+ ast::ScalarType::B16 | ast::ScalarType::U16 | ast::ScalarType::S16 => {
+ unsafe { LLVMInt16TypeInContext(self.context) }
+ }
+ ast::ScalarType::S32 | ast::ScalarType::B32 | ast::ScalarType::U32 => {
+ unsafe { LLVMInt32TypeInContext(self.context) }
+ }
+ ast::ScalarType::U64 | ast::ScalarType::S64 | ast::ScalarType::B64 => {
+ unsafe { LLVMInt64TypeInContext(self.context) }
+ }
+ ast::ScalarType::B128 => {
+ unsafe { LLVMInt128TypeInContext(self.context) }
+ }
+ ast::ScalarType::F16 => {
+ unsafe { LLVMHalfTypeInContext(self.context) }
+ }
+ ast::ScalarType::F32 => {
+ unsafe { LLVMFloatTypeInContext(self.context) }
+ }
+ ast::ScalarType::F64 => {
+ unsafe { LLVMDoubleTypeInContext(self.context) }
+ }
+ ast::ScalarType::BF16 => {
+ unsafe { LLVMBFloatTypeInContext(self.context) }
+ }
+ ast::ScalarType::U16x2 => todo!(),
+ ast::ScalarType::S16x2 => todo!(),
+ ast::ScalarType::F16x2 => todo!(),
+ ast::ScalarType::BF16x2 => todo!(),
+ },
+ ast::Type::Vector(_, _) => todo!(),
+ ast::Type::Array(_, _, _) => todo!(),
+ ast::Type::Pointer(_, _) => todo!(),
+ })
+ .collect::<Vec<_>>();
+ return unsafe {
+ LLVMFunctionType(
+ LLVMVoidTypeInContext(self.context),
+ input_args.as_mut_ptr(),
+ input_args.len() as u32,
+ 0,
+ )
+ };
+ }
+ todo!()
+ }
+}
+
+struct MethodEmitContext<'a, 'input> {
+ context: LLVMContextRef,
+ module: LLVMModuleRef,
+ method: LLVMValueRef,
+ builder: LLVMBuilderRef,
+ id_defs: &'a GlobalStringIdResolver<'input>,
+ variables_builder: Builder,
+ resolver: &'a mut ResolveIdent,
+}
+
+impl<'a, 'input> MethodEmitContext<'a, 'input> {
+ fn new<'x>(
+ parent: &'a mut ModuleEmitContext<'x, 'input>,
+ method: LLVMValueRef,
+ variables_builder: Builder,
+ ) -> MethodEmitContext<'a, 'input> {
+ MethodEmitContext {
+ context: parent.context,
+ module: parent.module,
+ builder: parent.builder.get(),
+ id_defs: parent.id_defs,
+ variables_builder,
+ resolver: &mut parent.resolver,
+ method,
+ }
+ }
+
+ fn emit_statement(
+ &mut self,
+ statement: Statement<ast::Instruction<SpirvWord>, SpirvWord>,
+ ) -> Result<(), TranslateError> {
+ Ok(match statement {
+ Statement::Variable(var) => self.emit_variable(var)?,
+ Statement::Label(label) => self.emit_label(label),
+ Statement::Instruction(inst) => self.emit_instruction(inst)?,
+ Statement::Conditional(_) => todo!(),
+ Statement::LoadVar(var) => self.emit_load_variable(var)?,
+ Statement::StoreVar(store) => self.emit_store_var(store)?,
+ Statement::Conversion(conversion) => self.emit_conversion(conversion)?,
+ Statement::Constant(constant) => self.emit_constant(constant)?,
+ Statement::RetValue(_, _) => todo!(),
+ Statement::PtrAccess(_) => todo!(),
+ Statement::RepackVector(_) => todo!(),
+ Statement::FunctionPointer(_) => todo!(),
+ })
+ }
+
+ fn emit_variable(&mut self, var: ast::Variable<SpirvWord>) -> Result<(), TranslateError> {
+ let alloca = unsafe {
+ LLVMZludaBuildAlloca(
+ self.variables_builder.get(),
+ get_type(self.context, &var.v_type)?,
+ get_state_space(var.state_space)?,
+ self.resolver.get_or_add_raw(var.name),
+ )
+ };
+ self.resolver.register(var.name, alloca);
+ if let Some(align) = var.align {
+ unsafe { LLVMSetAlignment(alloca, align) };
+ }
+ if !var.array_init.is_empty() {
+ todo!()
+ }
+ Ok(())
+ }
+
+ fn emit_label(&mut self, label: SpirvWord) {
+ let block = unsafe {
+ LLVMAppendBasicBlockInContext(
+ self.context,
+ self.method,
+ self.resolver.get_or_add_raw(label),
+ )
+ };
+ let last_block = unsafe { LLVMGetInsertBlock(self.builder) };
+ if unsafe { LLVMGetBasicBlockTerminator(last_block) } == ptr::null_mut() {
+ unsafe { LLVMBuildBr(self.builder, block) };
+ }
+ unsafe { LLVMPositionBuilderAtEnd(self.builder, block) };
+ }
+
+ fn emit_store_var(&mut self, store: StoreVarDetails) -> Result<(), TranslateError> {
+ let ptr = self.resolver.value(store.arg.src1)?;
+ let value = self.resolver.value(store.arg.src2)?;
+ unsafe { LLVMBuildStore(self.builder, value, ptr) };
+ Ok(())
+ }
+
+ fn emit_instruction(
+ &mut self,
+ inst: ast::Instruction<SpirvWord>,
+ ) -> Result<(), TranslateError> {
+ match inst {
+ ast::Instruction::Mov { data, arguments } => todo!(),
+ ast::Instruction::Ld { data, arguments } => self.emit_ld(data, arguments),
+ ast::Instruction::Add { data, arguments } => self.emit_add(data, arguments),
+ ast::Instruction::St { data, arguments } => self.emit_st(data, arguments),
+ ast::Instruction::Mul { data, arguments } => todo!(),
+ ast::Instruction::Setp { data, arguments } => todo!(),
+ ast::Instruction::SetpBool { data, arguments } => todo!(),
+ ast::Instruction::Not { data, arguments } => todo!(),
+ ast::Instruction::Or { data, arguments } => todo!(),
+ ast::Instruction::And { data, arguments } => todo!(),
+ ast::Instruction::Bra { arguments } => todo!(),
+ ast::Instruction::Call { data, arguments } => todo!(),
+ ast::Instruction::Cvt { data, arguments } => todo!(),
+ ast::Instruction::Shr { data, arguments } => todo!(),
+ ast::Instruction::Shl { data, arguments } => todo!(),
+ ast::Instruction::Ret { data } => Ok(self.emit_ret(data)),
+ ast::Instruction::Cvta { data, arguments } => todo!(),
+ ast::Instruction::Abs { data, arguments } => todo!(),
+ ast::Instruction::Mad { data, arguments } => todo!(),
+ ast::Instruction::Fma { data, arguments } => todo!(),
+ ast::Instruction::Sub { data, arguments } => todo!(),
+ ast::Instruction::Min { data, arguments } => todo!(),
+ ast::Instruction::Max { data, arguments } => todo!(),
+ ast::Instruction::Rcp { data, arguments } => todo!(),
+ ast::Instruction::Sqrt { data, arguments } => todo!(),
+ ast::Instruction::Rsqrt { data, arguments } => todo!(),
+ ast::Instruction::Selp { data, arguments } => todo!(),
+ ast::Instruction::Bar { data, arguments } => todo!(),
+ ast::Instruction::Atom { data, arguments } => todo!(),
+ ast::Instruction::AtomCas { data, arguments } => todo!(),
+ ast::Instruction::Div { data, arguments } => todo!(),
+ ast::Instruction::Neg { data, arguments } => todo!(),
+ ast::Instruction::Sin { data, arguments } => todo!(),
+ ast::Instruction::Cos { data, arguments } => todo!(),
+ ast::Instruction::Lg2 { data, arguments } => todo!(),
+ ast::Instruction::Ex2 { data, arguments } => todo!(),
+ ast::Instruction::Clz { data, arguments } => todo!(),
+ ast::Instruction::Brev { data, arguments } => todo!(),
+ ast::Instruction::Popc { data, arguments } => todo!(),
+ ast::Instruction::Xor { data, arguments } => todo!(),
+ ast::Instruction::Rem { data, arguments } => todo!(),
+ ast::Instruction::Bfe { data, arguments } => todo!(),
+ ast::Instruction::Bfi { data, arguments } => todo!(),
+ ast::Instruction::PrmtSlow { arguments } => todo!(),
+ ast::Instruction::Prmt { data, arguments } => todo!(),
+ ast::Instruction::Activemask { arguments } => todo!(),
+ ast::Instruction::Membar { data } => todo!(),
+ ast::Instruction::Trap {} => todo!(),
+ }
+ }
+
+ fn emit_ld(
+ &mut self,
+ data: ast::LdDetails,
+ arguments: ast::LdArgs<SpirvWord>,
+ ) -> Result<(), TranslateError> {
+ if data.non_coherent {
+ todo!()
+ }
+ if data.qualifier != ast::LdStQualifier::Weak {
+ todo!()
+ }
+ let builder = self.builder;
+ let type_ = get_type(self.context, &data.typ)?;
+ let ptr = self.resolver.value(arguments.src)?;
+ self.resolver.with_result(arguments.dst, |dst| unsafe {
+ LLVMBuildLoad2(builder, type_, ptr, dst)
+ });
+ Ok(())
+ }
+
+ fn emit_load_variable(&mut self, var: LoadVarDetails) -> Result<(), TranslateError> {
+ if var.member_index.is_some() {
+ todo!()
+ }
+ let builder = self.builder;
+ let type_ = get_type(self.context, &var.typ)?;
+ let ptr = self.resolver.value(var.arg.src)?;
+ self.resolver.with_result(var.arg.dst, |dst| unsafe {
+ LLVMBuildLoad2(builder, type_, ptr, dst)
+ });
+ Ok(())
+ }
+
+ fn emit_conversion(&mut self, conversion: ImplicitConversion) -> Result<(), TranslateError> {
+ let builder = self.builder;
+ match conversion.kind {
+ ConversionKind::Default => todo!(),
+ ConversionKind::SignExtend => todo!(),
+ ConversionKind::BitToPtr => {
+ let src = self.resolver.value(conversion.src)?;
+ let type_ = get_pointer_type(self.context, conversion.to_space)?;
+ self.resolver.with_result(conversion.dst, |dst| unsafe {
+ LLVMBuildIntToPtr(builder, src, type_, dst)
+ });
+ Ok(())
+ }
+ ConversionKind::PtrToPtr => todo!(),
+ ConversionKind::AddressOf => todo!(),
+ }
+ }
+
+ fn emit_constant(&mut self, constant: ConstantDefinition) -> Result<(), TranslateError> {
+ let type_ = get_scalar_type(self.context, constant.typ);
+ let value = match constant.value {
+ ast::ImmediateValue::U64(x) => unsafe { LLVMConstInt(type_, x, 0) },
+ ast::ImmediateValue::S64(x) => unsafe { LLVMConstInt(type_, x as u64, 0) },
+ ast::ImmediateValue::F32(x) => unsafe { LLVMConstReal(type_, x as f64) },
+ ast::ImmediateValue::F64(x) => unsafe { LLVMConstReal(type_, x) },
+ };
+ self.resolver.register(constant.dst, value);
+ Ok(())
+ }
+
+ fn emit_add(
+ &mut self,
+ data: ast::ArithDetails,
+ arguments: ast::AddArgs<SpirvWord>,
+ ) -> Result<(), TranslateError> {
+ let builder = self.builder;
+ let src1 = self.resolver.value(arguments.src1)?;
+ let src2 = self.resolver.value(arguments.src2)?;
+ let fn_ = match data {
+ ast::ArithDetails::Integer(integer) => LLVMBuildAdd,
+ ast::ArithDetails::Float(float) => LLVMBuildFAdd,
+ };
+ self.resolver.with_result(arguments.dst, |dst| unsafe {
+ fn_(builder, src1, src2, dst)
+ });
+ Ok(())
+ }
+
+ fn emit_st(
+ &self,
+ data: ptx_parser::StData,
+ arguments: ptx_parser::StArgs<SpirvWord>,
+ ) -> Result<(), TranslateError> {
+ let ptr = self.resolver.value(arguments.src1)?;
+ let value = self.resolver.value(arguments.src2)?;
+ if data.qualifier != ast::LdStQualifier::Weak {
+ todo!()
+ }
+ unsafe { LLVMBuildStore(self.builder, value, ptr) };
+ Ok(())
+ }
+
+ fn emit_ret(&self, _data: ptx_parser::RetData) {
+ unsafe { LLVMBuildRetVoid(self.builder) };
+ }
+}
+
+fn get_pointer_type<'ctx>(
+ context: LLVMContextRef,
+ to_space: ast::StateSpace,
+) -> Result<LLVMTypeRef, TranslateError> {
+ Ok(unsafe { LLVMPointerTypeInContext(context, get_state_space(to_space)?) })
+}
+
+fn get_type(context: LLVMContextRef, type_: &ast::Type) -> Result<LLVMTypeRef, TranslateError> {
+ Ok(match type_ {
+ ast::Type::Scalar(scalar) => get_scalar_type(context, *scalar),
+ ast::Type::Vector(size, scalar) => {
+ let base_type = get_scalar_type(context, *scalar);
+ unsafe { LLVMVectorType(base_type, *size as u32) }
+ }
+ ast::Type::Array(vec, scalar, dimensions) => {
+ let mut underlying_type = get_scalar_type(context, *scalar);
+ if let Some(size) = vec {
+ underlying_type = unsafe { LLVMVectorType(underlying_type, size.get() as u32) };
+ }
+ if dimensions.is_empty() {
+ return Ok(unsafe { LLVMArrayType2(underlying_type, 0) });
+ }
+ dimensions
+ .iter()
+ .rfold(underlying_type, |result, dimension| unsafe {
+ LLVMArrayType2(result, *dimension as u64)
+ })
+ }
+ ast::Type::Pointer(_, space) => get_pointer_type(context, *space)?,
+ })
+}
+
+fn get_scalar_type(context: LLVMContextRef, type_: ast::ScalarType) -> LLVMTypeRef {
+ match type_ {
+ ast::ScalarType::Pred => unsafe { LLVMInt1TypeInContext(context) },
+ ast::ScalarType::S8 | ast::ScalarType::B8 | ast::ScalarType::U8 => unsafe {
+ LLVMInt8TypeInContext(context)
+ },
+ ast::ScalarType::B16 | ast::ScalarType::U16 | ast::ScalarType::S16 => unsafe {
+ LLVMInt16TypeInContext(context)
+ },
+ ast::ScalarType::S32 | ast::ScalarType::B32 | ast::ScalarType::U32 => unsafe {
+ LLVMInt32TypeInContext(context)
+ },
+ ast::ScalarType::U64 | ast::ScalarType::S64 | ast::ScalarType::B64 => unsafe {
+ LLVMInt64TypeInContext(context)
+ },
+ ast::ScalarType::B128 => unsafe { LLVMInt128TypeInContext(context) },
+ ast::ScalarType::F16 => unsafe { LLVMHalfTypeInContext(context) },
+ ast::ScalarType::F32 => unsafe { LLVMFloatTypeInContext(context) },
+ ast::ScalarType::F64 => unsafe { LLVMDoubleTypeInContext(context) },
+ ast::ScalarType::BF16 => unsafe { LLVMBFloatTypeInContext(context) },
+ ast::ScalarType::U16x2 => todo!(),
+ ast::ScalarType::S16x2 => todo!(),
+ ast::ScalarType::F16x2 => todo!(),
+ ast::ScalarType::BF16x2 => todo!(),
+ }
+}
+
+fn get_state_space(space: ast::StateSpace) -> Result<u32, TranslateError> {
+ match space {
+ ast::StateSpace::Reg => Ok(PRIVATE_ADDRESS_SPACE),
+ ast::StateSpace::Generic => Ok(GENERIC_ADDRESS_SPACE),
+ ast::StateSpace::Sreg => Ok(PRIVATE_ADDRESS_SPACE),
+ ast::StateSpace::Param => Err(TranslateError::Todo),
+ ast::StateSpace::ParamEntry => Err(TranslateError::Todo),
+ ast::StateSpace::ParamFunc => Err(TranslateError::Todo),
+ ast::StateSpace::Local => Ok(PRIVATE_ADDRESS_SPACE),
+ ast::StateSpace::Global => Ok(GLOBAL_ADDRESS_SPACE),
+ ast::StateSpace::Const => Ok(CONSTANT_ADDRESS_SPACE),
+ ast::StateSpace::Shared => Ok(SHARED_ADDRESS_SPACE),
+ ast::StateSpace::SharedCta => Err(TranslateError::Todo),
+ ast::StateSpace::SharedCluster => Err(TranslateError::Todo),
+ }
+}
+
+struct ResolveIdent {
+ words: HashMap<SpirvWord, String>,
+ values: HashMap<SpirvWord, LLVMValueRef>,
+}
+
+impl ResolveIdent {
+ fn new<'input>(_id_defs: &GlobalStringIdResolver<'input>) -> Self {
+ ResolveIdent {
+ words: HashMap::new(),
+ values: HashMap::new(),
+ }
+ }
+
+ fn get_or_ad_impl<'a, T>(&'a mut self, word: SpirvWord, fn_: impl FnOnce(&'a str) -> T) -> T {
+ let str = match self.words.entry(word) {
+ hash_map::Entry::Occupied(entry) => entry.into_mut(),
+ hash_map::Entry::Vacant(entry) => {
+ let mut text = word.0.to_string();
+ text.push('\0');
+ entry.insert(text)
+ }
+ };
+ fn_(&str[..str.len() - 1])
+ }
+
+ fn get_or_add(&mut self, word: SpirvWord) -> &str {
+ self.get_or_ad_impl(word, |x| x)
+ }
+
+ fn get_or_add_raw(&mut self, word: SpirvWord) -> *const i8 {
+ self.get_or_add(word).as_ptr().cast()
+ }
+
+ fn register(&mut self, word: SpirvWord, v: LLVMValueRef) {
+ self.values.insert(word, v);
+ }
+
+ fn value(&self, word: SpirvWord) -> Result<LLVMValueRef, TranslateError> {
+ self.values
+ .get(&word)
+ .copied()
+ .ok_or_else(|| error_unreachable())
+ }
+
+ fn with_result(&mut self, word: SpirvWord, fn_: impl FnOnce(*const i8) -> LLVMValueRef) {
+ let t = self.get_or_ad_impl(word, |dst| fn_(dst.as_ptr().cast()));
+ self.register(word, t);
+ }
+}
diff --git a/ptx/src/pass/mod.rs b/ptx/src/pass/mod.rs index 2be6297..3aa3b0a 100644 --- a/ptx/src/pass/mod.rs +++ b/ptx/src/pass/mod.rs @@ -16,6 +16,7 @@ use std::{ mod convert_dynamic_shared_memory_usage;
mod convert_to_stateful_memory_access;
mod convert_to_typed;
+pub(crate) mod emit_llvm;
mod emit_spirv;
mod expand_arguments;
mod extract_globals;
@@ -30,7 +31,7 @@ static ZLUDA_PTX_IMPL_INTEL: &'static [u8] = include_bytes!("../../lib/zluda_ptx static ZLUDA_PTX_IMPL_AMD: &'static [u8] = include_bytes!("../../lib/zluda_ptx_impl.bc");
const ZLUDA_PTX_PREFIX: &'static str = "__zluda_ptx_impl__";
-pub fn to_spirv_module<'input>(ast: ast::Module<'input>) -> Result<Module, TranslateError> {
+pub fn to_llvm_module<'input>(ast: ast::Module<'input>) -> Result<Module, TranslateError> {
let mut id_defs = GlobalStringIdResolver::<'input>::new(SpirvWord(1));
let mut ptx_impl_imports = HashMap::new();
let directives = ast
@@ -56,17 +57,10 @@ pub fn to_spirv_module<'input>(ast: ast::Module<'input>) -> Result<Module, Trans })?;
normalize_variable_decls(&mut directives);
let denorm_information = compute_denorm_information(&directives);
- let (spirv, kernel_info, build_options) =
- emit_spirv::run(builder, &id_defs, call_map, denorm_information, directives)?;
+ let llvm_ir = emit_llvm::run(&id_defs, call_map, directives)?;
Ok(Module {
- spirv,
- kernel_info,
- should_link_ptx_impl: if must_link_ptx_impl {
- Some((ZLUDA_PTX_IMPL_INTEL, ZLUDA_PTX_IMPL_AMD))
- } else {
- None
- },
- build_options,
+ llvm_ir,
+ kernel_info: HashMap::new(),
})
}
@@ -187,22 +181,14 @@ fn to_ssa<'input, 'b>( }
pub struct Module {
- pub spirv: dr::Module,
+ pub llvm_ir: emit_llvm::MemoryBuffer,
pub kernel_info: HashMap<String, KernelInfo>,
- pub should_link_ptx_impl: Option<(&'static [u8], &'static [u8])>,
- pub build_options: CString,
-}
-
-impl Module {
- pub fn assemble(&self) -> Vec<u32> {
- self.spirv.assemble()
- }
}
struct GlobalStringIdResolver<'input> {
current_id: SpirvWord,
variables: HashMap<Cow<'input, str>, SpirvWord>,
- reverse_variables: HashMap<SpirvWord, &'input str>,
+ pub(crate) reverse_variables: HashMap<SpirvWord, &'input str>,
variables_type_check: HashMap<SpirvWord, Option<(ast::Type, ast::StateSpace, bool)>>,
special_registers: SpecialRegistersMap,
fns: HashMap<SpirvWord, FnSigMapper<'input>>,
@@ -611,6 +597,7 @@ fn error_unreachable() -> TranslateError { TranslateError::Unreachable
}
+#[cfg(debug_assertions)]
fn error_unknown_symbol() -> TranslateError {
panic!()
}
@@ -620,6 +607,7 @@ fn error_unknown_symbol() -> TranslateError { TranslateError::UnknownSymbol
}
+#[cfg(debug_assertions)]
fn error_mismatched_type() -> TranslateError {
panic!()
}
diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs index a798720..69dd206 100644 --- a/ptx/src/test/spirv_run/mod.rs +++ b/ptx/src/test/spirv_run/mod.rs @@ -31,7 +31,7 @@ macro_rules! test_ptx { ($fn_name:ident, $input:expr, $output:expr) => {
paste::item! {
#[test]
- fn [<$fn_name _ptx>]() -> Result<(), Box<dyn std::error::Error>> {
+ fn [<$fn_name _hip>]() -> Result<(), Box<dyn std::error::Error>> {
let ptx = include_str!(concat!(stringify!($fn_name), ".ptx"));
let input = $input;
let mut output = $output;
@@ -48,29 +48,9 @@ macro_rules! test_ptx { test_cuda_assert(stringify!($fn_name), ptx, &input, &mut output)
}
}
-
- paste::item! {
- #[test]
- fn [<$fn_name _spvtxt>]() -> Result<(), Box<dyn std::error::Error>> {
- let ptx_txt = include_str!(concat!(stringify!($fn_name), ".ptx"));
- let spirv_file_name = concat!(stringify!($fn_name), ".spvtxt");
- let spirv_txt = include_bytes!(concat!(stringify!($fn_name), ".spvtxt"));
- test_spvtxt_assert(ptx_txt, spirv_txt, spirv_file_name)
- }
- }
};
- ($fn_name:ident) => {
- paste::item! {
- #[test]
- fn [<$fn_name _spvtxt>]() -> Result<(), Box<dyn std::error::Error>> {
- let ptx_txt = include_str!(concat!(stringify!($fn_name), ".ptx"));
- let spirv_file_name = concat!(stringify!($fn_name), ".spvtxt");
- let spirv_txt = include_bytes!(concat!(stringify!($fn_name), ".spvtxt"));
- test_spvtxt_assert(ptx_txt, spirv_txt, spirv_file_name)
- }
- }
- };
+ ($fn_name:ident) => {};
}
test_ptx!(ld_st, [1u64], [1u64]);
@@ -255,13 +235,11 @@ fn test_hip_assert< input: &[Input],
output: &mut [Output],
) -> Result<(), Box<dyn error::Error + 'a>> {
- let mut errors = Vec::new();
- let ast = ptx::ModuleParser::new().parse(&mut errors, ptx_text)?;
- assert!(errors.len() == 0);
- let zluda_module = translate::to_spirv_module(ast)?;
+ let ast = ptx_parser::parse_module_checked(ptx_text).unwrap();
+ let llvm_ir = pass::to_llvm_module(ast).unwrap();
let name = CString::new(name)?;
- let result = run_hip(name.as_c_str(), zluda_module, input, output)
- .map_err(|err| DisplayError { err })?;
+ let result =
+ run_hip(name.as_c_str(), llvm_ir, input, output).map_err(|err| DisplayError { err })?;
assert_eq!(result.as_slice(), output);
Ok(())
}
@@ -283,18 +261,6 @@ fn test_cuda_assert< Ok(())
}
-macro_rules! hip_call {
- ($expr:expr) => {
- #[allow(unused_unsafe)]
- {
- let err = unsafe { $expr };
- if err != hip_runtime_sys::hipError_t::hipSuccess {
- return Result::Err(err);
- }
- }
- };
-}
-
macro_rules! cuda_call {
($expr:expr) => {
#[allow(unused_unsafe)]
@@ -344,124 +310,76 @@ fn run_cuda<Input: From<u8> + Copy + Debug, Output: From<u8> + Copy + Debug + De fn run_hip<Input: From<u8> + Copy + Debug, Output: From<u8> + Copy + Debug + Default>(
name: &CStr,
- module: translate::Module,
+ module: pass::Module,
input: &[Input],
output: &mut [Output],
) -> Result<Vec<Output>, hipError_t> {
use hip_runtime_sys::*;
- hip_call! { hipInit(0) };
- let spirv = module.spirv.assemble();
+ unsafe { hipInit(0) }.unwrap();
let mut result = vec![0u8.into(); output.len()];
{
let dev = 0;
let mut stream = ptr::null_mut();
- hip_call! { hipStreamCreate(&mut stream) };
+ unsafe { hipStreamCreate(&mut stream) }.unwrap();
let mut dev_props = unsafe { mem::zeroed() };
- hip_call! { hipGetDeviceProperties(&mut dev_props, dev) };
- let elf_module = compile_amd(&dev_props, &*spirv, module.should_link_ptx_impl)
- .map_err(|_| hipError_t::hipErrorUnknown)?;
+ unsafe { hipGetDevicePropertiesR0600(&mut dev_props, dev) }.unwrap();
+ let elf_module = comgr::compile_bitcode(
+ unsafe { CStr::from_ptr(dev_props.gcnArchName.as_ptr()) },
+ &*module.llvm_ir,
+ )
+ .unwrap();
let mut module = ptr::null_mut();
- hip_call! { hipModuleLoadData(&mut module, elf_module.as_ptr() as _) };
+ unsafe { hipModuleLoadData(&mut module, elf_module.as_ptr() as _) }.unwrap();
let mut kernel = ptr::null_mut();
- hip_call! { hipModuleGetFunction(&mut kernel, module, name.as_ptr()) };
+ unsafe { hipModuleGetFunction(&mut kernel, module, name.as_ptr()) }.unwrap();
let mut inp_b = ptr::null_mut();
- hip_call! { hipMalloc(&mut inp_b, input.len() * mem::size_of::<Input>()) };
+ unsafe { hipMalloc(&mut inp_b, input.len() * mem::size_of::<Input>()) }.unwrap();
let mut out_b = ptr::null_mut();
- hip_call! { hipMalloc(&mut out_b, output.len() * mem::size_of::<Output>()) };
- hip_call! { hipMemcpyWithStream(inp_b, input.as_ptr() as _, input.len() * mem::size_of::<Input>(), hipMemcpyKind::hipMemcpyHostToDevice, stream) };
- hip_call! { hipMemset(out_b, 0, output.len() * mem::size_of::<Output>()) };
+ unsafe { hipMalloc(&mut out_b, output.len() * mem::size_of::<Output>()) }.unwrap();
+ unsafe {
+ hipMemcpyWithStream(
+ inp_b,
+ input.as_ptr() as _,
+ input.len() * mem::size_of::<Input>(),
+ hipMemcpyKind::hipMemcpyHostToDevice,
+ stream,
+ )
+ }
+ .unwrap();
+ unsafe { hipMemset(out_b, 0, output.len() * mem::size_of::<Output>()) }.unwrap();
let mut args = [&inp_b, &out_b];
- hip_call! { hipModuleLaunchKernel(kernel, 1,1,1,1,1,1, 1024, stream, args.as_mut_ptr() as _, ptr::null_mut()) };
- hip_call! { hipMemcpyAsync(result.as_mut_ptr() as _, out_b, output.len() * mem::size_of::<Output>(), hipMemcpyKind::hipMemcpyDeviceToHost, stream) };
- hip_call! { hipStreamSynchronize(stream) };
- hip_call! { hipFree(inp_b) };
- hip_call! { hipFree(out_b) };
- hip_call! { hipModuleUnload(module) };
- }
- Ok(result)
-}
-
-fn test_spvtxt_assert<'a>(
- ptx_txt: &'a str,
- spirv_txt: &'a [u8],
- spirv_file_name: &'a str,
-) -> Result<(), Box<dyn error::Error + 'a>> {
- let ast = ptx_parser::parse_module_checked(ptx_txt).unwrap();
- let spirv_module = pass::to_spirv_module(ast)?;
- let spv_context =
- unsafe { spirv_tools::spvContextCreate(spv_target_env::SPV_ENV_UNIVERSAL_1_3) };
- assert!(spv_context != ptr::null_mut());
- let mut spv_binary: spv_binary = ptr::null_mut();
- let result = unsafe {
- spirv_tools::spvTextToBinary(
- spv_context,
- spirv_txt.as_ptr() as *const _,
- spirv_txt.len(),
- &mut spv_binary,
- ptr::null_mut(),
- )
- };
- if result != spv_result_t::SPV_SUCCESS {
- panic!("{:?}\n{}", result, unsafe {
- str::from_utf8_unchecked(spirv_txt)
- });
- }
- let mut parsed_spirv = Vec::<u32>::new();
- let result = unsafe {
- spirv_tools::spvBinaryParse(
- spv_context,
- &mut parsed_spirv as *mut _ as *mut _,
- (*spv_binary).code,
- (*spv_binary).wordCount,
- Some(parse_header_cb),
- Some(parse_instruction_cb),
- ptr::null_mut(),
- )
- };
- assert!(result == spv_result_t::SPV_SUCCESS);
- let mut loader = Loader::new();
- rspirv::binary::parse_words(&parsed_spirv, &mut loader)?;
- let spvtxt_mod = loader.module();
- unsafe { spirv_tools::spvBinaryDestroy(spv_binary) };
- if !is_spirv_fns_equal(&spirv_module.spirv.functions, &spvtxt_mod.functions) {
- // We could simply use ptx_mod.disassemble, but SPIRV-Tools text formattinmg is so much nicer
- let spv_from_ptx_binary = spirv_module.spirv.assemble();
- let mut spv_text: spirv_tools::spv_text = ptr::null_mut();
- let result = unsafe {
- spirv_tools::spvBinaryToText(
- spv_context,
- spv_from_ptx_binary.as_ptr(),
- spv_from_ptx_binary.len(),
- (spirv_tools::spv_binary_to_text_options_t::SPV_BINARY_TO_TEXT_OPTION_INDENT | spirv_tools::spv_binary_to_text_options_t::SPV_BINARY_TO_TEXT_OPTION_NO_HEADER | spirv_tools::spv_binary_to_text_options_t::SPV_BINARY_TO_TEXT_OPTION_FRIENDLY_NAMES).0,
- &mut spv_text as *mut _,
- ptr::null_mut()
+ unsafe {
+ hipModuleLaunchKernel(
+ kernel,
+ 1,
+ 1,
+ 1,
+ 1,
+ 1,
+ 1,
+ 1024,
+ stream,
+ args.as_mut_ptr() as _,
+ ptr::null_mut(),
+ )
+ }
+ .unwrap();
+ unsafe {
+ hipMemcpyAsync(
+ result.as_mut_ptr() as _,
+ out_b,
+ output.len() * mem::size_of::<Output>(),
+ hipMemcpyKind::hipMemcpyDeviceToHost,
+ stream,
)
- };
- unsafe { spirv_tools::spvContextDestroy(spv_context) };
- let spirv_text = if result == spv_result_t::SPV_SUCCESS {
- let raw_text = unsafe {
- std::slice::from_raw_parts((*spv_text).str_ as *const u8, (*spv_text).length)
- };
- let spv_from_ptx_text = unsafe { str::from_utf8_unchecked(raw_text) };
- // TODO: stop leaking kernel text
- Cow::Borrowed(spv_from_ptx_text)
- } else {
- Cow::Owned(spirv_module.spirv.disassemble())
- };
- if let Ok(dump_path) = env::var("ZLUDA_TEST_SPIRV_DUMP_DIR") {
- let mut path = PathBuf::from(dump_path);
- if let Ok(()) = fs::create_dir_all(&path) {
- path.push(spirv_file_name);
- #[allow(unused_must_use)]
- {
- fs::write(path, spirv_text.as_bytes());
- }
- }
}
- panic!("{}", spirv_text.to_string());
+ .unwrap();
+ unsafe { hipStreamSynchronize(stream) }.unwrap();
+ unsafe { hipFree(inp_b) }.unwrap();
+ unsafe { hipFree(out_b) }.unwrap();
+ unsafe { hipModuleUnload(module) }.unwrap();
}
- unsafe { spirv_tools::spvContextDestroy(spv_context) };
- Ok(())
+ Ok(result)
}
struct EqMap<T>
@@ -654,110 +572,6 @@ const AMDGPU_BITCODE: [&'static str; 8] = [ ];
const AMDGPU_BITCODE_DEVICE_PREFIX: &'static str = "oclc_isa_version_";
-fn compile_amd(
- device_pros: &hip::hipDeviceProp_t,
- spirv_il: &[u32],
- ptx_lib: Option<(&'static [u8], &'static [u8])>,
-) -> io::Result<Vec<u8>> {
- let null_terminator = device_pros
- .gcnArchName
- .iter()
- .position(|&x| x == 0)
- .unwrap();
- let gcn_arch_slice = unsafe {
- slice::from_raw_parts(device_pros.gcnArchName.as_ptr() as _, null_terminator + 1)
- };
- let device_name =
- if let Ok(Ok(name)) = CStr::from_bytes_with_nul(gcn_arch_slice).map(|x| x.to_str()) {
- name
- } else {
- return Err(io::Error::new(io::ErrorKind::Other, ""));
- };
- let dir = tempfile::tempdir()?;
- let mut spirv = NamedTempFile::new_in(&dir)?;
- let llvm = NamedTempFile::new_in(&dir)?;
- let spirv_il_u8 = unsafe {
- slice::from_raw_parts(
- spirv_il.as_ptr() as *const u8,
- spirv_il.len() * mem::size_of::<u32>(),
- )
- };
- spirv.write_all(spirv_il_u8)?;
- let llvm_spirv_path = match env::var("LLVM_SPIRV") {
- Ok(path) => Cow::Owned(path),
- Err(_) => Cow::Borrowed(LLVM_SPIRV),
- };
- let to_llvm_cmd = Command::new(&*llvm_spirv_path)
- .arg("-r")
- .arg("-o")
- .arg(llvm.path())
- .arg(spirv.path())
- .status()?;
- assert!(to_llvm_cmd.success());
- if cfg!(debug_assertions) {
- persist_file(llvm.path())?;
- }
- let linked_binary = NamedTempFile::new_in(&dir)?;
- let mut llvm_link = PathBuf::from(AMDGPU);
- llvm_link.push("llvm");
- llvm_link.push("bin");
- llvm_link.push("llvm-link");
- let mut linker_cmd = Command::new(&llvm_link);
- linker_cmd
- .arg("--only-needed")
- .arg("-o")
- .arg(linked_binary.path())
- .arg(llvm.path())
- .args(get_bitcode_paths(device_name));
- if cfg!(debug_assertions) {
- linker_cmd.arg("-v");
- }
- let status = linker_cmd.status()?;
- assert!(status.success());
- if cfg!(debug_assertions) {
- persist_file(linked_binary.path())?;
- }
- let mut ptx_lib_bitcode = NamedTempFile::new_in(&dir)?;
- let compiled_binary = NamedTempFile::new_in(&dir)?;
- let mut clang_exe = PathBuf::from(AMDGPU);
- clang_exe.push("llvm");
- clang_exe.push("bin");
- clang_exe.push("clang");
- let mut compiler_cmd = Command::new(&clang_exe);
- compiler_cmd
- .arg(format!("-mcpu={}", device_name))
- .arg("-ffp-contract=off")
- .arg("-nogpulib")
- .arg("-mno-wavefrontsize64")
- .arg("-O3")
- .arg("-Xlinker")
- .arg("--no-undefined")
- .arg("-target")
- .arg(AMDGPU_TARGET)
- .arg("-o")
- .arg(compiled_binary.path())
- .arg("-x")
- .arg("ir")
- .arg(linked_binary.path());
- if let Some((_, bitcode)) = ptx_lib {
- ptx_lib_bitcode.write_all(bitcode)?;
- compiler_cmd.arg(ptx_lib_bitcode.path());
- };
- if cfg!(debug_assertions) {
- compiler_cmd.arg("-v");
- }
- let status = compiler_cmd.status()?;
- assert!(status.success());
- let mut result = Vec::new();
- let compiled_bin_path = compiled_binary.path();
- let mut compiled_binary = File::open(compiled_bin_path)?;
- compiled_binary.read_to_end(&mut result)?;
- if cfg!(debug_assertions) {
- persist_file(compiled_bin_path)?;
- }
- Ok(result)
-}
-
fn persist_file(path: &Path) -> io::Result<()> {
let mut persistent = PathBuf::from("/tmp/zluda");
std::fs::create_dir_all(&persistent)?;
|