mirrors/ZLUDA
0
0
Fork 0
mirror of https://github.com/vosen/ZLUDA.git synced 2025-10-02 06:10:06 +00:00
Code Issues Projects Releases Packages Wiki Activity Actions 1

Add failing test

Browse source
This commit is contained in:
Andrzej Janik 2025-09-11 00:24:28 +00:00
commit c8f7bea30f
5 changed files with 34 additions and 9 deletions
Download patch file Download diff file Expand all files Collapse all files

2
ptx/src/pass/insert_implicit_conversions2.rs
View file

@ -186,7 +186,7 @@ fn default_implicit_conversion_space(
ast::Type::Scalar(ast::ScalarType::B32) ast::Type::Scalar(ast::ScalarType::B32)
| ast::Type::Scalar(ast::ScalarType::U32) | ast::Type::Scalar(ast::ScalarType::U32)
| ast::Type::Scalar(ast::ScalarType::S32) => match instruction_space { | ast::Type::Scalar(ast::ScalarType::S32) => match instruction_space {
ast::StateSpace::Const | ast::StateSpace::Local | ast::StateSpace::Shared => { ast::StateSpace::Local | ast::StateSpace::Shared => {
Ok(Some(ConversionKind::BitToPtr)) Ok(Some(ConversionKind::BitToPtr))
} }
_ => Err(error_mismatched_type()), _ => Err(error_mismatched_type()),

5
ptx/src/pass/llvm/emit.rs
View file

@ -693,6 +693,7 @@ impl<'a> MethodEmitContext<'a> {
} }
} }
(ast::Type::Vector(..), ast::Type::Scalar(..)) (ast::Type::Vector(..), ast::Type::Scalar(..))
| (ast::Type::Scalar(..), ast::Type::Vector(..))
| (ast::Type::Scalar(..), ast::Type::Array(..)) | (ast::Type::Scalar(..), ast::Type::Array(..))
| (ast::Type::Array(..), ast::Type::Scalar(..)) => { | (ast::Type::Array(..), ast::Type::Scalar(..)) => {
let dst_type = get_type(self.context, to_type)?; let dst_type = get_type(self.context, to_type)?;
@ -701,7 +702,7 @@ impl<'a> MethodEmitContext<'a> {
}); });
Ok(()) Ok(())
} }
_ => todo!(), _ => return Err(error_todo()),
} }
} }
@ -2409,7 +2410,7 @@ impl<'a> MethodEmitContext<'a> {
(control >> 12) & 0b1111, (control >> 12) & 0b1111,
]; ];
if components.iter().any(|&c| c > 7) { if components.iter().any(|&c| c > 7) {
return Err(TranslateError::Todo("".to_string())); return Err(error_todo());
} }
let u32_type = get_scalar_type(self.context, ast::ScalarType::U32); let u32_type = get_scalar_type(self.context, ast::ScalarType::U32);
let v4u8_type = get_type(self.context, &ast::Type::Vector(4, ast::ScalarType::U8))?; let v4u8_type = get_type(self.context, &ast::Type::Vector(4, ast::ScalarType::U8))?;

13
ptx/src/pass/llvm/mod.rs
View file

@ -169,8 +169,9 @@ fn get_scalar_type(context: LLVMContextRef, type_: ast::ScalarType) -> LLVMTypeR
ast::ScalarType::F32 => unsafe { LLVMFloatTypeInContext(context) }, ast::ScalarType::F32 => unsafe { LLVMFloatTypeInContext(context) },
ast::ScalarType::F64 => unsafe { LLVMDoubleTypeInContext(context) }, ast::ScalarType::F64 => unsafe { LLVMDoubleTypeInContext(context) },
ast::ScalarType::BF16 => unsafe { LLVMBFloatTypeInContext(context) }, ast::ScalarType::BF16 => unsafe { LLVMBFloatTypeInContext(context) },
ast::ScalarType::U16x2 => todo!(), ast::ScalarType::U16x2 | ast::ScalarType::S16x2 => unsafe {
ast::ScalarType::S16x2 => todo!(), LLVMVectorType(LLVMInt16TypeInContext(context), 2)
},
ast::ScalarType::F16x2 => unsafe { LLVMVectorType(LLVMHalfTypeInContext(context), 2) }, ast::ScalarType::F16x2 => unsafe { LLVMVectorType(LLVMHalfTypeInContext(context), 2) },
ast::ScalarType::BF16x2 => unsafe { LLVMVectorType(LLVMBFloatTypeInContext(context), 2) }, ast::ScalarType::BF16x2 => unsafe { LLVMVectorType(LLVMBFloatTypeInContext(context), 2) },
} }
@ -180,14 +181,14 @@ fn get_state_space(space: ast::StateSpace) -> Result<u32, TranslateError> {
match space { match space {
ast::StateSpace::Reg => Ok(PRIVATE_ADDRESS_SPACE), ast::StateSpace::Reg => Ok(PRIVATE_ADDRESS_SPACE),
ast::StateSpace::Generic => Ok(GENERIC_ADDRESS_SPACE), ast::StateSpace::Generic => Ok(GENERIC_ADDRESS_SPACE),
ast::StateSpace::Param => Err(TranslateError::Todo("".to_string())), ast::StateSpace::Param => Err(error_todo()),
ast::StateSpace::ParamEntry => Ok(CONSTANT_ADDRESS_SPACE), ast::StateSpace::ParamEntry => Ok(CONSTANT_ADDRESS_SPACE),
ast::StateSpace::ParamFunc => Err(TranslateError::Todo("".to_string())), ast::StateSpace::ParamFunc => Err(error_todo()),
ast::StateSpace::Local => Ok(PRIVATE_ADDRESS_SPACE), ast::StateSpace::Local => Ok(PRIVATE_ADDRESS_SPACE),
ast::StateSpace::Global => Ok(GLOBAL_ADDRESS_SPACE), ast::StateSpace::Global => Ok(GLOBAL_ADDRESS_SPACE),
ast::StateSpace::Const => Ok(CONSTANT_ADDRESS_SPACE), ast::StateSpace::Const => Ok(CONSTANT_ADDRESS_SPACE),
ast::StateSpace::Shared => Ok(SHARED_ADDRESS_SPACE), ast::StateSpace::Shared => Ok(SHARED_ADDRESS_SPACE),
ast::StateSpace::SharedCta => Err(TranslateError::Todo("".to_string())), ast::StateSpace::SharedCta => Err(error_todo()),
ast::StateSpace::SharedCluster => Err(TranslateError::Todo("".to_string())), ast::StateSpace::SharedCluster => Err(error_todo()),
} }
} }

1
ptx/src/test/spirv_run/mod.rs
View file

@ -339,6 +339,7 @@ test_ptx!(
[0x8e2da590u32, 0xedeaee14, 0x248a9f70], [0x8e2da590u32, 0xedeaee14, 0x248a9f70],
[613065134u32] [613065134u32]
); );
test_ptx!(param_is_addressable, [0xDEAD], [0u64]);
test_ptx!(assertfail); test_ptx!(assertfail);
// TODO: not yet supported // TODO: not yet supported

22
ptx/src/test/spirv_run/param_is_addressable.ptx Normal file
View file

@ -0,0 +1,22 @@
.version 6.5
.target sm_30
.address_size 64
.visible .entry param_is_addressable(
.param .u64 input,
.param .u64 output
)
{
.reg .u64 in_addr;
.reg .u64 out_addr;
.reg .b64 temp;
ld.param.u64 in_addr, [input];
ld.param.u64 out_addr, [output];
mov.b64 temp, input;
ld.param.b64 temp, [temp];
sub.u64 temp, temp, in_addr;
st.u64 [out_addr], temp;
ret;
}