mirrors/ZLUDA
0
0
Fork 0
mirror of https://github.com/vosen/ZLUDA.git synced 2025-08-03 14:50:53 +00:00
Code Issues Projects Releases Packages Wiki Activity Actions 1

Add supprot for tex.level

Browse source
This commit is contained in:
Andrzej Janik 2024-02-29 12:25:57 +01:00
commit 7d501f8d08
6 changed files with 173 additions and 65 deletions
Download patch file Download diff file Expand all files Collapse all files

BIN
ptx/lib/zluda_ptx_impl.bc
View file

Binary file not shown.

40
ptx/lib/zluda_ptx_impl.cpp
View file

@ -327,10 +327,20 @@ extern "C"
hipTextureObject_t textureObject = ptr->textureObject; \ hipTextureObject_t textureObject = ptr->textureObject; \
return tex2D<HIP_CHANNEL_TYPE##4>(textureObject, float(coord.x), float(coord.y)).data; \ return tex2D<HIP_CHANNEL_TYPE##4>(textureObject, float(coord.x), float(coord.y)).data; \
} \ } \
HIP_CHANNEL_TYPE##4 ::Native_vec_ FUNC(tex_level_2d_v4_##CHANNEL_TYPE##_##COORD_TYPE)(struct textureReference GLOBAL_SPACE * ptr, HIP_COORD_TYPE##2 ::Native_vec_ coord, HIP_COORD_TYPE lod) \
{ \
hipTextureObject_t textureObject = ptr->textureObject; \
return tex2DLod<HIP_CHANNEL_TYPE##4>(textureObject, float(coord.x), float(coord.y), float(lod)).data; \
} \
HIP_CHANNEL_TYPE##4 ::Native_vec_ FUNC(tex_indirect_2d_v4_##CHANNEL_TYPE##_##COORD_TYPE)(uint64_t texobj, HIP_COORD_TYPE##2 ::Native_vec_ coord) \ HIP_CHANNEL_TYPE##4 ::Native_vec_ FUNC(tex_indirect_2d_v4_##CHANNEL_TYPE##_##COORD_TYPE)(uint64_t texobj, HIP_COORD_TYPE##2 ::Native_vec_ coord) \
{ \ { \
hipTextureObject_t textureObject = (hipTextureObject_t)texobj; \ hipTextureObject_t textureObject = (hipTextureObject_t)texobj; \
return tex2D<HIP_CHANNEL_TYPE##4>(textureObject, float(coord.x), float(coord.y)).data; \ return tex2D<HIP_CHANNEL_TYPE##4>(textureObject, float(coord.x), float(coord.y)).data; \
} \
HIP_CHANNEL_TYPE##4 ::Native_vec_ FUNC(tex_level_indirect_2d_v4_##CHANNEL_TYPE##_##COORD_TYPE)(uint64_t texobj, HIP_COORD_TYPE##2 ::Native_vec_ coord, HIP_COORD_TYPE lod) \
{ \
hipTextureObject_t textureObject = (hipTextureObject_t)texobj; \
return tex2DLod<HIP_CHANNEL_TYPE##4>(textureObject, float(coord.x), float(coord.y), float(lod)).data; \
} }
__device__ half4 __ockl_image_sampleh_2D(unsigned int CONSTANT_SPACE *i, unsigned int ADDRESS_SPACE_CONSTANT *s, float2::Native_vec_ c); __device__ half4 __ockl_image_sampleh_2D(unsigned int CONSTANT_SPACE *i, unsigned int ADDRESS_SPACE_CONSTANT *s, float2::Native_vec_ c);
@ -363,10 +373,20 @@ extern "C"
hipTextureObject_t textureObject = ptr->textureObject; \ hipTextureObject_t textureObject = ptr->textureObject; \
return tex3D<HIP_CHANNEL_TYPE##4>(textureObject, float(coord.x), float(coord.y), float(coord.z)).data; \ return tex3D<HIP_CHANNEL_TYPE##4>(textureObject, float(coord.x), float(coord.y), float(coord.z)).data; \
} \ } \
HIP_CHANNEL_TYPE##4 ::Native_vec_ FUNC(tex_level_3d_v4_##CHANNEL_TYPE##_##COORD_TYPE)(struct textureReference GLOBAL_SPACE * ptr, HIP_COORD_TYPE##4 ::Native_vec_ coord, HIP_COORD_TYPE lod) \
{ \
hipTextureObject_t textureObject = ptr->textureObject; \
return tex3DLod<HIP_CHANNEL_TYPE##4>(textureObject, float(coord.x), float(coord.y), float(coord.z), float(lod)).data; \
} \
HIP_CHANNEL_TYPE##4 ::Native_vec_ FUNC(tex_indirect_3d_v4_##CHANNEL_TYPE##_##COORD_TYPE)(uint64_t texobj, HIP_COORD_TYPE##4 ::Native_vec_ coord) \ HIP_CHANNEL_TYPE##4 ::Native_vec_ FUNC(tex_indirect_3d_v4_##CHANNEL_TYPE##_##COORD_TYPE)(uint64_t texobj, HIP_COORD_TYPE##4 ::Native_vec_ coord) \
{ \ { \
hipTextureObject_t textureObject = (hipTextureObject_t)texobj; \ hipTextureObject_t textureObject = (hipTextureObject_t)texobj; \
return tex3D<HIP_CHANNEL_TYPE##4>(textureObject, float(coord.x), float(coord.y), float(coord.z)).data; \ return tex3D<HIP_CHANNEL_TYPE##4>(textureObject, float(coord.x), float(coord.y), float(coord.z)).data; \
} \
HIP_CHANNEL_TYPE##4 ::Native_vec_ FUNC(tex_level_indirect_3d_v4_##CHANNEL_TYPE##_##COORD_TYPE)(uint64_t texobj, HIP_COORD_TYPE##4 ::Native_vec_ coord, HIP_COORD_TYPE lod) \
{ \
hipTextureObject_t textureObject = (hipTextureObject_t)texobj; \
return tex3DLod<HIP_CHANNEL_TYPE##4>(textureObject, float(coord.x), float(coord.y), float(coord.z), float(lod)).data; \
} }
__device__ half4 __ockl_image_sampleh_3D(unsigned int CONSTANT_SPACE *i, unsigned int ADDRESS_SPACE_CONSTANT *s, float4::Native_vec_ c); __device__ half4 __ockl_image_sampleh_3D(unsigned int CONSTANT_SPACE *i, unsigned int ADDRESS_SPACE_CONSTANT *s, float4::Native_vec_ c);
@ -399,10 +419,20 @@ extern "C"
hipTextureObject_t textureObject = ptr->textureObject; \ hipTextureObject_t textureObject = ptr->textureObject; \
return tex1DLayered<HIP_CHANNEL_TYPE##4>(textureObject, float(x), int(layer)).data; \ return tex1DLayered<HIP_CHANNEL_TYPE##4>(textureObject, float(x), int(layer)).data; \
} \ } \
HIP_CHANNEL_TYPE##4 ::Native_vec_ FUNC(tex_level_a1d_v4_##CHANNEL_TYPE##_##COORD_TYPE)(struct textureReference GLOBAL_SPACE * ptr, uint32_t layer, HIP_COORD_TYPE x, HIP_COORD_TYPE lod) \
{ \
hipTextureObject_t textureObject = ptr->textureObject; \
return tex1DLayeredLod<HIP_CHANNEL_TYPE##4>(textureObject, float(x), int(layer), float(lod)).data; \
} \
HIP_CHANNEL_TYPE##4 ::Native_vec_ FUNC(tex_indirect_a1d_v4_##CHANNEL_TYPE##_##COORD_TYPE)(uint64_t texobj, uint32_t layer, HIP_COORD_TYPE x) \ HIP_CHANNEL_TYPE##4 ::Native_vec_ FUNC(tex_indirect_a1d_v4_##CHANNEL_TYPE##_##COORD_TYPE)(uint64_t texobj, uint32_t layer, HIP_COORD_TYPE x) \
{ \ { \
hipTextureObject_t textureObject = (hipTextureObject_t)texobj; \ hipTextureObject_t textureObject = (hipTextureObject_t)texobj; \
return tex1DLayered<HIP_CHANNEL_TYPE##4>(textureObject, float(x), int(layer)).data; \ return tex1DLayered<HIP_CHANNEL_TYPE##4>(textureObject, float(x), int(layer)).data; \
} \
HIP_CHANNEL_TYPE##4 ::Native_vec_ FUNC(tex_level_indirect_a1d_v4_##CHANNEL_TYPE##_##COORD_TYPE)(uint64_t texobj, uint32_t layer, HIP_COORD_TYPE x, HIP_COORD_TYPE lod) \
{ \
hipTextureObject_t textureObject = (hipTextureObject_t)texobj; \
return tex1DLayeredLod<HIP_CHANNEL_TYPE##4>(textureObject, float(x), int(layer), float(lod)).data; \
} }
__device__ half4 __ockl_image_sampleh_1Da(unsigned int CONSTANT_SPACE *i, unsigned int ADDRESS_SPACE_CONSTANT *s, float2::Native_vec_ c); __device__ half4 __ockl_image_sampleh_1Da(unsigned int CONSTANT_SPACE *i, unsigned int ADDRESS_SPACE_CONSTANT *s, float2::Native_vec_ c);
@ -435,10 +465,20 @@ extern "C"
hipTextureObject_t textureObject = ptr->textureObject; \ hipTextureObject_t textureObject = ptr->textureObject; \
return tex2DLayered<HIP_CHANNEL_TYPE##4>(textureObject, float(x), float(y), int(layer)).data; \ return tex2DLayered<HIP_CHANNEL_TYPE##4>(textureObject, float(x), float(y), int(layer)).data; \
} \ } \
HIP_CHANNEL_TYPE##4 ::Native_vec_ FUNC(tex_level_a2d_v4_##CHANNEL_TYPE##_##COORD_TYPE)(struct textureReference GLOBAL_SPACE * ptr, uint32_t layer, HIP_COORD_TYPE x, HIP_COORD_TYPE y, HIP_COORD_TYPE lod) \
{ \
hipTextureObject_t textureObject = ptr->textureObject; \
return tex2DLayeredLod<HIP_CHANNEL_TYPE##4>(textureObject, float(x), float(y), int(layer), float(lod)).data; \
} \
HIP_CHANNEL_TYPE##4 ::Native_vec_ FUNC(tex_indirect_a2d_v4_##CHANNEL_TYPE##_##COORD_TYPE)(uint64_t texobj, uint32_t layer, HIP_COORD_TYPE x, HIP_COORD_TYPE y) \ HIP_CHANNEL_TYPE##4 ::Native_vec_ FUNC(tex_indirect_a2d_v4_##CHANNEL_TYPE##_##COORD_TYPE)(uint64_t texobj, uint32_t layer, HIP_COORD_TYPE x, HIP_COORD_TYPE y) \
{ \ { \
hipTextureObject_t textureObject = (hipTextureObject_t)texobj; \ hipTextureObject_t textureObject = (hipTextureObject_t)texobj; \
return tex2DLayered<HIP_CHANNEL_TYPE##4>(textureObject, float(x), float(y), int(layer)).data; \ return tex2DLayered<HIP_CHANNEL_TYPE##4>(textureObject, float(x), float(y), int(layer)).data; \
} \
HIP_CHANNEL_TYPE##4 ::Native_vec_ FUNC(tex_level_indirect_a2d_v4_##CHANNEL_TYPE##_##COORD_TYPE)(uint64_t texobj, uint32_t layer, HIP_COORD_TYPE x, HIP_COORD_TYPE y, HIP_COORD_TYPE lod) \
{ \
hipTextureObject_t textureObject = (hipTextureObject_t)texobj; \
return tex2DLayeredLod<HIP_CHANNEL_TYPE##4>(textureObject, float(x), float(y), int(layer), float(lod)).data; \
} }
__device__ half4 __ockl_image_sampleh_2Da(unsigned int CONSTANT_SPACE *i, unsigned int ADDRESS_SPACE_CONSTANT *s, float4::Native_vec_ c); __device__ half4 __ockl_image_sampleh_2Da(unsigned int CONSTANT_SPACE *i, unsigned int ADDRESS_SPACE_CONSTANT *s, float4::Native_vec_ c);

21
ptx/src/ast.rs
View file

@ -459,8 +459,8 @@ pub enum Instruction<P: ArgParams> {
Membar { Membar {
level: MemScope, level: MemScope,
}, },
Tex(TexDetails, Arg4Tex<P>), Tex(TexDetails, Arg5Tex<P>),
Suld(SurfaceDetails, Arg4Tex<P>), Suld(SurfaceDetails, Arg5Tex<P>),
Sust(SurfaceDetails, Arg4Sust<P>), Sust(SurfaceDetails, Arg4Sust<P>),
Shfl(ShflMode, Arg5Shfl<P>), Shfl(ShflMode, Arg5Shfl<P>),
Shf(FunnelShift, Arg4<P>), Shf(FunnelShift, Arg4<P>),
@ -616,13 +616,6 @@ pub struct Arg4Setp<P: ArgParams> {
pub src2: P::Operand, pub src2: P::Operand,
} }
pub struct Arg4Tex<P: ArgParams> {
pub dst: P::Operand,
pub image: P::Operand,
pub layer: Option<P::Operand>,
pub coordinates: P::Operand,
}
pub struct Arg4Sust<P: ArgParams> { pub struct Arg4Sust<P: ArgParams> {
pub image: P::Operand, pub image: P::Operand,
pub coordinates: P::Operand, pub coordinates: P::Operand,
@ -638,6 +631,14 @@ pub struct Arg5<P: ArgParams> {
pub src4: P::Operand, pub src4: P::Operand,
} }
pub struct Arg5Tex<P: ArgParams> {
pub dst: P::Operand,
pub image: P::Operand,
pub layer: Option<P::Operand>,
pub coordinates: P::Operand,
pub lod: Option<P::Operand>,
}
pub struct Arg5Setp<P: ArgParams> { pub struct Arg5Setp<P: ArgParams> {
pub dst1: P::Id, pub dst1: P::Id,
pub dst2: Option<P::Id>, pub dst2: Option<P::Id>,
@ -1316,7 +1317,7 @@ pub enum TuningDirective {
MaxNtid(u32, u32, u32), MaxNtid(u32, u32, u32),
ReqNtid(u32, u32, u32), ReqNtid(u32, u32, u32),
MinNCtaPerSm(u32), MinNCtaPerSm(u32),
Noreturn Noreturn,
} }
#[repr(u8)] #[repr(u8)]

62
ptx/src/ptx.lalrpop
View file

@ -97,6 +97,7 @@ match {
".l", ".l",
".le", ".le",
".leu", ".leu",
".level",
".lo", ".lo",
".loc", ".loc",
".local", ".local",
@ -2065,11 +2066,23 @@ InstMembar: ast::Instruction<ast::ParsedArgParams<'input>> = {
// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#texture-instructions-tex // https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#texture-instructions-tex
InstTex: ast::Instruction<ast::ParsedArgParams<'input>> = { InstTex: ast::Instruction<ast::ParsedArgParams<'input>> = {
"tex" <geometry: UnlayeredTextureGeometry> ".v4" <channel_type:TextureChannelType> <coordinate_type:TextureCoordinateType> <dst:DstOperandVec> "," "[" <image:Operand> "," <coordinates:SrcOperandVec> "]" => { "tex" <geometry: UnlayeredTextureGeometry> ".v4" <channel_type:TextureChannelType> <coordinate_type:TextureCoordinateType> <dst:DstOperandVec> "," "[" <image:Operand> "," <coordinates:SrcOperandVec> "]" => {
let args = ast::Arg4Tex { let args = ast::Arg5Tex {
dst, dst,
image, image,
coordinates, coordinates,
layer: None layer: None,
lod: None
};
let details = ast::TexDetails { geometry, channel_type, coordinate_type, direct: false };
ast::Instruction::Tex(details, args)
},
"tex" ".level" <geometry: UnlayeredTextureGeometry> ".v4" <channel_type:TextureChannelType> <coordinate_type:TextureCoordinateType> <dst:DstOperandVec> "," "[" <image:Operand> "," <coordinates:SrcOperandVec> "]" "," <lod:Operand> => {
let args = ast::Arg5Tex {
dst,
image,
coordinates,
layer: None,
lod: Some(lod)
}; };
let details = ast::TexDetails { geometry, channel_type, coordinate_type, direct: false }; let details = ast::TexDetails { geometry, channel_type, coordinate_type, direct: false };
ast::Instruction::Tex(details, args) ast::Instruction::Tex(details, args)
@ -2080,11 +2093,25 @@ InstTex: ast::Instruction<ast::ParsedArgParams<'input>> = {
"tex" ".a1d" ".v4" <channel_type:TextureChannelType> <coordinate_type:TextureCoordinateType> "tex" ".a1d" ".v4" <channel_type:TextureChannelType> <coordinate_type:TextureCoordinateType>
<dst:DstOperandVec> "," "[" <image:SrcOperand> "," "{" <layer:SrcOperand> "," <x:RegOrImmediate> "}" "]" => { <dst:DstOperandVec> "," "[" <image:SrcOperand> "," "{" <layer:SrcOperand> "," <x:RegOrImmediate> "}" "]" => {
let geometry = ast::TextureGeometry::Array1D; let geometry = ast::TextureGeometry::Array1D;
let args = ast::Arg4Tex { let args = ast::Arg5Tex {
dst, dst,
image, image,
coordinates: ast::Operand::VecPack(vec![x]), coordinates: ast::Operand::VecPack(vec![x]),
layer: Some(layer) layer: Some(layer),
lod: None
};
let details = ast::TexDetails { geometry, channel_type, coordinate_type, direct: false };
ast::Instruction::Tex(details, args)
},
"tex" ".level" ".a1d" ".v4" <channel_type:TextureChannelType> <coordinate_type:TextureCoordinateType>
<dst:DstOperandVec> "," "[" <image:SrcOperand> "," "{" <layer:SrcOperand> "," <x:RegOrImmediate> "}" "]" "," <lod:Operand> => {
let geometry = ast::TextureGeometry::Array1D;
let args = ast::Arg5Tex {
dst,
image,
coordinates: ast::Operand::VecPack(vec![x]),
layer: Some(layer),
lod: Some(lod)
}; };
let details = ast::TexDetails { geometry, channel_type, coordinate_type, direct: false }; let details = ast::TexDetails { geometry, channel_type, coordinate_type, direct: false };
ast::Instruction::Tex(details, args) ast::Instruction::Tex(details, args)
@ -2092,11 +2119,25 @@ InstTex: ast::Instruction<ast::ParsedArgParams<'input>> = {
"tex" ".a2d" ".v4" <channel_type:TextureChannelType> <coordinate_type:TextureCoordinateType> "tex" ".a2d" ".v4" <channel_type:TextureChannelType> <coordinate_type:TextureCoordinateType>
<dst:DstOperandVec> "," "[" <image:SrcOperand> "," "{" <layer:SrcOperand> "," <x:RegOrImmediate> "," <y:RegOrImmediate> "," RegOrImmediate "}" "]" => { <dst:DstOperandVec> "," "[" <image:SrcOperand> "," "{" <layer:SrcOperand> "," <x:RegOrImmediate> "," <y:RegOrImmediate> "," RegOrImmediate "}" "]" => {
let geometry = ast::TextureGeometry::Array2D; let geometry = ast::TextureGeometry::Array2D;
let args = ast::Arg4Tex { let args = ast::Arg5Tex {
dst, dst,
image, image,
coordinates: ast::Operand::VecPack(vec![x, y]), coordinates: ast::Operand::VecPack(vec![x, y]),
layer: Some(layer) layer: Some(layer),
lod: None
};
let details = ast::TexDetails { geometry, channel_type, coordinate_type, direct: false };
ast::Instruction::Tex(details, args)
},
"tex" ".level" ".a2d" ".v4" <channel_type:TextureChannelType> <coordinate_type:TextureCoordinateType>
<dst:DstOperandVec> "," "[" <image:SrcOperand> "," "{" <layer:SrcOperand> "," <x:RegOrImmediate> "," <y:RegOrImmediate> "," RegOrImmediate "}" "]" "," <lod:Operand> => {
let geometry = ast::TextureGeometry::Array2D;
let args = ast::Arg5Tex {
dst,
image,
coordinates: ast::Operand::VecPack(vec![x, y]),
layer: Some(layer),
lod: Some(lod)
}; };
let details = ast::TexDetails { geometry, channel_type, coordinate_type, direct: false }; let details = ast::TexDetails { geometry, channel_type, coordinate_type, direct: false };
ast::Instruction::Tex(details, args) ast::Instruction::Tex(details, args)
@ -2106,33 +2147,36 @@ InstTex: ast::Instruction<ast::ParsedArgParams<'input>> = {
// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#surface-instructions-suld // https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#surface-instructions-suld
InstSuld: ast::Instruction<ast::ParsedArgParams<'input>> = { InstSuld: ast::Instruction<ast::ParsedArgParams<'input>> = {
"suld" ".b" <geometry: UnlayeredTextureGeometry> <vector:VectorPrefix?> <type_:AnyBitType> ".trap" <dst:DstOperandVec> "," "[" <image:Operand> "," <coordinates:SrcOperandVec> "]" => { "suld" ".b" <geometry: UnlayeredTextureGeometry> <vector:VectorPrefix?> <type_:AnyBitType> ".trap" <dst:DstOperandVec> "," "[" <image:Operand> "," <coordinates:SrcOperandVec> "]" => {
let args = ast::Arg4Tex { let args = ast::Arg5Tex {
dst, dst,
image, image,
coordinates, coordinates,
layer: None, layer: None,
lod: None,
}; };
let details = ast::SurfaceDetails { geometry, vector, type_, direct: false, }; let details = ast::SurfaceDetails { geometry, vector, type_, direct: false, };
ast::Instruction::Suld(details, args) ast::Instruction::Suld(details, args)
}, },
"suld" ".b" ".a1d" <vector:VectorPrefix?> <type_:AnyBitType> ".trap" <dst:DstOperandVec> "," "[" <image:Operand> "," "{" <layer:SrcOperand> "," <x:RegOrImmediate> "}" "]" => { "suld" ".b" ".a1d" <vector:VectorPrefix?> <type_:AnyBitType> ".trap" <dst:DstOperandVec> "," "[" <image:Operand> "," "{" <layer:SrcOperand> "," <x:RegOrImmediate> "}" "]" => {
let geometry = ast::TextureGeometry::Array1D; let geometry = ast::TextureGeometry::Array1D;
let args = ast::Arg4Tex { let args = ast::Arg5Tex {
dst, dst,
image, image,
coordinates: ast::Operand::VecPack(vec![x]), coordinates: ast::Operand::VecPack(vec![x]),
layer: Some(layer), layer: Some(layer),
lod: None,
}; };
let details = ast::SurfaceDetails { geometry, vector, type_, direct: false, }; let details = ast::SurfaceDetails { geometry, vector, type_, direct: false, };
ast::Instruction::Suld(details, args) ast::Instruction::Suld(details, args)
}, },
"suld" ".b" ".a2d" <vector:VectorPrefix?> <type_:AnyBitType> ".trap" <dst:DstOperandVec> "," "[" <image:Operand> "," "{" <layer:SrcOperand> "," <x:RegOrImmediate> "," <y:RegOrImmediate> "," RegOrImmediate "}" "]" => { "suld" ".b" ".a2d" <vector:VectorPrefix?> <type_:AnyBitType> ".trap" <dst:DstOperandVec> "," "[" <image:Operand> "," "{" <layer:SrcOperand> "," <x:RegOrImmediate> "," <y:RegOrImmediate> "," RegOrImmediate "}" "]" => {
let geometry = ast::TextureGeometry::Array2D; let geometry = ast::TextureGeometry::Array2D;
let args = ast::Arg4Tex { let args = ast::Arg5Tex {
dst, dst,
image, image,
coordinates: ast::Operand::VecPack(vec![x, y]), coordinates: ast::Operand::VecPack(vec![x, y]),
layer: Some(layer), layer: Some(layer),
lod: None,
}; };
let details = ast::SurfaceDetails { geometry, vector, type_, direct: false, }; let details = ast::SurfaceDetails { geometry, vector, type_, direct: false, };
ast::Instruction::Suld(details, args) ast::Instruction::Suld(details, args)

6
ptx/src/raytracing.rs
View file

@ -2996,11 +2996,12 @@ fn convert_optix_builtin_variable_and_attribute_access_single_function<'input>(
} }
Statement::Instruction(ast::Instruction::Tex( Statement::Instruction(ast::Instruction::Tex(
tex, tex,
ast::Arg4Tex { ast::Arg5Tex {
dst, dst,
image, image,
layer, layer,
coordinates, coordinates,
lod,
}, },
)) => { )) => {
if let Some(StateSpaceRemapping::ToBlock(id, ast::StateSpace::Global, offset)) = if let Some(StateSpaceRemapping::ToBlock(id, ast::StateSpace::Global, offset)) =
@ -3014,11 +3015,12 @@ fn convert_optix_builtin_variable_and_attribute_access_single_function<'input>(
)?; )?;
result.push(Statement::Instruction(ast::Instruction::Tex( result.push(Statement::Instruction(ast::Instruction::Tex(
tex, tex,
ast::Arg4Tex { ast::Arg5Tex {
dst, dst,
image, image,
layer, layer,
coordinates, coordinates,
lod,
}, },
))); )));
} else { } else {

29
ptx/src/translate.rs
View file

@ -2775,9 +2775,14 @@ fn replace_instructions_with_builtins_impl<'input>(
} }
Statement::Instruction(ast::Instruction::Tex(tex, arg)) => { Statement::Instruction(ast::Instruction::Tex(tex, arg)) => {
let geometry = tex.geometry.as_ptx(); let geometry = tex.geometry.as_ptx();
let op_name = if arg.lod.is_none() {
"tex"
} else {
"tex_level"
};
let fn_name = [ let fn_name = [
ZLUDA_PTX_PREFIX, ZLUDA_PTX_PREFIX,
"tex", op_name,
tex.suffix(), tex.suffix(),
"_", "_",
geometry, geometry,
@ -8089,7 +8094,7 @@ fn texture_geometry_to_vec_length(geometry: ast::TextureGeometry) -> u8 {
} }
} }
impl<T: ArgParamsEx> ast::Arg4Tex<T> { impl<T: ArgParamsEx> ast::Arg5Tex<T> {
fn map<U: ArgParamsEx, V: ArgumentMapVisitor<T, U>>( fn map<U: ArgParamsEx, V: ArgumentMapVisitor<T, U>>(
self, self,
visitor: &mut V, visitor: &mut V,
@ -8097,7 +8102,7 @@ impl<T: ArgParamsEx> ast::Arg4Tex<T> {
geometry: ast::TextureGeometry, geometry: ast::TextureGeometry,
value_type: ast::Type, value_type: ast::Type,
coordinate_type: ast::ScalarType, coordinate_type: ast::ScalarType,
) -> Result<ast::Arg4Tex<U>, TranslateError> { ) -> Result<ast::Arg5Tex<U>, TranslateError> {
let dst = visitor.operand( let dst = visitor.operand(
ArgumentDescriptor { ArgumentDescriptor {
op: self.dst, op: self.dst,
@ -8144,11 +8149,27 @@ impl<T: ArgParamsEx> ast::Arg4Tex<T> {
&ast::Type::Vector(coordinate_type, coord_length), &ast::Type::Vector(coordinate_type, coord_length),
ast::StateSpace::Reg, ast::StateSpace::Reg,
)?; )?;
Ok(ast::Arg4Tex { let lod = self
.lod
.map(|lod| {
visitor.operand(
ArgumentDescriptor {
op: lod,
is_dst: false,
is_memory_access: false,
non_default_implicit_conversion: None,
},
&ast::Type::Scalar(coordinate_type),
ast::StateSpace::Reg,
)
})
.transpose()?;
Ok(ast::Arg5Tex {
dst, dst,
image, image,
layer, layer,
coordinates, coordinates,
lod,
}) })
} }
} }