diff --git a/src/airconv/air_operations.cpp b/src/airconv/air_operations.cpp
deleted file mode 100644
index 795d0cbd5..000000000
--- a/src/airconv/air_operations.cpp
+++ /dev/null
@@ -1,493 +0,0 @@
-#include "air_operations.hpp"
-
-namespace dxmt::air {
-
-ReaderIO<AIRBuilderContext, AIRBuilderContext> get_context() {
-  return ReaderIO<AIRBuilderContext, AIRBuilderContext>(
-    [](struct AIRBuilderContext ctx) { return ctx; }
-  );
-};
-
-AIRBuilderResult get_int(uint32_t value) {
-  return make_op([=](auto ctx) { return ctx.builder.getInt32(value); });
-};
-
-AIRBuilderResult get_float(float value) {
-  return make_op([=](auto ctx) {
-    return llvm::ConstantFP::get(ctx.llvm, llvm::APFloat{value});
-  });
-};
-
-AIRBuilderResult call_unpack_impl(
-  std::string op, pvalue src, llvm::Type *src_type, llvm::Type *dst_type
-) {
-  using namespace llvm;
-  auto ctx = co_yield get_context();
-  auto &context = ctx.llvm;
-  auto &module = ctx.module;
-  auto att = AttributeList::get(
-    context, {{~0U, Attribute::get(context, Attribute::AttrKind::NoUnwind)},
-              {~0U, Attribute::get(context, Attribute::AttrKind::WillReturn)},
-              {~0U, Attribute::get(context, Attribute::AttrKind::MustProgress)},
-              {~0U, Attribute::get(context, Attribute::AttrKind::NoFree)},
-              {~0U, Attribute::get(context, Attribute::AttrKind::NoSync)},
-              {~0U, Attribute::get(context, Attribute::AttrKind::ReadNone)}}
-  );
-  auto fn = (module.getOrInsertFunction(
-    "air.unpack." + op, llvm::FunctionType::get(dst_type, {src_type}, false),
-    att
-  ));
-  co_return ctx.builder.CreateCall(fn, {src});
-};
-
-AIRBuilderResult load_from_device_buffer(
-  llvm::Type *value_type, pvalue base_addr, pvalue aligned_offset,
-  uint32_t element_offset, uint32_t align
-) {
-  auto ctx = co_yield get_context();
-  co_return ctx.builder.CreateAlignedLoad(
-    value_type,
-    ctx.builder.CreateConstGEP1_32(
-      value_type,
-      ctx.builder.CreateBitCast(
-        ctx.builder.CreateGEP(ctx.types._byte, base_addr, {aligned_offset}),
-        value_type->getPointerTo((uint32_t)AddressSpace::device)
-      ),
-      element_offset
-    ),
-    llvm::MaybeAlign(align)
-  );
-};
-
-AIRBuilderResult unpack_fvec4_from_addr(
-  MTLAttributeFormat format, pvalue base_addr, pvalue aligned_offset
-) {
-  auto ctx = co_yield get_context();
-  auto &types = ctx.types;
-  auto &builder = ctx.builder;
-
-  if (format == MTLAttributeFormat::UShort4Normalized) {
-    pvalue _xy = co_yield call_unpack_impl(
-        "unorm2x16.v2f32", co_yield load_from_device_buffer(types._int, base_addr, aligned_offset, 0, 2), types._int,
-        types._float2
-    );
-    pvalue _zw = co_yield call_unpack_impl(
-        "unorm2x16.v2f32", co_yield load_from_device_buffer(types._int, base_addr, aligned_offset, 1, 2), types._int,
-        types._float2
-    );
-    co_return builder.CreateShuffleVector(_xy, _zw, {0, 1, 2, 3});
-  }
-  if (format == MTLAttributeFormat::Short4Normalized) {
-    pvalue _xy = co_yield call_unpack_impl(
-        "snorm2x16.v2f32", co_yield load_from_device_buffer(types._int, base_addr, aligned_offset, 0, 2), types._int,
-        types._float2
-    );
-    pvalue _zw = co_yield call_unpack_impl(
-        "snorm2x16.v2f32", co_yield load_from_device_buffer(types._int, base_addr, aligned_offset, 1, 2), types._int,
-        types._float2
-    );
-    co_return builder.CreateShuffleVector(_xy, _zw, {0, 1, 2, 3});
-  }
-
-  std::string op;
-  llvm::Type *src_type = nullptr;
-  uint32_t align = 0;
-  llvm::Type *dst_type = nullptr;
-  switch (format) {
-  case MTLAttributeFormat::UCharNormalized:
-    op = "unorm1x8.f32";
-    src_type = types._byte;
-    align = 1;
-    dst_type = types._float;
-    break;
-  case MTLAttributeFormat::UChar2Normalized:
-    op = "unorm2x8.v2f32";
-    src_type = types._short;
-    align = 1;
-    dst_type = types._float2;
-    break;
-  case MTLAttributeFormat::UChar4Normalized:
-  case MTLAttributeFormat::UChar4Normalized_BGRA:
-    op = "unorm4x8.v4f32";
-    src_type = types._int;
-    align = 1;
-    dst_type = types._float4;
-    break;
-  case MTLAttributeFormat::CharNormalized:
-    op = "snorm1x8.f32";
-    src_type = types._byte;
-    align = 1;
-    dst_type = types._float;
-    break;
-  case MTLAttributeFormat::Char2Normalized:
-    op = "snorm2x8.v2f32";
-    src_type = types._short;
-    align = 1;
-    dst_type = types._float2;
-    break;
-  case MTLAttributeFormat::Char4Normalized:
-    op = "snorm4x8.v4f32";
-    src_type = types._int;
-    align = 1;
-    dst_type = types._float4;
-    break;
-  case MTLAttributeFormat::UShortNormalized:
-    op = "unorm1x16.f32";
-    src_type = types._short;
-    align = 2;
-    dst_type = types._float;
-    break;
-  case MTLAttributeFormat::UShort2Normalized:
-    op = "unorm2x16.v2f32";
-    src_type = types._int;
-    align = 2;
-    dst_type = types._float2;
-    break;
-  case MTLAttributeFormat::ShortNormalized:
-    op = "snorm1x16.f32";
-    src_type = types._short;
-    align = 2;
-    dst_type = types._float;
-    break;
-  case MTLAttributeFormat::Short2Normalized:
-    op = "snorm2x16.v2f32";
-    src_type = types._int;
-    align = 2;
-    dst_type = types._float2;
-    break;
-  case MTLAttributeFormat::UInt1010102Normalized:
-    op = "unorm.rgb10a2.v4f32";
-    src_type = types._int;
-    align = 4;
-    dst_type = types._float4;
-    break;
-  case MTLAttributeFormat::Int1010102Normalized:
-    assert(0 && "unused");
-    // op = "unorm.rgb10a2.v4f32";
-    // src_type = types._int;
-    // align = 4;
-    // dst_type = types._float4;
-    break;
-  case MTLAttributeFormat::FloatRG11B10:
-    op = "unorm.rg11b10f.v3f32";
-    src_type = types._int;
-    align = 4;
-    dst_type = types._float3;
-    break;
-  case MTLAttributeFormat::FloatRGB9E5:
-    op = "unorm.rgb9e5.v3f32";
-    src_type = types._int;
-    align = 4;
-    dst_type = types._float3;
-    break;
-  default:
-    break;
-  }
-  assert(src_type && dst_type);
-  pvalue ret = co_yield call_unpack_impl(
-    op,
-    co_yield load_from_device_buffer(
-      src_type, base_addr, aligned_offset, 0, align
-    ),
-    src_type, dst_type
-  );
-  if (format == MTLAttributeFormat::UChar4Normalized_BGRA) {
-    ret = builder.CreateShuffleVector(ret, {2, 1, 0, 3});
-  }
-  if (dst_type == types._float) {
-    ret = builder.CreateInsertElement(
-      llvm::PoisonValue::get(types._float4), ret, (int)0
-    );
-    ret = builder.CreateInsertElement(ret, co_yield get_float(0), (int)1);
-    ret = builder.CreateInsertElement(ret, co_yield get_float(0), (int)2);
-    ret = builder.CreateInsertElement(ret, co_yield get_float(1), (int)3);
-  } else if (dst_type == types._float2) {
-    ret = builder.CreateShuffleVector(ret, {0, 1, -1, -1});
-    ret = builder.CreateInsertElement(ret, co_yield get_float(0), (int)2);
-    ret = builder.CreateInsertElement(ret, co_yield get_float(1), (int)3);
-  } else if (dst_type == types._float3) {
-    ret = builder.CreateShuffleVector(ret, {0, 1, 2, -1});
-    ret = builder.CreateInsertElement(ret, co_yield get_float(1), (int)3);
-  }
-  co_return ret;
-};
-
-AIRBuilderResult pull_vec4_from_addr_checked(
-  MTLAttributeFormat format, pvalue base_addr, pvalue byte_offset
-) {
-  auto ctx = co_yield get_context();
-  auto &types = ctx.types;
-  auto &builder = ctx.builder;
-
-  pvalue value = nullptr;
-
-  switch (format) {
-  case MTLAttributeFormat::Char:
-    value = co_yield load_from_device_buffer(
-      types._byte, base_addr, byte_offset, 0, 1
-    );
-    value = builder.CreateSExt(value, types._int);
-    value = builder.CreateInsertElement(
-      llvm::PoisonValue::get(types._int4), value, (int)0
-    );
-    value = builder.CreateInsertElement(value, co_yield get_int(0), (int)1);
-    value = builder.CreateInsertElement(value, co_yield get_int(0), (int)2);
-    value = builder.CreateInsertElement(value, co_yield get_int(1), (int)3);
-    break;
-  case MTLAttributeFormat::Char2:
-    value = co_yield load_from_device_buffer(
-      types._char2, base_addr, byte_offset, 0, 1
-    );
-    value = builder.CreateSExt(value, types._int2);
-    value = builder.CreateShuffleVector(value, {0, 1, -1, -1});
-    value = builder.CreateInsertElement(value, co_yield get_int(0), (int)2);
-    value = builder.CreateInsertElement(value, co_yield get_int(1), (int)3);
-    break;
-  case MTLAttributeFormat::Char4:
-    value = co_yield load_from_device_buffer(
-      types._char4, base_addr, byte_offset, 0, 1
-    );
-    value = builder.CreateSExt(value, types._int4);
-    break;
-  case MTLAttributeFormat::UChar:
-    value = co_yield load_from_device_buffer(
-      types._byte, base_addr, byte_offset, 0, 1
-    );
-    value = builder.CreateZExt(value, types._int);
-    value = builder.CreateInsertElement(
-      llvm::PoisonValue::get(types._int4), value, (int)0
-    );
-    value = builder.CreateInsertElement(value, co_yield get_int(0), (int)1);
-    value = builder.CreateInsertElement(value, co_yield get_int(0), (int)2);
-    value = builder.CreateInsertElement(value, co_yield get_int(1), (int)3);
-    break;
-  case MTLAttributeFormat::UChar2:
-    value = co_yield load_from_device_buffer(
-      types._char2, base_addr, byte_offset, 0, 1
-    );
-    value = builder.CreateZExt(value, types._int2);
-    value = builder.CreateShuffleVector(value, {0, 1, -1, -1});
-    value = builder.CreateInsertElement(value, co_yield get_int(0), (int)2);
-    value = builder.CreateInsertElement(value, co_yield get_int(1), (int)3);
-    break;
-  case MTLAttributeFormat::UChar4:
-    value = co_yield load_from_device_buffer(
-      types._char4, base_addr, byte_offset, 0, 1
-    );
-    value = builder.CreateZExt(value, types._int4);
-    break;
-  case MTLAttributeFormat::Short:
-    value = co_yield load_from_device_buffer(
-      types._short, base_addr, byte_offset, 0, 2
-    );
-    value = builder.CreateSExt(value, types._int);
-    value = builder.CreateInsertElement(
-      llvm::PoisonValue::get(types._int4), value, (int)0
-    );
-    value = builder.CreateInsertElement(value, co_yield get_int(0), (int)1);
-    value = builder.CreateInsertElement(value, co_yield get_int(0), (int)2);
-    value = builder.CreateInsertElement(value, co_yield get_int(1), (int)3);
-    break;
-  case MTLAttributeFormat::Short2:
-    value = co_yield load_from_device_buffer(
-      types._short2, base_addr, byte_offset, 0, 2
-    );
-    value = builder.CreateSExt(value, types._int2);
-    value = builder.CreateShuffleVector(value, {0, 1, -1, -1});
-    value = builder.CreateInsertElement(value, co_yield get_int(0), (int)2);
-    value = builder.CreateInsertElement(value, co_yield get_int(1), (int)3);
-    break;
-  case MTLAttributeFormat::Short4:
-    value = co_yield load_from_device_buffer(
-      types._short4, base_addr, byte_offset, 0, 2
-    );
-    value = builder.CreateSExt(value, types._int4);
-    break;
-  case MTLAttributeFormat::UShort:
-    value = co_yield load_from_device_buffer(
-      types._short, base_addr, byte_offset, 0, 2
-    );
-    value = builder.CreateZExt(value, types._int);
-    value = builder.CreateInsertElement(
-      llvm::PoisonValue::get(types._int4), value, (int)0
-    );
-    value = builder.CreateInsertElement(value, co_yield get_int(0), (int)1);
-    value = builder.CreateInsertElement(value, co_yield get_int(0), (int)2);
-    value = builder.CreateInsertElement(value, co_yield get_int(1), (int)3);
-    break;
-  case MTLAttributeFormat::UShort2:
-    value = co_yield load_from_device_buffer(
-      types._short2, base_addr, byte_offset, 0, 2
-    );
-    value = builder.CreateZExt(value, types._int2);
-    value = builder.CreateShuffleVector(value, {0, 1, -1, -1});
-    value = builder.CreateInsertElement(value, co_yield get_int(0), (int)2);
-    value = builder.CreateInsertElement(value, co_yield get_int(1), (int)3);
-    break;
-  case MTLAttributeFormat::UShort4:
-    value = co_yield load_from_device_buffer(
-      types._short4, base_addr, byte_offset, 0, 2
-    );
-    value = builder.CreateZExt(value, types._int4);
-    break;
-  case MTLAttributeFormat::Half:
-    value = co_yield load_from_device_buffer(
-      types._half, base_addr, byte_offset, 0, 2
-    );
-    value = ctx.air.CreateConvertToFloat(value);
-    value = builder.CreateInsertElement(
-      llvm::PoisonValue::get(types._float4), value, (int)0
-    );
-    value = builder.CreateInsertElement(value, co_yield get_float(0), (int)1);
-    value = builder.CreateInsertElement(value, co_yield get_float(0), (int)2);
-    value = builder.CreateInsertElement(value, co_yield get_float(1), (int)3);
-    break;
-  case MTLAttributeFormat::Half2:
-    value = co_yield load_from_device_buffer(
-      types._half2, base_addr, byte_offset, 0, 2
-    );
-    value = ctx.air.CreateConvertToFloat(value);
-    value = builder.CreateShuffleVector(value, {0, 1, -1, -1});
-    value = builder.CreateInsertElement(value, co_yield get_float(0), (int)2);
-    value = builder.CreateInsertElement(value, co_yield get_float(1), (int)3);
-    break;
-  case MTLAttributeFormat::Half4:
-    value = co_yield load_from_device_buffer(
-      types._half4, base_addr, byte_offset, 0, 2
-    );
-    value = ctx.air.CreateConvertToFloat(value);
-    break;
-  case MTLAttributeFormat::Float:
-    value = co_yield load_from_device_buffer(
-      types._float, base_addr, byte_offset, 0, 4
-    );
-    value = builder.CreateInsertElement(
-      llvm::PoisonValue::get(types._float4), value, (int)0
-    );
-    value = builder.CreateInsertElement(value, co_yield get_float(0), (int)1);
-    value = builder.CreateInsertElement(value, co_yield get_float(0), (int)2);
-    value = builder.CreateInsertElement(value, co_yield get_float(1), (int)3);
-    break;
-  case MTLAttributeFormat::Float2:
-    value = co_yield load_from_device_buffer(
-      types._float2, base_addr, byte_offset, 0, 4
-    );
-    value = builder.CreateShuffleVector(value, {0, 1, -1, -1});
-    value = builder.CreateInsertElement(value, co_yield get_float(0), (int)2);
-    value = builder.CreateInsertElement(value, co_yield get_float(1), (int)3);
-    break;
-  case MTLAttributeFormat::Float3:
-    value = co_yield load_from_device_buffer(
-      types._float3, base_addr, byte_offset, 0, 4
-    );
-    value = builder.CreateShuffleVector(value, {0, 1, 2, -1});
-    value = builder.CreateInsertElement(value, co_yield get_float(1), (int)3);
-    break;
-  case MTLAttributeFormat::Float4:
-    value = co_yield load_from_device_buffer(
-      types._float4, base_addr, byte_offset, 0, 4
-    );
-    break;
-  case MTLAttributeFormat::UInt:
-  case MTLAttributeFormat::Int:
-    value = co_yield load_from_device_buffer(
-      types._int, base_addr, byte_offset, 0, 4
-    );
-    value = builder.CreateInsertElement(
-      llvm::PoisonValue::get(types._int4), value, (int)0
-    );
-    value = builder.CreateInsertElement(value, co_yield get_int(0), (int)1);
-    value = builder.CreateInsertElement(value, co_yield get_int(0), (int)2);
-    value = builder.CreateInsertElement(value, co_yield get_int(1), (int)3);
-    break;
-  case MTLAttributeFormat::UInt2:
-  case MTLAttributeFormat::Int2:
-    value = co_yield load_from_device_buffer(
-      types._int2, base_addr, byte_offset, 0, 4
-    );
-    value = builder.CreateShuffleVector(value, {0, 1, -1, -1});
-    value = builder.CreateInsertElement(value, co_yield get_int(0), (int)2);
-    value = builder.CreateInsertElement(value, co_yield get_int(1), (int)3);
-    break;
-  case MTLAttributeFormat::UInt3:
-  case MTLAttributeFormat::Int3:
-    value = co_yield load_from_device_buffer(
-      types._int3, base_addr, byte_offset, 0, 4
-    );
-    value = builder.CreateShuffleVector(value, {0, 1, 2, -1});
-    value = builder.CreateInsertElement(value, co_yield get_int(1), (int)3);
-    break;
-  case MTLAttributeFormat::UInt4:
-  case MTLAttributeFormat::Int4:
-    value = co_yield load_from_device_buffer(
-      types._int4, base_addr, byte_offset, 0, 4
-    );
-    break;
-  case MTLAttributeFormat::UCharNormalized:
-  case MTLAttributeFormat::UChar2Normalized:
-  case MTLAttributeFormat::UChar4Normalized:
-  case MTLAttributeFormat::UChar4Normalized_BGRA:
-  case MTLAttributeFormat::CharNormalized:
-  case MTLAttributeFormat::Char2Normalized:
-  case MTLAttributeFormat::Char4Normalized:
-  case MTLAttributeFormat::UShortNormalized:
-  case MTLAttributeFormat::UShort2Normalized:
-  case MTLAttributeFormat::UShort4Normalized:
-  case MTLAttributeFormat::ShortNormalized:
-  case MTLAttributeFormat::Short2Normalized:
-  case MTLAttributeFormat::Short4Normalized:
-  case MTLAttributeFormat::FloatRG11B10:
-  case MTLAttributeFormat::FloatRGB9E5:
-  case MTLAttributeFormat::Int1010102Normalized:
-  case MTLAttributeFormat::UInt1010102Normalized:
-    value = co_yield unpack_fvec4_from_addr(format, base_addr, byte_offset);
-    break;
-  case MTLAttributeFormat::Invalid:
-    break;
-  }
-  assert(value);
-
-  co_return value;
-};
-
-AIRBuilderResult pull_vec4_from_addr(
-  MTLAttributeFormat format, pvalue base_addr, pvalue byte_offset
-) {
-  return make_op([=](struct AIRBuilderContext ctx) -> llvm::Expected<pvalue> {
-    using namespace llvm;
-    auto &context = ctx.llvm;
-    auto &builder = ctx.builder;
-
-    auto current = builder.GetInsertBlock();
-
-    auto is_null_binding = builder.CreateICmpEQ(
-      builder.CreatePtrToInt(base_addr, ctx.types._long), builder.getInt64(0)
-    );
-
-    auto pull_vertex =
-      llvm::BasicBlock::Create(context, "", current->getParent());
-    auto continuation =
-      llvm::BasicBlock::Create(context, "", current->getParent());
-
-    builder.CreateCondBr(is_null_binding, continuation, pull_vertex);
-    builder.SetInsertPoint(pull_vertex);
-
-    auto ret =
-      pull_vec4_from_addr_checked(format, base_addr, byte_offset).build(ctx);
-    if (auto err = ret.takeError()) {
-      return err;
-    }
-    auto value = ret.get();
-
-    builder.CreateBr(continuation);
-    builder.SetInsertPoint(continuation);
-
-    auto phi = builder.CreatePHI(value->getType(), 2);
-    phi->addIncoming(ConstantAggregateZero::get(value->getType()), current);
-    phi->addIncoming(value, pull_vertex);
-
-    return phi;
-  });
-};
-
-}; // namespace dxmt::air
\ No newline at end of file
diff --git a/src/airconv/air_operations.hpp b/src/airconv/air_operations.hpp
deleted file mode 100644
index 2754169c6..000000000
--- a/src/airconv/air_operations.hpp
+++ /dev/null
@@ -1,95 +0,0 @@
-#pragma once
-
-#include "air_signature.hpp"
-#include "air_type.hpp"
-#include "monad.hpp"
-#include "nt/air_builder.hpp"
-#include "llvm/IR/IRBuilder.h"
-#include "llvm/IR/LLVMContext.h"
-#include "llvm/IR/Module.h"
-#include "llvm/IR/Value.h"
-#include "llvm/Support/Error.h"
-
-namespace dxmt::air {
-
-using pvalue = llvm::Value *;
-using epvalue = llvm::Expected<pvalue>;
-
-struct AIRBuilderContext;
-
-using AIRBuilderResult = ReaderIO<AIRBuilderContext, pvalue>;
-
-struct AIRBuilderContext {
-  llvm::LLVMContext &llvm;
-  llvm::Module &module;
-  llvm::IRBuilder<> &builder;
-  AirType &types;
-  llvm::air::AIRBuilder& air;
-};
-
-template <typename S> AIRBuilderResult make_op(S &&fs) {
-  return AIRBuilderResult(std::forward<S>(fs));
-}
-
-enum class MTLAttributeFormat {
-  Invalid = 0,
-  UChar2 = 1,
-  // UChar3 = 2,
-  UChar4 = 3,
-  Char2 = 4,
-  // Char3 = 5,
-  Char4 = 6,
-  UChar2Normalized = 7,
-  // UChar3Normalized = 8,
-  UChar4Normalized = 9,
-  Char2Normalized = 10,
-  // Char3Normalized = 11,
-  Char4Normalized = 12,
-  UShort2 = 13,
-  // UShort3 = 14,
-  UShort4 = 15,
-  Short2 = 16,
-  // Short3 = 17,
-  Short4 = 18,
-  UShort2Normalized = 19,
-  // UShort3Normalized = 20,
-  UShort4Normalized = 21,
-  Short2Normalized = 22,
-  // Short3Normalized = 23,
-  Short4Normalized = 24,
-  Half2 = 25,
-  // Half3 = 26,
-  Half4 = 27,
-  Float = 28,
-  Float2 = 29,
-  Float3 = 30,
-  Float4 = 31,
-  Int = 32,
-  Int2 = 33,
-  Int3 = 34,
-  Int4 = 35,
-  UInt = 36,
-  UInt2 = 37,
-  UInt3 = 38,
-  UInt4 = 39,
-  Int1010102Normalized = 40,
-  UInt1010102Normalized = 41,
-  UChar4Normalized_BGRA = 42,
-  UChar = 45,
-  Char = 46,
-  UCharNormalized = 47,
-  CharNormalized = 48,
-  UShort = 49,
-  Short = 50,
-  UShortNormalized = 51,
-  ShortNormalized = 52,
-  Half = 53,
-  FloatRG11B10 = 54,
-  FloatRGB9E5 = 55,
-};
-
-AIRBuilderResult pull_vec4_from_addr(
-  MTLAttributeFormat format, pvalue base_addr, pvalue byte_offset
-);
-
-}; // namespace dxmt::air
\ No newline at end of file
diff --git a/src/airconv/dxbc_converter.cpp b/src/airconv/dxbc_converter.cpp
index b0704ac75..80b4eb017 100644
--- a/src/airconv/dxbc_converter.cpp
+++ b/src/airconv/dxbc_converter.cpp
@@ -31,11 +31,6 @@ class SM50CompiledBitcodeInternal {
   llvm::SmallVector<char, 0> vec;
 };
 
-class SM50ErrorInternal {
-public:
-  llvm::SmallVector<char, 0> buf;
-};
-
 namespace dxmt::dxbc {
 
 inline dxmt::shader::common::ResourceType
@@ -350,14 +345,8 @@ llvm::Error convert_dxbc_pixel_shader(
   SM50_SHADER_ROOT_SIGNATURE_DATA *rootsig = nullptr;
   args_get_data<SM50_SHADER_ROOT_SIGNATURE, SM50_SHADER_ROOT_SIGNATURE_DATA>(pArgs, &rootsig);
 
-  IREffect prologue([](auto) { return std::monostate(); });
-  IRValue epilogue([](struct context ctx) -> pvalue {
-    auto retTy = ctx.function->getReturnType();
-    if (retTy->isVoidTy()) {
-      return nullptr;
-    }
-    return llvm::ConstantAggregateZero::get(retTy);
-  });
+  Prologues prologue;
+  Epilogues epilogue;
 
   io_binding_map resource_map;
   air::AirType types(context);
@@ -368,6 +357,7 @@ llvm::Error convert_dxbc_pixel_shader(
     sig_ctx.disable_depth_output = pso_disable_depth_output;
     sig_ctx.pull_mode_reg_mask = shader_info->pull_mode_reg_mask;
     sig_ctx.unorm_output_reg_mask = pso_unorm_output_reg_mask;
+    sig_ctx.pso_sample_mask = pso_sample_mask;
     if (pso_pixel_formats)
       memcpy(sig_ctx.pixel_formats, pso_pixel_formats, sizeof(sig_ctx.pixel_formats));
     for (auto &p : pShaderInternal->signature_handlers) {
@@ -377,16 +367,9 @@ llvm::Error convert_dxbc_pixel_shader(
   if (pso_sample_mask != 0xffffffff) {
     auto assigned_index =
       func_signature.DefineOutput(air::OutputCoverageMask{});
-    epilogue >> [=](pvalue value) -> IRValue {
-      return make_irvalue([=](struct context ctx) {
-        auto &builder = ctx.builder;
-        if (ctx.resource.coverage_mask_reg)
-          return value;
-        return builder.CreateInsertValue(
-          value, builder.getInt32(ctx.pso_sample_mask), {assigned_index}
-        );
-      });
-    };
+    epilogue.Add([=](llvm::Value *value, IOHelper &io) -> llvm::Expected<llvm::Value *> {
+      return io.PopulateDefaultCoverageMask(value, assigned_index, pso_sample_mask);
+    });
   }
 
   auto binding_map = rootsig ? setup_binding_rootsig(
@@ -447,7 +430,9 @@ llvm::Error convert_dxbc_pixel_shader(
       .metal_version = metal_version,
   };
 
-  if (auto err = prologue.build(ctx).takeError()) {
+  IOHelper io(air, function, types, resource_map);
+
+  if (auto err = prologue.Run(io)) {
     return err;
   }
   auto real_entry = convert_basicblocks(pShaderInternal->entry(), ctx, epilogue_bb);
@@ -457,7 +442,7 @@ llvm::Error convert_dxbc_pixel_shader(
   builder.CreateBr(real_entry.get());
 
   builder.SetInsertPoint(epilogue_bb);
-  auto epilogue_result = epilogue.build(ctx);
+  auto epilogue_result = epilogue.Run(function->getReturnType(), io);
   if (auto err = epilogue_result.takeError()) {
     return err;
   }
@@ -494,14 +479,8 @@ llvm::Error convert_dxbc_compute_shader(
   SM50_SHADER_ROOT_SIGNATURE_DATA *rootsig = nullptr;
   args_get_data<SM50_SHADER_ROOT_SIGNATURE, SM50_SHADER_ROOT_SIGNATURE_DATA>(pArgs, &rootsig);
 
-  IREffect prologue([](auto) { return std::monostate(); });
-  IRValue epilogue([](struct context ctx) -> pvalue {
-    auto retTy = ctx.function->getReturnType();
-    if (retTy->isVoidTy()) {
-      return nullptr;
-    }
-    return llvm::ConstantAggregateZero::get(retTy);
-  });
+  Prologues prologue;
+  Epilogues epilogue;
 
   io_binding_map resource_map;
   air::AirType types(context);
@@ -555,7 +534,9 @@ llvm::Error convert_dxbc_compute_shader(
       .metal_version = metal_version,
   };
 
-  if (auto err = prologue.build(ctx).takeError()) {
+  IOHelper io(air, function, types, resource_map);
+
+  if (auto err = prologue.Run(io)) {
     return err;
   }
   auto real_entry = convert_basicblocks(pShaderInternal->entry(), ctx, epilogue_bb);
@@ -566,7 +547,8 @@ llvm::Error convert_dxbc_compute_shader(
 
   builder.SetInsertPoint(epilogue_bb);
 
-  if (auto err = epilogue.build(ctx).takeError()) {
+  auto epilogue_result = epilogue.Run(function->getReturnType(), io);
+  if (auto err = epilogue_result.takeError()) {
     return err;
   }
   builder.CreateRetVoid();
@@ -609,14 +591,8 @@ llvm::Error convert_dxbc_vertex_shader(
   SM50_SHADER_ROOT_SIGNATURE_DATA *rootsig = nullptr;
   args_get_data<SM50_SHADER_ROOT_SIGNATURE, SM50_SHADER_ROOT_SIGNATURE_DATA>(pArgs, &rootsig);
 
-  IREffect prologue([](auto) { return std::monostate(); });
-  IRValue epilogue([](struct context ctx) -> pvalue {
-    auto retTy = ctx.function->getReturnType();
-    if (retTy->isVoidTy()) {
-      return nullptr;
-    }
-    return llvm::ConstantAggregateZero::get(retTy);
-  });
+  Prologues prologue;
+  Epilogues epilogue;
 
   io_binding_map resource_map;
   air::AirType types(context);
@@ -650,44 +626,8 @@ llvm::Error convert_dxbc_vertex_shader(
       .arg_name = "stream_outputs",
       .raster_order_group = {}
     });
-    epilogue << make_irvalue([=](struct context ctx) {
-      auto &builder = ctx.builder;
-      auto base_vertex = ctx.function->getArg(bv);
-      auto vertex_id = ctx.function->getArg(vid);
-      auto so_entries_type = ctx.types._dxmt_stream_output_buffer_entry;
-      auto so_entries = builder.CreateBitCast(
-          ctx.function->getArg(so_table), so_entries_type->getPointerTo((uint32_t)air::AddressSpace::constant)
-      );
-      auto slot0_entry =
-          builder.CreateLoad(so_entries_type, builder.CreateConstInBoundsGEP1_32(so_entries_type, so_entries, 0));
-      auto slot0 = builder.CreateExtractValue(slot0_entry, {0});
-      auto adjusted_vertex_id = builder.CreateSub(vertex_id, base_vertex);
-      auto output_regs = builder.CreateBitOrPointerCast(
-        ctx.resource.output.ptr_int4, llvm::PointerType::get(ctx.types._int, 0)
-      );
-      for (unsigned i = 0; i < vertex_so->num_elements; i++) {
-        auto &element = vertex_so->elements[i];
-        if (element.reg_id == 0xffffffff)
-          continue;
-        auto ptr = ctx.builder.CreateConstGEP1_32(
-          ctx.types._int, output_regs,
-          (unsigned)(element.reg_id * 4 + element.component)
-        );
-        auto target_offset = ctx.builder.CreateAdd(
-          ctx.builder.CreateMul(
-            adjusted_vertex_id,
-            ctx.builder.getInt32(vertex_so->strides[element.output_slot /* expected to be 0 */])
-          ),
-          ctx.builder.getInt32(element.offset)
-        );
-        auto target_ptr = ctx.builder.CreateGEP(
-          ctx.types._int, slot0, {ctx.builder.CreateLShr(target_offset, 2)}
-        );
-        ctx.builder.CreateStore(
-          ctx.builder.CreateLoad(ctx.types._int, ptr), target_ptr, true
-        );
-      }
-      return nullptr;
+    epilogue.Add([=](llvm::Value *value, IOHelper &io) -> llvm::Expected<llvm::Value *> {
+      return io.PopulateOutputStreamOutput(value, vertex_so, bv, vid, so_table);
     });
   }
 
@@ -785,7 +725,9 @@ llvm::Error convert_dxbc_vertex_shader(
       .metal_version = metal_version,
   };
 
-  if (auto err = prologue.build(ctx).takeError()) {
+  IOHelper io(air, function, types, resource_map);
+
+  if (auto err = prologue.Run(io)) {
     return err;
   }
   auto real_entry = convert_basicblocks(pShaderInternal->entry(), ctx, epilogue_bb);
@@ -795,7 +737,7 @@ llvm::Error convert_dxbc_vertex_shader(
   builder.CreateBr(real_entry.get());
 
   builder.SetInsertPoint(epilogue_bb);
-  auto epilogue_result = epilogue.build(ctx);
+  auto epilogue_result = epilogue.Run(function->getReturnType(), io);
   if (auto err = epilogue_result.takeError()) {
     return err;
   }
@@ -992,20 +934,18 @@ AIRCONV_API int SM50Initialize(
   if (ppError) {
     *ppError = nullptr;
   }
-  auto errorObj = new SM50ErrorInternal();
+  auto errorObj = std::make_unique<SM50ErrorInternal>();
   llvm::raw_svector_ostream errorOut(errorObj->buf);
 
   if (ppShader == nullptr) {
     errorOut << "ppShader can not be null\0";
-    *ppError = (sm50_error_t)errorObj;
-    return 1;
+    SM50_RETURN_ERROR(errorObj, ppError, 1);
   }
 
   CDXBCParser DXBCParser;
   if (DXBCParser.ReadDXBC(pBytecode, BytecodeSize) != S_OK) {
     errorOut << "Invalid DXBC bytecode\0";
-    *ppError = (sm50_error_t)errorObj;
-    return 1;
+    SM50_RETURN_ERROR(errorObj, ppError, 1);
   }
 
   UINT codeBlobIdx = DXBCParser.FindNextMatchingBlob(DXBC_GenericShaderEx);
@@ -1014,8 +954,7 @@ AIRCONV_API int SM50Initialize(
   }
   if (codeBlobIdx == DXBC_BLOB_NOT_FOUND) {
     errorOut << "Invalid DXBC bytecode: shader blob not found\0";
-    *ppError = (sm50_error_t)errorObj;
-    return 1;
+    SM50_RETURN_ERROR(errorObj, ppError, 1);
   }
   const void *codeBlob = DXBCParser.GetBlob(codeBlobIdx);
 
@@ -1025,17 +964,15 @@ AIRCONV_API int SM50Initialize(
   CSignatureParser inputParser;
   if (DXBCGetInputSignature(pBytecode, &inputParser) != S_OK) {
     errorOut << "Invalid DXBC bytecode: input signature not found\0";
-    *ppError = (sm50_error_t)errorObj;
-    return 1;
+    SM50_RETURN_ERROR(errorObj, ppError, 1);
   }
   CSignatureParser5 outputParser;
   if (DXBCGetOutputSignature(pBytecode, &outputParser) != S_OK) {
     errorOut << "Invalid DXBC bytecode: output signature not found\0";
-    *ppError = (sm50_error_t)errorObj;
-    return 1;
+    SM50_RETURN_ERROR(errorObj, ppError, 1);
   }
 
-  auto sm50_shader = new SM50ShaderInternal();
+  auto sm50_shader = std::make_unique<SM50ShaderInternal>();
   sm50_shader->shader_type = CodeParser.ShaderType();
   auto shader_info = &(sm50_shader->shader_info);
 
@@ -1047,7 +984,7 @@ AIRCONV_API int SM50Initialize(
     }
   }
 
-  sm50_shader->bbs = read_control_flow(CodeParser, sm50_shader, inputParser, outputParser);
+  sm50_shader->bbs = read_control_flow(CodeParser, sm50_shader.get(), inputParser, outputParser);
 
   auto &binding_table = shader_info->binding_table;
   auto &binding_table_cbuffer = shader_info->binding_table_cbuffer;
@@ -1252,20 +1189,18 @@ AIRCONV_API int SM50Initialize(
       auto threads_per_patch = next_pow2(sm50_shader->hull_maximum_threads_per_patch);
       if (threads_per_patch > 32) {
         errorOut << "Threadgroup size of tessellation pipeline is too large.";
-        *ppError = (sm50_error_t)errorObj;
-        return 1;
+        SM50_RETURN_ERROR(errorObj, ppError, 1);
       }
       auto patch_per_group = 32 / threads_per_patch;
       float max_tesselation_factor = sm50_shader->max_tesselation_factor;
 
-      while (estimate_payload_size(sm50_shader, max_tesselation_factor, patch_per_group) > 16384) {
+      while (estimate_payload_size(sm50_shader.get(), max_tesselation_factor, patch_per_group) > 16384) {
         if (patch_per_group == 1) {
           if (max_tesselation_factor > 1.0f) {
             max_tesselation_factor = std::max(max_tesselation_factor - 2.0f, 1.0f);
           } else {
             errorOut << "Payload size of tessellation pipeline is too large.";
-            *ppError = (sm50_error_t)errorObj;
-            return 1;
+            SM50_RETURN_ERROR(errorObj, ppError, 1);
           }
         } else {
           patch_per_group = patch_per_group >> 1;
@@ -1286,7 +1221,7 @@ AIRCONV_API int SM50Initialize(
       uint32_t max_potential_tess_factor = 1;
 
       for (int tess_factor = 1; tess_factor <= 64; tess_factor++) {
-        auto x = estimate_mesh_size(sm50_shader, tess_factor);
+        auto x = estimate_mesh_size(sm50_shader.get(), tess_factor);
         if (x > 32768)
           break;
         max_potential_tess_factor = tess_factor;
@@ -1311,7 +1246,7 @@ AIRCONV_API int SM50Initialize(
     pRefl->ArgumentTableQwords = binding_table.Size();
   }
 
-  *ppShader = sm50_shader;
+  *ppShader = (sm50_shader_t)sm50_shader.release();
   return 0;
 };
 
@@ -1344,22 +1279,21 @@ AIRCONV_API int SM50Compile(
 ) {
   using namespace llvm;
   using namespace dxmt;
+  using namespace dxmt::dxbc;
 
   if (ppError) {
     *ppError = nullptr;
   }
-  auto errorObj = new SM50ErrorInternal();
+  auto errorObj = std::make_unique<SM50ErrorInternal>();
   llvm::raw_svector_ostream errorOut(errorObj->buf);
   if (ppBitcode == nullptr) {
     errorOut << "ppBitcode can not be null\0";
-    *ppError = (sm50_error_t)errorObj;
-    return 1;
+    SM50_RETURN_ERROR(errorObj, ppError, 1);
   }
 
-  // pArgs is ignored for now
   LLVMContext context;
 
-  context.setOpaquePointers(false); // I suspect Metal uses LLVM 14...
+  context.setOpaquePointers(false);
 
   auto &shader_info = ((dxmt::dxbc::SM50ShaderInternal *)pShader)->shader_info;
 
@@ -1372,8 +1306,7 @@ AIRCONV_API int SM50Compile(
     llvm::handleAllErrors(std::move(err), [&](const UnsupportedFeature &u) {
       errorOut << u.msg;
     });
-    *ppError = (sm50_error_t)errorObj;
-    return 1;
+    SM50_RETURN_ERROR(errorObj, ppError, 1);
   }
 
   if (shader_info.use_msad)
@@ -1405,22 +1338,21 @@ AIRCONV_API int SM50CompileTessellationPipelineHull(
 ) {
   using namespace llvm;
   using namespace dxmt;
+  using namespace dxmt::dxbc;
 
   if (ppError) {
     *ppError = nullptr;
   }
-  auto errorObj = new SM50ErrorInternal();
+  auto errorObj = std::make_unique<SM50ErrorInternal>();
   llvm::raw_svector_ostream errorOut(errorObj->buf);
   if (ppBitcode == nullptr) {
     errorOut << "ppBitcode can not be null\0";
-    *ppError = (sm50_error_t)errorObj;
-    return 1;
+    SM50_RETURN_ERROR(errorObj, ppError, 1);
   }
 
-  // pArgs is ignored for now
   LLVMContext context;
 
-  context.setOpaquePointers(false); // I suspect Metal uses LLVM 14...
+  context.setOpaquePointers(false);
 
   auto &shader_info =
     ((dxmt::dxbc::SM50ShaderInternal *)pHullShader)->shader_info;
@@ -1435,8 +1367,7 @@ AIRCONV_API int SM50CompileTessellationPipelineHull(
     llvm::handleAllErrors(std::move(err), [&](const UnsupportedFeature &u) {
       errorOut << u.msg;
     });
-    *ppError = (sm50_error_t)errorObj;
-    return 1;
+    SM50_RETURN_ERROR(errorObj, ppError, 1);
   }
 
   if (shader_info.use_msad)
@@ -1469,22 +1400,21 @@ AIRCONV_API int SM50CompileTessellationPipelineDomain(
 ) {
   using namespace llvm;
   using namespace dxmt;
+  using namespace dxmt::dxbc;
 
   if (ppError) {
     *ppError = nullptr;
   }
-  auto errorObj = new SM50ErrorInternal();
+  auto errorObj = std::make_unique<SM50ErrorInternal>();
   llvm::raw_svector_ostream errorOut(errorObj->buf);
   if (ppBitcode == nullptr) {
     errorOut << "ppBitcode can not be null\0";
-    *ppError = (sm50_error_t)errorObj;
-    return 1;
+    SM50_RETURN_ERROR(errorObj, ppError, 1);
   }
 
-  // pArgs is ignored for now
   LLVMContext context;
 
-  context.setOpaquePointers(false); // I suspect Metal uses LLVM 14...
+  context.setOpaquePointers(false);
 
   auto &shader_info =
     ((dxmt::dxbc::SM50ShaderInternal *)pDomainShader)->shader_info;
@@ -1500,8 +1430,7 @@ AIRCONV_API int SM50CompileTessellationPipelineDomain(
     llvm::handleAllErrors(std::move(err), [&](const UnsupportedFeature &u) {
       errorOut << u.msg;
     });
-    *ppError = (sm50_error_t)errorObj;
-    return 1;
+    SM50_RETURN_ERROR(errorObj, ppError, 1);
   }
 
   if (shader_info.use_msad)
@@ -1534,22 +1463,21 @@ AIRCONV_API int SM50CompileGeometryPipelineVertex(
 ) {
   using namespace llvm;
   using namespace dxmt;
+  using namespace dxmt::dxbc;
 
   if (ppError) {
     *ppError = nullptr;
   }
-  auto errorObj = new SM50ErrorInternal();
+  auto errorObj = std::make_unique<SM50ErrorInternal>();
   llvm::raw_svector_ostream errorOut(errorObj->buf);
   if (ppBitcode == nullptr) {
     errorOut << "ppBitcode can not be null\0";
-    *ppError = (sm50_error_t)errorObj;
-    return 1;
+    SM50_RETURN_ERROR(errorObj, ppError, 1);
   }
 
-  // pArgs is ignored for now
   LLVMContext context;
 
-  context.setOpaquePointers(false); // I suspect Metal uses LLVM 14...
+  context.setOpaquePointers(false);
 
   auto &shader_info = ((dxmt::dxbc::SM50ShaderInternal *)pGeometryShader)->shader_info;
 
@@ -1564,8 +1492,7 @@ AIRCONV_API int SM50CompileGeometryPipelineVertex(
     llvm::handleAllErrors(std::move(err), [&](const UnsupportedFeature &u) {
       errorOut << u.msg;
     });
-    *ppError = (sm50_error_t)errorObj;
-    return 1;
+    SM50_RETURN_ERROR(errorObj, ppError, 1);
   }
 
   if (shader_info.use_msad)
@@ -1597,22 +1524,21 @@ AIRCONV_API int SM50CompileGeometryPipelineGeometry(
 ) {
   using namespace llvm;
   using namespace dxmt;
+  using namespace dxmt::dxbc;
 
   if (ppError) {
     *ppError = nullptr;
   }
-  auto errorObj = new SM50ErrorInternal();
+  auto errorObj = std::make_unique<SM50ErrorInternal>();
   llvm::raw_svector_ostream errorOut(errorObj->buf);
   if (ppBitcode == nullptr) {
     errorOut << "ppBitcode can not be null\0";
-    *ppError = (sm50_error_t)errorObj;
-    return 1;
+    SM50_RETURN_ERROR(errorObj, ppError, 1);
   }
 
-  // pArgs is ignored for now
   LLVMContext context;
 
-  context.setOpaquePointers(false); // I suspect Metal uses LLVM 14...
+  context.setOpaquePointers(false);
 
   auto &shader_info =
     ((dxmt::dxbc::SM50ShaderInternal *)pGeometryShader)->shader_info;
@@ -1628,8 +1554,7 @@ AIRCONV_API int SM50CompileGeometryPipelineGeometry(
     llvm::handleAllErrors(std::move(err), [&](const UnsupportedFeature &u) {
       errorOut << u.msg;
     });
-    *ppError = (sm50_error_t)errorObj;
-    return 1;
+    SM50_RETURN_ERROR(errorObj, ppError, 1);
   }
 
   if (shader_info.use_msad)
@@ -1668,7 +1593,7 @@ AIRCONV_API void SM50DestroyBitcode(sm50_bitcode_t pBitcode) {
 }
 
 AIRCONV_API size_t SM50GetErrorMessage(sm50_error_t pError, char *pBuffer, size_t BufferSize) {
-  auto pInternal = (SM50ErrorInternal *)pError;
+  auto pInternal = (dxmt::dxbc::SM50ErrorInternal *)pError;
   auto str_len = std::min(pInternal->buf.size(), BufferSize - 1);
   memcpy(pBuffer, pInternal->buf.data(), str_len);
   pBuffer[str_len] = '\0';
@@ -1678,6 +1603,6 @@ AIRCONV_API size_t SM50GetErrorMessage(sm50_error_t pError, char *pBuffer, size_
 AIRCONV_API void SM50FreeError(sm50_error_t pError) {
   if (pError == nullptr)
     return;
-  auto pInternal = (SM50ErrorInternal *)pError;
+  auto pInternal = (dxmt::dxbc::SM50ErrorInternal *)pError;
   delete pInternal;
 }
diff --git a/src/airconv/dxbc_converter.hpp b/src/airconv/dxbc_converter.hpp
index a64f6e0f9..d4ed25392 100644
--- a/src/airconv/dxbc_converter.hpp
+++ b/src/airconv/dxbc_converter.hpp
@@ -1,6 +1,7 @@
 #pragma once
 
 #include <cstdint>
+#include <functional>
 #include <map>
 #include <optional>
 #include <string>
@@ -9,12 +10,13 @@
 #include <vector>
 
 #include "DXBCParser/DXBCUtils.h"
-#include "air_operations.hpp"
 #include "air_signature.hpp"
+#include "air_type.hpp"
 #include "dxbc_constants.hpp"
 #include "dxbc_instructions.hpp"
 #include "nt/air_builder.hpp"
 #include "nt/dxbc_binding_map.hpp"
+#include "nt/dxbc_io_helper.hpp"
 #include "shader_common.hpp"
 
 #include "airconv_public.h"
@@ -111,15 +113,12 @@ Instruction readInstruction(
   ShaderInfo &shader_info, uint32_t phase
 );
 
-using pvalue = dxmt::air::pvalue;
+using pvalue = llvm::Value *;
 using epvalue = llvm::Expected<pvalue>;
 using dxbc::Swizzle;
 using dxbc::swizzle_identity;
 
 struct context;
-using IRValue = ReaderIO<context, pvalue>;
-using IREffect = ReaderIO<context, std::monostate>;
-using IndexedIRValue = std::function<IRValue(pvalue)>;
 
 struct register_file {
   llvm::Value *ptr_int4 = nullptr;
@@ -137,28 +136,8 @@ struct phase_temp {
   std::unordered_map<uint32_t, indexable_register_file> indexable_temp_map{};
 };
 
-struct sampler_descriptor {
-  IndexedIRValue handle;
-  IndexedIRValue handle_cube;
-  IndexedIRValue bias;
-};
-
-struct texture_descriptor {
-  air::MSLTexture texture_info;
-  IndexedIRValue resource_id;
-  IndexedIRValue metadata;
-  bool global_coherent;
-};
-
-struct buffer_descriptor {
-  uint32_t structure_stride;
-  IndexedIRValue resource_id;
-  IndexedIRValue metadata;
-  bool global_coherent;
-};
-
 struct interpolant_descriptor {
-  IndexedIRValue interpolant;
+  uint32_t arg_index;
   bool perspective;
 };
 
@@ -216,8 +195,8 @@ struct io_binding_map {
 
   // geometry shader ops
   llvm::Value *mesh = nullptr;
-  std::function<IREffect()> call_emit;
-  std::function<IREffect()> call_cut;
+  std::function<void()> call_emit;
+  std::function<void()> call_cut;
 };
 
 struct context {
@@ -234,60 +213,6 @@ struct context {
   SM50_SHADER_METAL_VERSION metal_version;
 };
 
-template <typename S> IRValue make_irvalue(S &&fs) {
-  return IRValue(std::forward<S>(fs));
-}
-
-template <typename S> IRValue make_irvalue_bind(S &&fs) {
-  return IRValue([fs = std::forward<S>(fs)](auto ctx) {
-    return fs(ctx).build(ctx);
-  });
-}
-
-template <typename S> IREffect make_effect(S &&fs) {
-  return IREffect(std::forward<S>(fs));
-}
-
-template <typename S> IREffect make_effect_bind(S &&fs) {
-  return IREffect([fs = std::forward<S>(fs)](auto ctx) mutable {
-    return fs(ctx).build(ctx);
-  });
-}
-
-IREffect store_at_vec4_array_masked(
-  llvm::Value *array, pvalue index, pvalue maybe_vec4, uint32_t mask
-);
-
-IREffect init_input_reg(
-  uint32_t with_fnarg_at, uint32_t to_reg, uint32_t mask,
-  bool fix_w_component = false
-);
-
-IREffect init_input_reg_with_interpolation(
-  uint32_t with_fnarg_at, uint32_t to_reg, uint32_t mask,
-  air::Interpolation interpolation, uint32_t sampleidx_at
-);
-
-std::function<IRValue(pvalue)>
-pop_output_reg(uint32_t from_reg, uint32_t mask, uint32_t to_element);
-
-std::function<IRValue(pvalue)>
-pop_output_reg_fix_unorm(uint32_t from_reg, uint32_t mask, uint32_t to_element);
-
-std::function<IRValue(pvalue)>
-pop_output_reg_sanitize_pos(uint32_t from_reg, uint32_t mask, uint32_t to_element);
-
-IREffect pull_vertex_input(
-  air::FunctionSignatureBuilder &func_signature, uint32_t to_reg, uint32_t mask,
-  SM50_IA_INPUT_ELEMENT element_info, uint32_t slot_mask
-);
-
-IREffect pop_mesh_output_render_target_array_index(uint32_t from_reg, uint32_t mask, pvalue primitive_id);
-IREffect pop_mesh_output_viewport_array_index(uint32_t from_reg, uint32_t mask, pvalue primitive_id);
-IREffect pop_mesh_output_position(uint32_t from_reg, uint32_t mask, pvalue vertex_id);
-IREffect
-pop_mesh_output_vertex_data(uint32_t from_reg, uint32_t mask, uint32_t idx, pvalue vertex_id, air::MSLScalerOrVectorType desired_type);
-
 llvm::Expected<llvm::BasicBlock *> convert_basicblocks(
   BasicBlock *entry, context &ctx, llvm::BasicBlock *return_bb
 );
@@ -321,8 +246,8 @@ struct PatchConstantScalarInfo {
 };
 
 struct SignatureContext {
-  IREffect &prologue;
-  IRValue &epilogue;
+  Prologues &prologue;
+  Epilogues &epilogue;
   air::FunctionSignatureBuilder &func_signature;
   io_binding_map &resource;
   SM50_SHADER_IA_INPUT_LAYOUT_DATA *ia_layout;
@@ -331,11 +256,12 @@ struct SignatureContext {
   bool skip_vertex_output;
   uint32_t pull_mode_reg_mask;
   uint32_t unorm_output_reg_mask;
+  uint32_t pso_sample_mask;
   // it is considered optional (as a hint)
   air::MTLPixelFormat pixel_formats[8];
 
   SignatureContext(
-      IREffect &prologue, IRValue &epilogue, air::FunctionSignatureBuilder &func_signature, io_binding_map &resource
+      Prologues &prologue, Epilogues &epilogue, air::FunctionSignatureBuilder &func_signature, io_binding_map &resource
   ) :
       prologue(prologue),
       epilogue(epilogue),
@@ -346,7 +272,8 @@ struct SignatureContext {
       disable_depth_output(false),
       skip_vertex_output(false),
       pull_mode_reg_mask(0),
-      unorm_output_reg_mask(0) {
+      unorm_output_reg_mask(0),
+      pso_sample_mask(0xffffffff) {
     memset(pixel_formats, 0, sizeof(pixel_formats));
   };
 };
@@ -442,7 +369,7 @@ class SM50ShaderInternal {
   std::vector<ScalarInfo> clip_distance_scalars;
   std::vector<ScalarInfo> cull_distance_scalars;
   microsoft::D3D10_SB_PRIMITIVE gs_input_primitive = {};
-  std::vector<std::function<IREffect(MeshOutputContext &)>> mesh_output_handlers;
+  std::vector<std::function<void(MeshOutputContext &, IOHelper &)>> mesh_output_handlers;
   uint32_t num_mesh_vertex_data = 0;
   microsoft::D3D10_SB_PRIMITIVE_TOPOLOGY gs_output_topology = {};
   uint32_t gs_max_vertex_output = 0;
@@ -454,6 +381,18 @@ class SM50ShaderInternal {
   }
 };
 
+class SM50ErrorInternal {
+public:
+  llvm::SmallVector<char, 0> buf;
+};
+
+#define SM50_RETURN_ERROR(errorObj, ppError, errorCode) \
+  do { \
+    if (ppError) \
+      *(ppError) = (sm50_error_t)(errorObj).release(); \
+    return errorCode; \
+  } while (0)
+
 void handle_signature(
   microsoft::CSignatureParser &inputParser,
   microsoft::CSignatureParser5 &outputParser,
diff --git a/src/airconv/dxbc_converter_basicblock.cpp b/src/airconv/dxbc_converter_basicblock.cpp
index a289a582f..b237e8c16 100644
--- a/src/airconv/dxbc_converter_basicblock.cpp
+++ b/src/airconv/dxbc_converter_basicblock.cpp
@@ -1,11 +1,10 @@
-#include "air_operations.hpp"
 #include "air_signature.hpp"
 #include "air_type.hpp"
 #include "airconv_error.hpp"
 #include "airconv_public.h"
 #include "dxbc_converter.hpp"
-#include "ftl.hpp"
 #include "nt/air_builder.hpp"
+#include "nt/air_vertex_puller.hpp"
 #include "nt/dxbc_converter_base.hpp"
 #include "llvm/ADT/APFloat.h"
 #include "llvm/ADT/APInt.h"
@@ -30,427 +29,6 @@ char dxmt::UnsupportedFeature::ID;
 
 namespace dxmt::dxbc {
 
-template <typename T = std::monostate>
-ReaderIO<context, T> throwUnsupported(llvm::StringRef ref) {
-  return ReaderIO<context, T>(
-    [msg = std::string(ref)](struct context ctx) mutable -> llvm::Expected<T> {
-      // CAUTION: create llvm::Error as lazy as possible
-      // DON'T create it and store somewhere like plain data
-      // ^ which make fking perfect sense but NO DONT DO THAT!
-      // BECAUSE IT IS STATEFUL! IT PRESEVED A 'CHECKED' STATE
-      // to enforce programmer handle any created error
-      // ^ imaging you move-capture the error in a lambda
-      //   but it's never called and eventually deconstructed
-      //   then BOOM
-      // what the fuck why should I use it
-      // `Monadic error handling w/ a gun shoot yourself in the foot`
-      // ^ never claimed to be a monad actually 👆🤓
-      // DO I REALLY DESERVE THIS?
-      return llvm::Expected<T>(llvm::make_error<UnsupportedFeature>(msg));
-    }
-  );
-};
-
-ReaderIO<context, context> get_context() {
-  return ReaderIO<context, context>([](struct context ctx) { return ctx; });
-};
-
-auto get_float4_splat(float value) -> IRValue {
-  return make_irvalue([=](context ctx) {
-    return llvm::ConstantVector::get(
-      {llvm::ConstantFP::get(ctx.llvm, llvm::APFloat{value}),
-       llvm::ConstantFP::get(ctx.llvm, llvm::APFloat{value}),
-       llvm::ConstantFP::get(ctx.llvm, llvm::APFloat{value}),
-       llvm::ConstantFP::get(ctx.llvm, llvm::APFloat{value})}
-    );
-  });
-};
-
-auto extend_to_vec4(pvalue value) {
-  return make_irvalue([=](context ctx) {
-    auto ty = value->getType();
-    if (ty->isVectorTy()) {
-      auto vecTy = llvm::cast<llvm::FixedVectorType>(ty);
-      if (vecTy->getNumElements() == 4)
-        return ctx.builder.CreateShuffleVector(value, {0, 1, 2, 3});
-      if (vecTy->getNumElements() == 3)
-        return ctx.builder.CreateShuffleVector(value, {0, 1, 2, 2});
-      if (vecTy->getNumElements() == 2)
-        return ctx.builder.CreateShuffleVector(value, {0, 1, 1, 1});
-      if (vecTy->getNumElements() == 1)
-        return ctx.builder.CreateShuffleVector(value, {0, 0, 0, 0});
-      assert(0 && "?");
-    } else {
-      return ctx.builder.CreateVectorSplat(4, value);
-    }
-  });
-}
-
-auto to_desired_type_from_int_vec4(pvalue vec4, llvm::Type *desired, uint32_t mask) {
-  return make_irvalue([=](context ctx) {
-    assert(vec4->getType() == ctx.types._int4);
-    std::function<pvalue(pvalue, llvm::Type *)> convert =
-      [&ctx, &convert, mask](pvalue vec4, llvm::Type *desired) {
-        auto masked = [mask](int i) {
-          return (mask & (1 << i)) ? i : llvm::UndefMaskElem;
-        };
-        if (desired == ctx.types._int4)
-          return ctx.builder.CreateShuffleVector(
-            vec4, {masked(0), masked(1), masked(2), masked(3)}
-          );
-        if (desired == ctx.types._float4)
-          return ctx.builder.CreateShuffleVector(
-            ctx.builder.CreateBitCast(vec4, ctx.types._float4),
-            {masked(0), masked(1), masked(2), masked(3)}
-          );
-        // FIXME: 3d/2d vector implementations are unused and probably wrong!
-        if (desired == ctx.types._int3)
-          return ctx.builder.CreateShuffleVector(
-            vec4, {masked(0), masked(1), masked(2)}
-          );
-        if (desired == ctx.types._float3)
-          return ctx.builder.CreateShuffleVector(
-            convert(vec4, ctx.types._float4), {masked(0), masked(1), masked(2)}
-          );
-        if (desired == ctx.types._int2)
-          return ctx.builder.CreateShuffleVector(vec4, {masked(0), masked(1)});
-        if (desired == ctx.types._float2)
-          return ctx.builder.CreateShuffleVector(
-            convert(vec4, ctx.types._float4), {masked(0), masked(1)}
-          );
-        if (desired == ctx.types._int)
-          return ctx.builder.CreateExtractElement(vec4, (uint64_t)__builtin_ctz(mask));
-        if (desired == ctx.types._float)
-          return ctx.builder.CreateExtractElement(
-            ctx.builder.CreateBitCast(vec4, ctx.types._float4), (uint64_t)__builtin_ctz(mask)
-          );
-        assert(0 && "unhandled vec4");
-      };
-    return convert(vec4, desired);
-  });
-};
-
-auto load_from_array_at(llvm::Value *array, pvalue index) -> IRValue {
-  return make_irvalue([=](context ctx) {
-    auto array_ty = llvm::cast<llvm::ArrayType>( // force line break
-      llvm::cast<llvm::PointerType>(array->getType())
-        ->getNonOpaquePointerElementType()
-    );
-    auto ptr = ctx.builder.CreateGEP(
-      array_ty, array, {ctx.builder.getInt32(0), index}, "",
-      llvm::isa<llvm::ConstantInt>(index)
-    );
-    return ctx.builder.CreateLoad(array_ty->getElementType(), ptr);
-  });
-};
-
-IRValue
-load_from_vec4_array_masked(llvm::Value *array, pvalue index, uint32_t mask) {
-  if (mask == 0b1111) {
-    return load_from_array_at(array, index);
-  }
-  return make_irvalue([=](context ctx) {
-    auto array_type = llvm::cast<llvm::PointerType>(array->getType())->getNonOpaquePointerElementType();
-    auto vec4_type = llvm::cast<llvm::ArrayType>(array_type)->getArrayElementType();
-    auto ele_type = llvm::cast<llvm::VectorType>(vec4_type)->getElementType();
-    pvalue value = llvm::ConstantAggregateZero::get(vec4_type);
-    for (unsigned i = 0; i < 4; i++) {
-      if ((mask & (1 << i)) == 0)
-        continue;
-      auto component_ptr =
-          ctx.builder.CreateGEP(array_type, array, {ctx.builder.getInt32(0), index, ctx.builder.getInt32(i)});
-      value = ctx.builder.CreateInsertElement(value, ctx.builder.CreateLoad(ele_type, component_ptr), uint64_t(i));
-    }
-    return value;
-  });
-};
-
-auto store_to_array_at(
-  llvm::Value *array, pvalue index, pvalue vec4_type_matched
-) -> IREffect {
-  return make_effect([=](context ctx) {
-    auto ptr = ctx.builder.CreateInBoundsGEP(
-      llvm::cast<llvm::PointerType>(array->getType())
-        ->getNonOpaquePointerElementType(),
-      array, {ctx.builder.getInt32(0), index}
-    );
-    ctx.builder.CreateStore(vec4_type_matched, ptr);
-    return std::monostate();
-  });
-};
-
-IREffect store_at_vec4_array_masked(
-  llvm::Value *array, pvalue index, pvalue maybe_vec4, uint32_t mask
-) {
-  return extend_to_vec4(maybe_vec4) >>= [=](pvalue vec4) {
-    if (mask == 0b1111) {
-      return store_to_array_at(array, index, vec4);
-    }
-    return make_effect([=](context ctx) {
-      for (unsigned i = 0; i < 4; i++) {
-        if ((mask & (1 << i)) == 0)
-          continue;
-        auto component_ptr =  ctx.builder.CreateGEP(
-          llvm::cast<llvm::PointerType>(array->getType())
-            ->getNonOpaquePointerElementType(),
-          array, {ctx.builder.getInt32(0), index, ctx.builder.getInt32(i)}
-        );
-        ctx.builder.CreateStore(
-          ctx.builder.CreateExtractElement(vec4, i), component_ptr
-        );
-      }
-      return std::monostate();
-    });
-  };
-};
-
-IREffect init_input_reg(
-  uint32_t with_fnarg_at, uint32_t to_reg, uint32_t mask,
-  bool fix_w_component
-) {
-  // no it doesn't work like this
-  // regular input can be masked like .zw
-  // assert((mask & 1) && "todo: handle input register sharing correctly");
-  // FIXME: it's buggy as hell
-  return make_effect_bind([=](context ctx) {
-    pvalue arg = ctx.function->getArg(with_fnarg_at);
-    auto const_index =
-      llvm::ConstantInt::get(ctx.llvm, llvm::APInt{32, to_reg, false});
-    if (arg->getType()->getScalarType()->isIntegerTy()) {
-      if (arg->getType()->getScalarType() == ctx.types._bool) {
-        // front_facing
-        return store_at_vec4_array_masked(
-          ctx.resource.input.ptr_int4, const_index,
-          ctx.builder.CreateSExt(arg, ctx.types._int), mask
-        );
-      }
-      if (arg->getType()->getScalarType() != ctx.types._int) {
-        return throwUnsupported("TODO: handle non 32 bit integer");
-      }
-      return store_at_vec4_array_masked(
-        ctx.resource.input.ptr_int4, const_index, arg, mask
-      );
-    } else if (arg->getType()->getScalarType()->isFloatTy()) {
-      if (fix_w_component && isa<llvm::FixedVectorType>(arg->getType()) &&
-          cast<llvm::FixedVectorType>(arg->getType())->getNumElements() == 4) {
-        auto w_component = ctx.builder.CreateExtractElement(arg, 3);
-        auto rcp_w = ctx.builder.CreateFDiv(
-          llvm::ConstantFP::get(ctx.types._float, 1), w_component
-        );
-        arg = ctx.builder.CreateInsertElement(arg, rcp_w, 3);
-      }
-      return store_at_vec4_array_masked(
-        ctx.resource.input.ptr_float4, const_index, arg, mask
-      );
-    } else {
-      return throwUnsupported(
-        "TODO: unhandled input type? might be interpolant..."
-      );
-    }
-  });
-};
-
-IREffect init_input_reg_with_interpolation(
-  uint32_t with_fnarg_at, uint32_t to_reg, uint32_t mask,
-  air::Interpolation interpolation, uint32_t sampleidx_at
-) {
-  return make_effect_bind([=](context ctx) -> IREffect {
-    pvalue interpolant = ctx.function->getArg(with_fnarg_at);
-    auto const_index = llvm::ConstantInt::get(ctx.llvm, llvm::APInt{32, to_reg, false});
-    pvalue interpolated_val = nullptr;
-    switch(interpolation) {
-    case air::Interpolation::center_perspective:
-      interpolated_val = ctx.air.CreateInterpolateAtCenter(interpolant, true);
-      break;
-    case air::Interpolation::center_no_perspective:
-      interpolated_val = ctx.air.CreateInterpolateAtCenter(interpolant, false);
-      break;
-    case air::Interpolation::centroid_perspective:
-      interpolated_val = ctx.air.CreateInterpolateAtCentroid(interpolant, true);
-      break;
-    case air::Interpolation::centroid_no_perspective:
-      interpolated_val = ctx.air.CreateInterpolateAtCentroid(interpolant, false);
-      break;
-    case air::Interpolation::sample_perspective:
-      assert(~sampleidx_at);
-      interpolated_val = ctx.air.CreateInterpolateAtSample(interpolant, ctx.function->getArg(sampleidx_at), true);
-      break;
-    case air::Interpolation::sample_no_perspective:
-      assert(~sampleidx_at);
-      interpolated_val = ctx.air.CreateInterpolateAtSample(interpolant, ctx.function->getArg(sampleidx_at), false);
-      break;
-    case air::Interpolation::flat:
-      assert(0 && "unhandled interpolant mode");
-      break;
-    }
-    co_return co_yield store_at_vec4_array_masked(
-        ctx.resource.input.ptr_float4, const_index, interpolated_val, mask
-    );
-  });
-}
-
-std::function<IRValue(pvalue)>
-pop_output_reg(uint32_t from_reg, uint32_t mask, uint32_t to_element) {
-  return [=](pvalue ret) {
-    return make_irvalue_bind([=](context ctx) {
-      auto const_index =
-        llvm::ConstantInt::get(ctx.llvm, llvm::APInt{32, from_reg, false});
-      return load_from_array_at(
-               ctx.resource.output.ptr_int4, const_index
-             ) >>= [=, &ctx](auto ivec4) {
-        auto desired_type =
-          ctx.function->getReturnType()->getStructElementType(to_element);
-        return to_desired_type_from_int_vec4(ivec4, desired_type, mask) |
-               [=, &ctx](auto value) {
-                 return ctx.builder.CreateInsertValue(ret, value, {to_element});
-               };
-      };
-    });
-  };
-}
-
-std::function<IRValue(pvalue)> pop_output_reg_fix_unorm(uint32_t from_reg, uint32_t mask, uint32_t to_element) {
-  return [=](pvalue ret) {
-    return make_irvalue_bind([=](context ctx) -> IRValue {
-      auto const_index = llvm::ConstantInt::get(ctx.llvm, llvm::APInt{32, from_reg, false});
-      auto fvec4 = co_yield load_from_array_at(ctx.resource.output.ptr_float4, const_index);
-      auto fixed = ctx.builder.CreateFSub(fvec4, co_yield get_float4_splat(1.0f / 127500.0f) /* magic delta ?! */);
-      co_return ctx.builder.CreateInsertValue(ret, fixed, {to_element});
-    });
-  };
-}
-
-std::function<IRValue(pvalue)> pop_output_reg_sanitize_pos(uint32_t from_reg, uint32_t mask, uint32_t to_element) {
-  return [=](pvalue ret) {
-    return make_irvalue_bind([=](context ctx) -> IRValue {
-      auto const_index = llvm::ConstantInt::get(ctx.llvm, llvm::APInt{32, from_reg, false});
-      auto fvec4 = co_yield load_from_array_at(ctx.resource.output.ptr_float4, const_index);
-      auto fixed = ctx.air.SanitizePosition(fvec4);
-      co_return ctx.builder.CreateInsertValue(ret, fixed, {to_element});
-    });
-  };
-}
-
-IREffect
-pop_mesh_output_render_target_array_index(uint32_t from_reg, uint32_t mask, pvalue primitive_id) {
-  auto ctx = co_yield get_context();
-  auto result = co_yield to_desired_type_from_int_vec4(
-      co_yield load_from_array_at(ctx.resource.output.ptr_int4, ctx.builder.getInt32(from_reg)), ctx.types._int, mask
-  );
-  ctx.air.CreateSetMeshRenderTargetArrayIndex(primitive_id, result);
-  co_return {};
-}
-
-IREffect
-pop_mesh_output_viewport_array_index(uint32_t from_reg, uint32_t mask, pvalue primitive_id) {
-  auto ctx = co_yield get_context();
-  auto result = co_yield to_desired_type_from_int_vec4(
-      co_yield load_from_array_at(ctx.resource.output.ptr_int4, ctx.builder.getInt32(from_reg)), ctx.types._int, mask
-  );
-  ctx.air.CreateSetMeshViewportArrayIndex(primitive_id, result);
-  co_return {};
-}
-
-IREffect
-pop_mesh_output_position(uint32_t from_reg, uint32_t mask, pvalue vertex_id) {
-  auto ctx = co_yield get_context();
-  auto result = co_yield load_from_array_at(ctx.resource.output.ptr_float4, ctx.builder.getInt32(from_reg));
-  ctx.air.CreateSetMeshPosition(vertex_id, result);
-  co_return {};
-}
-
-IREffect
-pop_mesh_output_vertex_data(
-    uint32_t from_reg, uint32_t mask, uint32_t idx, pvalue vertex_id, air::MSLScalerOrVectorType desired_type
-) {
-  auto ctx = co_yield get_context();
-  auto result = co_yield to_desired_type_from_int_vec4(
-      co_yield load_from_vec4_array_masked(ctx.resource.output.ptr_int4, ctx.builder.getInt32(from_reg), mask),
-      air::get_llvm_type(desired_type, ctx.llvm), mask
-  );
-  ctx.air.CreateSetMeshVertexData(vertex_id, idx, result);
-  co_return {};
-}
-
-IREffect pull_vertex_input(
-  air::FunctionSignatureBuilder &func_signature, uint32_t to_reg, uint32_t mask,
-  SM50_IA_INPUT_ELEMENT element_info, uint32_t slot_mask
-) {
-  auto vbuf_table = func_signature.DefineInput(air::ArgumentBindingBuffer{
-    .buffer_size = {},
-    .location_index = SM50_BINDING_INDEX_VERTEX_BUFFER,
-    .array_size = 0,
-    .memory_access = air::MemoryAccess::read,
-    .address_space = air::AddressSpace::constant,
-    .type = air::msl_uint,
-    .arg_name = "vertex_buffers",
-    .raster_order_group = {},
-  });
-  return make_effect_bind([=](context ctx) -> IREffect {
-    auto &types = ctx.types;
-    auto &builder = ctx.builder;
-    pvalue index;                   // uint
-    if (element_info.step_function) // per_instance
-    {
-      if (element_info.step_rate) {
-        index = builder.CreateAdd(
-          ctx.resource.base_instance_id,
-          builder.CreateUDiv(
-            ctx.resource.instance_id,
-            builder.getInt32(element_info.step_rate)
-          )
-        );
-      } else {
-        // 0 takes special meaning, that the instance data should never be
-        // stepped at all
-        index = ctx.resource.base_instance_id;
-      }
-    } else {
-      index = ctx.resource.vertex_id_with_base;
-    }
-    if (!ctx.resource.vertex_buffer_table) {
-      ctx.resource.vertex_buffer_table = ctx.builder.CreateBitCast(
-        ctx.function->getArg(vbuf_table),
-        ctx.types._dxmt_vertex_buffer_entry->getPointerTo((uint32_t
-        )air::AddressSpace::constant)
-      );
-    }
-    unsigned int shift = 32u - element_info.slot;
-    unsigned int vertex_buffer_entry_index =
-      element_info.slot ? __builtin_popcount((slot_mask << shift) >> shift) : 0;
-    auto vertex_buffer_table = ctx.resource.vertex_buffer_table;
-    auto vertex_buffer_entry = builder.CreateLoad(
-      types._dxmt_vertex_buffer_entry,
-      builder.CreateConstGEP1_32(
-        types._dxmt_vertex_buffer_entry, vertex_buffer_table,
-        vertex_buffer_entry_index
-      )
-    );
-    auto base_addr = builder.CreateExtractValue(vertex_buffer_entry, {0});
-    auto stride = builder.CreateExtractValue(vertex_buffer_entry, {1});
-    auto byte_offset = builder.CreateAdd(
-      builder.CreateMul(stride, index),
-      builder.getInt32(element_info.aligned_byte_offset)
-    );
-    auto vec4 = co_yield air::pull_vec4_from_addr(
-      (air::MTLAttributeFormat)element_info.format, base_addr, byte_offset
-    );
-    if (vec4->getType() == types._float4) {
-      co_yield store_at_vec4_array_masked(
-        ctx.resource.input.ptr_float4, builder.getInt32(element_info.reg), vec4,
-        mask
-      );
-    } else {
-      assert(vec4->getType() == types._int4);
-      co_yield store_at_vec4_array_masked(
-        ctx.resource.input.ptr_int4, builder.getInt32(element_info.reg), vec4,
-        mask
-      );
-    }
-    co_return {};
-  });
-};
-
 llvm::Expected<llvm::BasicBlock *> convert_basicblocks(
   BasicBlock *entry, context &ctx, llvm::BasicBlock *return_bb
 ) {
@@ -626,10 +204,3 @@ llvm::Expected<llvm::BasicBlock *> convert_basicblocks(
 }
 
 } // namespace dxmt::dxbc
-
-template <>
-struct environment_cast<::dxmt::dxbc::context, ::dxmt::air::AIRBuilderContext> {
-  ::dxmt::air::AIRBuilderContext cast(const ::dxmt::dxbc::context &src) {
-    return {src.llvm, src.module, src.builder, src.types, src.air};
-  };
-};
\ No newline at end of file
diff --git a/src/airconv/dxbc_converter_gs.cpp b/src/airconv/dxbc_converter_gs.cpp
index d237b9deb..17ac78fb2 100644
--- a/src/airconv/dxbc_converter_gs.cpp
+++ b/src/airconv/dxbc_converter_gs.cpp
@@ -1,4 +1,3 @@
-#include "air_operations.hpp"
 #include "air_signature.hpp"
 #include "airconv_error.hpp"
 #include "dxbc_converter.hpp"
@@ -121,8 +120,8 @@ convert_dxbc_geometry_shader(
   SM50_SHADER_ROOT_SIGNATURE_DATA *rootsig = nullptr;
   args_get_data<SM50_SHADER_ROOT_SIGNATURE, SM50_SHADER_ROOT_SIGNATURE_DATA>(pArgs, &rootsig);
 
-  IREffect prologue([](auto) { return std::monostate(); });
-  IRValue epilogue([](struct context ctx) -> pvalue { return nullptr; });
+  Prologues prologue;
+  Epilogues epilogue;
 
   io_binding_map resource_map;
   air::AirType types(context);
@@ -216,8 +215,10 @@ convert_dxbc_geometry_shader(
     }
   };
 
+  IOHelper io(air, function, types, resource_map);
+
   if (topology == air::MeshOutputTopology::Triangle) {
-    resource_map.call_emit = [&]() -> IREffect {
+    resource_map.call_emit = [&]() {
       auto current_write_vertex = builder.CreateLoad(types._int, next_write_vertex);
       auto current_vertex_offset = builder.CreateLoad(types._int, vertex_offset);
       auto current_vertex_with_offset = builder.CreateAdd(current_vertex_offset, current_write_vertex);
@@ -231,7 +232,7 @@ convert_dxbc_geometry_shader(
 
       MeshOutputContext gs_out_ctx{current_vertex_with_offset, current_primitive_idx};
       for (auto &h : gs_output_handlers) {
-        co_yield h(gs_out_ctx);
+        h(gs_out_ctx, io);
       }
       emit_clip_distances(current_vertex_with_offset);
 
@@ -249,10 +250,8 @@ convert_dxbc_geometry_shader(
       );
 
       builder.CreateStore(builder.CreateAdd(one_const, current_write_vertex), next_write_vertex);
-
-      co_return {};
     };
-    resource_map.call_cut = [&]() -> IREffect {
+    resource_map.call_cut = [&]() {
       auto current_write_vertex = builder.CreateLoad(types._int, next_write_vertex);
       builder.CreateStore(zero_const, next_write_vertex);
 
@@ -267,10 +266,9 @@ convert_dxbc_geometry_shader(
       builder.CreateStore(
           builder.CreateAdd(builder.CreateLoad(types._int, primitive_count), add_primitive_count), primitive_count
       );
-      co_return {};
     };
   } else if (topology == air::MeshOutputTopology::Line) {
-    resource_map.call_emit = [&]() -> IREffect {
+    resource_map.call_emit = [&]() {
       auto current_write_vertex = builder.CreateLoad(types._int, next_write_vertex);
       builder.CreateStore(builder.CreateAdd(one_const, current_write_vertex), next_write_vertex);
 
@@ -282,7 +280,7 @@ convert_dxbc_geometry_shader(
 
       MeshOutputContext gs_out_ctx{current_vertex_with_offset, current_primitive_idx};
       for (auto &h : gs_output_handlers) {
-        co_yield h(gs_out_ctx);
+        h(gs_out_ctx, io);
       }
       emit_clip_distances(current_vertex_with_offset);
 
@@ -294,10 +292,8 @@ convert_dxbc_geometry_shader(
           builder.CreateAdd(double_primitive_idx, one_const),
           builder.CreateAdd(current_vertex_with_offset, one_const)
       );
-
-      co_return {};
     };
-    resource_map.call_cut = [&]() -> IREffect {
+    resource_map.call_cut = [&]() {
       auto current_write_vertex = builder.CreateLoad(types._int, next_write_vertex);
       builder.CreateStore(zero_const, next_write_vertex);
 
@@ -312,25 +308,22 @@ convert_dxbc_geometry_shader(
       builder.CreateStore(
           builder.CreateAdd(builder.CreateLoad(types._int, primitive_count), add_primitive_count), primitive_count
       );
-      co_return {};
     };
   } else {
-    resource_map.call_emit = [&]() -> IREffect {
+    resource_map.call_emit = [&]() {
       // only one accumulator to maintain, simple one ~
       auto current_write_vertex = builder.CreateLoad(types._int, next_write_vertex);
       builder.CreateStore(builder.CreateAdd(one_const, current_write_vertex), next_write_vertex);
 
       MeshOutputContext gs_out_ctx{current_write_vertex, current_write_vertex};
       for (auto &h : gs_output_handlers) {
-        co_yield h(gs_out_ctx);
+        h(gs_out_ctx, io);
       }
       emit_clip_distances(current_write_vertex);
       air.CreateSetMeshIndex(current_write_vertex, current_write_vertex);
-      co_return {};
     };
-    resource_map.call_cut = []() -> IREffect {
+    resource_map.call_cut = []() {
       // there is nothing to cut!
-      co_return {};
     };
   }
 
@@ -492,7 +485,7 @@ convert_dxbc_geometry_shader(
       .metal_version = metal_version,
   };
 
-  if (auto err = prologue.build(ctx).takeError()) {
+  if (auto err = prologue.Run(io)) {
     return err;
   }
 
@@ -503,12 +496,11 @@ convert_dxbc_geometry_shader(
   builder.CreateBr(real_entry.get());
   builder.SetInsertPoint(epilogue_bb);
 
-  if (auto err = epilogue.build(ctx).takeError()) {
-    return err;
-  }
-  if (auto err = resource_map.call_cut().build(ctx).takeError()) {
+  auto epilogue_result = epilogue.Run(function->getReturnType(), io);
+  if (auto err = epilogue_result.takeError()) {
     return err;
   }
+  resource_map.call_cut();
   air.CreateSetMeshPrimitiveCount(builder.CreateLoad(types._int, primitive_count));
 
   builder.CreateRetVoid();
@@ -548,14 +540,8 @@ convert_dxbc_vertex_for_geometry_shader(
   bool is_triadj_strip = is_strip && pGeometryStage->gs_input_primitive == D3D10_SB_PRIMITIVE_TRIANGLE_ADJ;
   bool is_indexed_draw = ia_layout && ia_layout->index_buffer_format > 0;
 
-  IREffect prologue([](auto) { return std::monostate(); });
-  IRValue epilogue([](struct context ctx) -> pvalue {
-    auto retTy = ctx.function->getReturnType();
-    if (retTy->isVoidTy()) {
-      return nullptr;
-    }
-    return llvm::UndefValue::get(retTy);
-  });
+  Prologues prologue;
+  Epilogues epilogue;
 
   io_binding_map resource_map;
   air::AirType types(context);
@@ -774,7 +760,9 @@ convert_dxbc_vertex_for_geometry_shader(
   resource_map.base_instance_id = builder.CreateExtractValue(draw_arguments, is_indexed_draw ? 4 : 3);
   resource_map.instance_id_with_base = builder.CreateAdd(resource_map.instance_id, resource_map.base_instance_id);
 
-  if (auto err = prologue.build(ctx).takeError()) {
+  IOHelper io(air, function, types, resource_map);
+
+  if (auto err = prologue.Run(io)) {
     return err;
   }
 
@@ -785,7 +773,7 @@ convert_dxbc_vertex_for_geometry_shader(
   builder.CreateBr(real_entry.get());
 
   builder.SetInsertPoint(epilogue_bb);
-  auto epilogue_result = epilogue.build(ctx);
+  auto epilogue_result = epilogue.Run(function->getReturnType(), io);
   if (auto err = epilogue_result.takeError()) {
     return err;
   }
@@ -883,10 +871,3 @@ convert_dxbc_vertex_for_geometry_shader(
 };
 
 } // namespace dxmt::dxbc
-
-template <> struct environment_cast<::dxmt::dxbc::context, ::dxmt::air::AIRBuilderContext> {
-  ::dxmt::air::AIRBuilderContext
-  cast(const ::dxmt::dxbc::context &src) {
-    return {src.llvm, src.module, src.builder, src.types, src.air};
-  };
-};
diff --git a/src/airconv/dxbc_converter_ts.cpp b/src/airconv/dxbc_converter_ts.cpp
index fd32c70bb..d97bac0da 100644
--- a/src/airconv/dxbc_converter_ts.cpp
+++ b/src/airconv/dxbc_converter_ts.cpp
@@ -216,18 +216,10 @@ convert_dxbc_vertex_hull_shader(
 
   bool is_indexed_draw = ia_layout && ia_layout->index_buffer_format > 0;
 
-  IREffect prologue_vs([](auto) { return std::monostate(); });
-  IREffect prologue_hs([](auto) { return std::monostate(); });
-  IRValue epilogue_vs([](struct context ctx) -> pvalue {
-    // just return null since the function returns void
-    // and we handle output registers directly
-    return nullptr;
-  });
-  IRValue epilogue_hs([](struct context ctx) -> pvalue {
-    // just return null since the function returns void
-    // and we handle output registers directly
-    return nullptr;
-  });
+  Prologues prologue_vs;
+  Prologues prologue_hs;
+  Epilogues epilogue_vs;
+  Epilogues epilogue_hs;
 
   io_binding_map resource_map_vs;
   io_binding_map resource_map_hs;
@@ -465,7 +457,9 @@ convert_dxbc_vertex_hull_shader(
     resource_map.base_instance_id = builder.CreateExtractValue(draw_arguments, is_indexed_draw ? 4 : 3);
     resource_map.instance_id_with_base = builder.CreateAdd(resource_map.instance_id, resource_map.base_instance_id);
 
-    if (auto err = prologue.build(ctx).takeError()) {
+    IOHelper io(air, function, types, resource_map);
+
+    if (auto err = prologue.Run(io)) {
       return err;
     }
     auto real_entry = convert_basicblocks(pVertexStage->entry(), ctx, epilogue_bb);
@@ -475,7 +469,7 @@ convert_dxbc_vertex_hull_shader(
     builder.CreateBr(real_entry.get());
 
     builder.SetInsertPoint(epilogue_bb);
-    auto epilogue_result = epilogue.build(ctx);
+    auto epilogue_result = epilogue.Run(function->getReturnType(), io);
     if (auto err = epilogue_result.takeError()) {
       return err;
     }
@@ -609,7 +603,9 @@ convert_dxbc_vertex_hull_shader(
     };
     dxbc::Converter dxbc(ctx.air, ctx, ctx.resource);
 
-    if (auto err = prologue.build(ctx).takeError()) {
+    IOHelper io(air, function, types, resource_map);
+
+    if (auto err = prologue.Run(io)) {
       return err;
     }
     auto real_entry = convert_basicblocks(pHullStage->entry(), ctx, epilogue_bb);
@@ -632,7 +628,8 @@ convert_dxbc_vertex_hull_shader(
     );
 
     builder.SetInsertPoint(write_patch_constant);
-    if (auto err = epilogue.build(ctx).takeError()) {
+    auto epilogue_result = epilogue.Run(function->getReturnType(), io);
+    if (auto err = epilogue_result.takeError()) {
       return err;
     }
 
@@ -756,10 +753,8 @@ convert_dxbc_tesselator_domain_shader(
 
   auto [final_maxtessfactor, factor_int] = get_final_maxtessfactor(pHullStage, pArgs);
 
-  IREffect prologue([](auto) { return std::monostate(); });
-  IRValue epilogue([](struct context ctx) -> pvalue {
-    return nullptr; // a mesh shader...
-  });
+  Prologues prologue;
+  Epilogues epilogue;
 
   io_binding_map resource_map;
   air::AirType types(context);
@@ -993,10 +988,9 @@ convert_dxbc_tesselator_domain_shader(
   auto vertex_id = actual_thread_index;
   auto primitive_id = actual_thread_index; // 
   MeshOutputContext gs_out_ctx{vertex_id, primitive_id};
+  IOHelper io(air, function, types, resource_map);
   for (auto &h : ds_output_handlers) {
-    if (auto err = h(gs_out_ctx).build(ctx).takeError()) {
-      return err;
-    }
+    h(gs_out_ctx, io);
   }
 
   for (auto x : llvm::enumerate(pShaderInternal->clip_distance_scalars)) {
diff --git a/src/airconv/dxbc_signature.cpp b/src/airconv/dxbc_signature.cpp
index 72d67f94b..79b70dda3 100644
--- a/src/airconv/dxbc_signature.cpp
+++ b/src/airconv/dxbc_signature.cpp
@@ -90,22 +90,16 @@ void handle_signature_vs(
     switch (sgv) {
     case D3D10_SB_NAME_VERTEX_ID: {
       signature_handlers.push_back([=](SignatureContext &ctx) {
-        ctx.prologue << make_effect_bind([=](struct context ctx) {
-          return store_at_vec4_array_masked(
-            ctx.resource.input.ptr_int4, ctx.builder.getInt32(reg),
-            ctx.resource.vertex_id, mask
-          );
+        ctx.prologue.Add([=](IOHelper &io) -> llvm::Error {
+          return io.InitializeVertexIDInput(reg, mask);
         });
       });
       break;
     }
     case D3D10_SB_NAME_INSTANCE_ID: {
       signature_handlers.push_back([=](SignatureContext &ctx) {
-        ctx.prologue << make_effect_bind([=](struct context ctx) {
-          return store_at_vec4_array_masked(
-            ctx.resource.input.ptr_int4, ctx.builder.getInt32(reg),
-            ctx.resource.instance_id, mask
-          );
+        ctx.prologue.Add([=](IOHelper &io) -> llvm::Error {
+          return io.InitializeInstanceIDInput(reg, mask);
         });
       });
       break;
@@ -137,10 +131,21 @@ void handle_signature_vs(
           if (ctx.ia_layout) {
             for (unsigned i = 0; i < ctx.ia_layout->num_elements; i++) {
               if (ctx.ia_layout->elements[i].reg == reg) {
-                ctx.prologue << pull_vertex_input(
-                  ctx.func_signature, reg, mask, ctx.ia_layout->elements[i],
-                  ctx.ia_layout->slot_mask
-                );
+                auto element_info = ctx.ia_layout->elements[i];
+                auto slot_mask = ctx.ia_layout->slot_mask;
+                auto vbuf_table = ctx.func_signature.DefineInput(air::ArgumentBindingBuffer{
+                  .buffer_size = {},
+                  .location_index = SM50_BINDING_INDEX_VERTEX_BUFFER,
+                  .array_size = 0,
+                  .memory_access = air::MemoryAccess::read,
+                  .address_space = air::AddressSpace::constant,
+                  .type = air::msl_uint,
+                  .arg_name = "vertex_buffers",
+                  .raster_order_group = {},
+                });
+                ctx.prologue.Add([=](IOHelper &io) -> llvm::Error {
+                  return io.PullVertexInput(vbuf_table, reg, mask, element_info, slot_mask);
+                });
                 break;
               }
             }
@@ -148,7 +153,9 @@ void handle_signature_vs(
             auto assigned_index = ctx.func_signature.DefineInput(
               InputVertexStageIn{.attribute = reg, .type = type, .name = name}
             );
-            ctx.prologue << init_input_reg(assigned_index, reg, mask);
+            ctx.prologue.Add([=](IOHelper &io) -> llvm::Error {
+              return io.InitializeInput(assigned_index, reg, mask);
+            });
           }
         }
       );
@@ -180,7 +187,9 @@ void handle_signature_vs(
       signature_handlers.push_back([=](SignatureContext &ctx) {
         if (ctx.skip_vertex_output)
           return;
-        ctx.epilogue >> pop_output_reg(reg, mask, assigned_index);
+        ctx.epilogue.Add([=](llvm::Value *ret, IOHelper &io) -> llvm::Expected<llvm::Value *> {
+          return io.PopulateOutput(reg, mask, ret, assigned_index);
+        });
       });
       max_output_register = std::max(reg + 1, max_output_register);
       break;
@@ -200,7 +209,9 @@ void handle_signature_vs(
       signature_handlers.push_back([=](SignatureContext &ctx) {
         if (ctx.skip_vertex_output)
           return;
-        ctx.epilogue >> pop_output_reg(reg, mask, assigned_index);
+        ctx.epilogue.Add([=](llvm::Value *ret, IOHelper &io) -> llvm::Expected<llvm::Value *> {
+          return io.PopulateOutput(reg, mask, ret, assigned_index);
+        });
       });
       max_output_register = std::max(reg + 1, max_output_register);
       break;
@@ -211,7 +222,9 @@ void handle_signature_vs(
       signature_handlers.push_back([=](SignatureContext &ctx) {
         if (ctx.skip_vertex_output)
           return;
-        ctx.epilogue >> pop_output_reg_sanitize_pos(reg, mask, assigned_index);
+        ctx.epilogue.Add([=](llvm::Value *ret, IOHelper &io) -> llvm::Expected<llvm::Value *> {
+          return io.PopulateOutputSanitizePosition(reg, mask, ret, assigned_index);
+        });
       });
       break;
     }
@@ -221,7 +234,9 @@ void handle_signature_vs(
       signature_handlers.push_back([=](SignatureContext &ctx) {
         if (ctx.skip_vertex_output)
           return;
-        ctx.epilogue >> pop_output_reg(reg, mask, assigned_index);
+        ctx.epilogue.Add([=](llvm::Value *ret, IOHelper &io) -> llvm::Expected<llvm::Value *> {
+          return io.PopulateOutput(reg, mask, ret, assigned_index);
+        });
       });
       break;
     }
@@ -231,7 +246,9 @@ void handle_signature_vs(
       signature_handlers.push_back([=](SignatureContext &ctx) {
         if (ctx.skip_vertex_output)
           return;
-        ctx.epilogue >> pop_output_reg(reg, mask, assigned_index);
+        ctx.epilogue.Add([=](llvm::Value *ret, IOHelper &io) -> llvm::Expected<llvm::Value *> {
+          return io.PopulateOutput(reg, mask, ret, assigned_index);
+        });
       });
       break;
     }
@@ -261,7 +278,9 @@ void handle_signature_vs(
       signature_handlers.push_back([=](SignatureContext &ctx) {
         if (ctx.skip_vertex_output)
           return;
-        ctx.epilogue >> pop_output_reg(reg, mask, assigned_index);
+        ctx.epilogue.Add([=](llvm::Value *ret, IOHelper &io) -> llvm::Expected<llvm::Value *> {
+          return io.PopulateOutput(reg, mask, ret, assigned_index);
+        });
       });
       break;
     }
@@ -326,10 +345,8 @@ void handle_signature_ps(
     case D3D11_SB_OPERAND_TYPE_INPUT_COVERAGE_MASK: {
       auto assigned_index = func_signature.DefineInput(InputInputCoverage{});
       signature_handlers.push_back([=](SignatureContext &ctx) {
-        ctx.prologue << make_effect([=](struct context ctx) {
-          auto attr = ctx.function->getArg(assigned_index);
-          ctx.resource.coverage_mask_arg = attr;
-          return std::monostate{};
+        ctx.prologue.Add([=](IOHelper &io) -> llvm::Error {
+          return io.InitializeInputAttribute(InputAttribute::CoverageMask, assigned_index);
         });
       });
       break;
@@ -392,9 +409,9 @@ void handle_signature_ps(
       break;
     }
     signature_handlers.push_back([=](SignatureContext &ctx) {
-      ctx.prologue << init_input_reg(
-        assigned_index, reg, mask, siv == D3D10_SB_NAME_POSITION
-      );
+      ctx.prologue.Add([=](IOHelper &io) -> llvm::Error {
+        return io.InitializeInput(assigned_index, reg, mask, siv == D3D10_SB_NAME_POSITION);
+      });
     });
     max_input_register = std::max(reg + 1, max_input_register);
     break;
@@ -423,7 +440,9 @@ void handle_signature_ps(
       break;
     }
     signature_handlers.push_back([=](SignatureContext &ctx) {
-      ctx.prologue << init_input_reg(assigned_index, reg, mask);
+      ctx.prologue.Add([=](IOHelper &io) -> llvm::Error {
+        return io.InitializeInput(assigned_index, reg, mask);
+      });
     });
     max_input_register = std::max(reg + 1, max_input_register);
     break;
@@ -447,11 +466,7 @@ void handle_signature_ps(
       if (pull_mode) {
         assert(type == RegisterComponentType::Float && "otherwise the input register contains mixed data type");
         ctx.resource.interpolant_map[reg] = interpolant_descriptor{
-          [=](auto) {
-            return make_irvalue([=](struct context ctx) {
-              return ctx.function->getArg(assigned_index);
-            });
-          },
+          assigned_index,
           (interpolation == air::Interpolation::sample_perspective ||
            interpolation == air::Interpolation::center_perspective ||
            interpolation == air::Interpolation::centroid_perspective)
@@ -465,9 +480,13 @@ void handle_signature_ps(
         default:
           break;
         }
-        ctx.prologue << init_input_reg_with_interpolation(assigned_index, reg, mask, interpolation, sampleidx_at);
+        ctx.prologue.Add([=](IOHelper &io) -> llvm::Error {
+          return io.InitializeInterpolatedInput(assigned_index, reg, mask, interpolation, sampleidx_at);
+        });
       } else {
-        ctx.prologue << init_input_reg(assigned_index, reg, mask);
+        ctx.prologue.Add([=](IOHelper &io) -> llvm::Error {
+          return io.InitializeInput(assigned_index, reg, mask);
+        });
       }
     });
     break;
@@ -496,14 +515,8 @@ void handle_signature_ps(
     case D3D11_SB_OPERAND_TYPE_OUTPUT_DEPTH_GREATER_EQUAL:
     case D3D11_SB_OPERAND_TYPE_OUTPUT_DEPTH_LESS_EQUAL: {
       signature_handlers.push_back([=](SignatureContext &ctx) {
-        ctx.prologue << make_effect([](struct context ctx) -> std::monostate {
-          assert(
-            ctx.resource.depth_output_reg == nullptr &&
-            "otherwise oDepth is defined twice"
-          );
-          ctx.resource.depth_output_reg =
-            ctx.builder.CreateAlloca(ctx.types._float);
-          return {};
+        ctx.prologue.Add([](IOHelper &io) -> llvm::Error {
+          return io.CreateDepthOutputRegister();
         });
       });
       signature_handlers.push_back([=](SignatureContext &ctx) {
@@ -517,75 +530,37 @@ void handle_signature_ps(
               ? DepthArgument::less
               : DepthArgument::any
         });
-        ctx.epilogue >> [=](pvalue v) {
-          return make_irvalue([=](struct context ctx) {
-            return ctx.builder.CreateInsertValue(
-              v,
-              ctx.builder.CreateLoad(
-                ctx.types._float,
-                ctx.builder.CreateConstInBoundsGEP1_32(
-                  ctx.types._float, ctx.resource.depth_output_reg, 0
-                )
-              ),
-              {assigned_index}
-            );
-          });
-        };
+        ctx.epilogue.Add([=](llvm::Value *ret, IOHelper &io) -> llvm::Expected<llvm::Value *> {
+          return io.PopulateOutputDepth(ret, assigned_index);
+        });
       });
       break;
     }
     case D3D11_SB_OPERAND_TYPE_OUTPUT_STENCIL_REF: {
       signature_handlers.push_back([=](SignatureContext &ctx) {
-        ctx.prologue << make_effect([](struct context ctx) -> std::monostate {
-          assert(ctx.resource.stencil_ref_reg == nullptr && "otherwise oStencil is defined twice");
-          ctx.resource.stencil_ref_reg = ctx.builder.CreateAlloca(ctx.types._int);
-          return {};
+        ctx.prologue.Add([](IOHelper &io) -> llvm::Error {
+          return io.CreateStencilRefOutputRegister();
         });
       });
       auto assigned_index = func_signature.DefineOutput(OutputStencilRef{});
       signature_handlers.push_back([=](SignatureContext &ctx) {
-        ctx.epilogue >> [=](pvalue v) {
-          return make_irvalue([=](struct context ctx) {
-            auto ostencil = ctx.builder.CreateLoad(
-                ctx.types._int, ctx.builder.CreateConstInBoundsGEP1_32(ctx.types._int, ctx.resource.stencil_ref_reg, 0)
-            );
-            return ctx.builder.CreateInsertValue(v, ostencil, {assigned_index});
-          });
-        };
+        ctx.epilogue.Add([=](llvm::Value *ret, IOHelper &io) -> llvm::Expected<llvm::Value *> {
+          return io.PopulateOutputStencilRef(ret, assigned_index);
+        });
       });
       break;
     }
     case D3D10_SB_OPERAND_TYPE_OUTPUT_COVERAGE_MASK: {
       signature_handlers.push_back([=](SignatureContext &ctx) {
-        ctx.prologue << make_effect([](struct context ctx) -> std::monostate {
-          assert(
-            ctx.resource.coverage_mask_reg == nullptr &&
-            "otherwise oMask is defined twice"
-          );
-          ctx.resource.coverage_mask_reg =
-            ctx.builder.CreateAlloca(ctx.types._int);
-          return {};
+        ctx.prologue.Add([](IOHelper &io) -> llvm::Error {
+          return io.CreateCoverageMaskOutputRegister();
         });
       });
       auto assigned_index = func_signature.DefineOutput(OutputCoverageMask{});
       signature_handlers.push_back([=](SignatureContext &ctx) {
-        ctx.epilogue >> [=](pvalue v) {
-          return make_irvalue([=](struct context ctx) {
-            auto odepth = ctx.builder.CreateLoad(
-              ctx.types._int,
-              ctx.builder.CreateConstInBoundsGEP1_32(
-                ctx.types._int, ctx.resource.coverage_mask_reg, 0
-              )
-            );
-            return ctx.builder.CreateInsertValue(
-              v,
-              ctx.pso_sample_mask != 0xffffffff
-                ? ctx.builder.CreateAnd(odepth, ctx.pso_sample_mask)
-                : odepth,
-              {assigned_index}
-            );
-          });
-        };
+        ctx.epilogue.Add([=, sample_mask = ctx.pso_sample_mask](llvm::Value *ret, IOHelper &io) -> llvm::Expected<llvm::Value *> {
+          return io.PopulateOutputCoverageMask(ret, assigned_index, sample_mask);
+        });
       });
       sm50_shader->ps_has_coverage_output = 1;
       break;
@@ -623,9 +598,13 @@ void handle_signature_ps(
           });
         }
         if (type == RegisterComponentType::Float && ctx.unorm_output_reg_mask & (1 << reg))
-          ctx.epilogue >> pop_output_reg_fix_unorm(reg, mask, assigned_index);
+          ctx.epilogue.Add([=](llvm::Value *ret, IOHelper &io) -> llvm::Expected<llvm::Value *> {
+            return io.PopulateOutputFixUnorm(reg, mask, ret, assigned_index);
+          });
         else
-          ctx.epilogue >> pop_output_reg(reg, mask, assigned_index);
+          ctx.epilogue.Add([=](llvm::Value *ret, IOHelper &io) -> llvm::Expected<llvm::Value *> {
+            return io.PopulateOutput(reg, mask, ret, assigned_index);
+          });
       });
       break;
     }
@@ -915,22 +894,22 @@ void handle_signature_ds(
     }
     case D3D10_SB_NAME_POSITION: {
       func_signature.DefineMeshVertexOutput(OutputPosition{.type = msl_float4});
-      mesh_output_handlers.push_back([=](MeshOutputContext& output) -> IREffect {
-        return pop_mesh_output_position(reg, mask, output.vertex_id);
+      mesh_output_handlers.push_back([=](MeshOutputContext &output, IOHelper &io) {
+        io.PopulateMeshOutputPosition(reg, mask, output.vertex_id);
       });
       break;
     }
     case D3D10_SB_NAME_RENDER_TARGET_ARRAY_INDEX: {
       func_signature.DefineMeshPrimitiveOutput(OutputRenderTargetArrayIndex{});
-      mesh_output_handlers.push_back([=](MeshOutputContext& output) -> IREffect {
-        return pop_mesh_output_render_target_array_index(reg, mask, output.primitive_id);
+      mesh_output_handlers.push_back([=](MeshOutputContext &output, IOHelper &io) {
+        io.PopulateMeshOutputRenderTargetArrayIndex(reg, mask, output.primitive_id);
       });
       break;
     }
     case D3D10_SB_NAME_VIEWPORT_ARRAY_INDEX: {
       func_signature.DefineMeshPrimitiveOutput(OutputViewportArrayIndex{});
-      mesh_output_handlers.push_back([=](MeshOutputContext& output) -> IREffect {
-        return pop_mesh_output_viewport_array_index(reg, mask, output.primitive_id);
+      mesh_output_handlers.push_back([=](MeshOutputContext &output, IOHelper &io) {
+        io.PopulateMeshOutputViewportArrayIndex(reg, mask, output.primitive_id);
       });
       break;
     }
@@ -960,8 +939,8 @@ void handle_signature_ds(
         .type = type,
         .index = mesh_vertex_data_index
       });
-      mesh_output_handlers.push_back([=](MeshOutputContext& output) -> IREffect {
-        return pop_mesh_output_vertex_data(reg, mask, mesh_vertex_data_index, output.vertex_id, type);
+      mesh_output_handlers.push_back([=](MeshOutputContext &output, IOHelper &io) {
+        io.PopulateMeshOutputVertexData(reg, mask, mesh_vertex_data_index, output.vertex_id, type);
       });
       break;
     }
@@ -998,10 +977,8 @@ void handle_signature_cs(
       auto assigned_index =
         func_signature.DefineInput(InputThreadPositionInGrid{});
       signature_handlers.push_back([=](SignatureContext &ctx) {
-        ctx.prologue << make_effect([=](struct context ctx) {
-          auto attr = ctx.function->getArg(assigned_index);
-          ctx.resource.thread_id_arg = attr;
-          return std::monostate{};
+        ctx.prologue.Add([=](IOHelper &io) -> llvm::Error {
+          return io.InitializeInputAttribute(InputAttribute::ThreadId, assigned_index);
         });
       });
       break;
@@ -1010,10 +987,8 @@ void handle_signature_cs(
       auto assigned_index =
         func_signature.DefineInput(InputThreadgroupPositionInGrid{});
       signature_handlers.push_back([=](SignatureContext &ctx) {
-        ctx.prologue << make_effect([=](struct context ctx) {
-          auto attr = ctx.function->getArg(assigned_index);
-          ctx.resource.thread_group_id_arg = attr;
-          return std::monostate{};
+        ctx.prologue.Add([=](IOHelper &io) -> llvm::Error {
+          return io.InitializeInputAttribute(InputAttribute::ThreadGroupId, assigned_index);
         });
       });
       break;
@@ -1022,10 +997,8 @@ void handle_signature_cs(
       auto assigned_index =
         func_signature.DefineInput(InputThreadPositionInThreadgroup{});
       signature_handlers.push_back([=](SignatureContext &ctx) {
-        ctx.prologue << make_effect([=](struct context ctx) {
-          auto attr = ctx.function->getArg(assigned_index);
-          ctx.resource.thread_id_in_group_arg = attr;
-          return std::monostate{};
+        ctx.prologue.Add([=](IOHelper &io) -> llvm::Error {
+          return io.InitializeInputAttribute(InputAttribute::ThreadIdInGroup, assigned_index);
         });
       });
       break;
@@ -1034,10 +1007,8 @@ void handle_signature_cs(
       auto assigned_index =
         func_signature.DefineInput(InputThreadIndexInThreadgroup{});
       signature_handlers.push_back([=](SignatureContext &ctx) {
-        ctx.prologue << make_effect([=](struct context ctx) {
-          auto attr = ctx.function->getArg(assigned_index);
-          ctx.resource.thread_id_in_group_flat_arg = attr;
-          return std::monostate{};
+        ctx.prologue.Add([=](IOHelper &io) -> llvm::Error {
+          return io.InitializeInputAttribute(InputAttribute::ThreadIdInGroupFlatten, assigned_index);
         });
       });
       break;
@@ -1147,22 +1118,22 @@ handle_signature_gs(
     }
     case D3D10_SB_NAME_POSITION: {
       func_signature.DefineMeshVertexOutput(OutputPosition{.type = msl_float4});
-      mesh_output_handlers.push_back([=](MeshOutputContext& output) -> IREffect {
-        return pop_mesh_output_position(reg, mask, output.vertex_id);
+      mesh_output_handlers.push_back([=](MeshOutputContext &output, IOHelper &io) {
+        io.PopulateMeshOutputPosition(reg, mask, output.vertex_id);
       });
       break;
     }
     case D3D10_SB_NAME_RENDER_TARGET_ARRAY_INDEX: {
       func_signature.DefineMeshPrimitiveOutput(OutputRenderTargetArrayIndex{});
-      mesh_output_handlers.push_back([=](MeshOutputContext& output) -> IREffect {
-        return pop_mesh_output_render_target_array_index(reg, mask, output.primitive_id);
+      mesh_output_handlers.push_back([=](MeshOutputContext &output, IOHelper &io) {
+        io.PopulateMeshOutputRenderTargetArrayIndex(reg, mask, output.primitive_id);
       });
       break;
     }
     case D3D10_SB_NAME_VIEWPORT_ARRAY_INDEX: {
       func_signature.DefineMeshPrimitiveOutput(OutputViewportArrayIndex{});
-      mesh_output_handlers.push_back([=](MeshOutputContext& output) -> IREffect {
-        return pop_mesh_output_viewport_array_index(reg, mask, output.primitive_id);
+      mesh_output_handlers.push_back([=](MeshOutputContext &output, IOHelper &io) {
+        io.PopulateMeshOutputViewportArrayIndex(reg, mask, output.primitive_id);
       });
       break;
     }
@@ -1192,8 +1163,8 @@ handle_signature_gs(
         .type = type,
         .index = mesh_vertex_data_index
       });
-      mesh_output_handlers.push_back([=](MeshOutputContext& output) -> IREffect {
-        return pop_mesh_output_vertex_data(reg, mask, mesh_vertex_data_index, output.vertex_id, type);
+      mesh_output_handlers.push_back([=](MeshOutputContext &output, IOHelper &io) {
+        io.PopulateMeshOutputVertexData(reg, mask, mesh_vertex_data_index, output.vertex_id, type);
       });
       break;
     }
diff --git a/src/airconv/meson.build b/src/airconv/meson.build
index 8e0c8bedb..40eb7266c 100644
--- a/src/airconv/meson.build
+++ b/src/airconv/meson.build
@@ -2,7 +2,6 @@ airconv_src = files([
  'airconv_context.cpp',
  'air_type.cpp',
  'air_signature.cpp',
- 'air_operations.cpp',
  'dxbc_binding_sm50.cpp',
  'dxbc_binding_rootsig.cpp',
  'dxbc_converter.cpp',
@@ -14,7 +13,9 @@ airconv_src = files([
  'dxbc_signature.cpp',
  'metallib_writer.cpp',
  'nt/air_builder.cpp',
+ 'nt/air_vertex_puller.cpp',
  'nt/dxbc_converter_base.cpp',
+ 'nt/dxbc_io_helper.cpp',
  'transforms/lower_16bit_texread.cpp',
  'transforms/simdgroup_implicit_membarrier.cpp',
 ])
diff --git a/src/airconv/monad.hpp b/src/airconv/monad.hpp
deleted file mode 100644
index a0db1b3b8..000000000
--- a/src/airconv/monad.hpp
+++ /dev/null
@@ -1,248 +0,0 @@
-#pragma once
-#include "llvm/Support/Error.h"
-#include <coroutine>
-#include <functional>
-#include <type_traits>
-#include <utility>
-
-template <typename T, typename G>
-concept MoveAndInvocable =
-  std::is_object_v<T> && std::move_constructible<T> && std::invocable<T, G>;
-
-template <typename Test, template <typename...> class Ref>
-struct is_specialization : std::false_type {};
-
-template <template <typename...> class Ref, typename... Args>
-struct is_specialization<Ref<Args...>, Ref> : std::true_type {};
-
-template <typename Src, typename Dst> struct environment_cast {
-  Dst cast(Src &src);
-};
-
-template <typename Env, typename V> class ReaderIO {
-
-  class BaseFunction {
-  public:
-    virtual llvm::Expected<V> invoke(Env env) = 0;
-    virtual ~BaseFunction() {};
-  };
-
-  template <MoveAndInvocable<Env> Fn>
-  class ErasureFunction : public BaseFunction {
-  public:
-    llvm::Expected<V> invoke(Env env) override { return std::invoke(fn, env); };
-    ErasureFunction(Fn &&ff) : fn(std::forward<Fn>(ff)) {}
-    ~ErasureFunction() = default;
-    ErasureFunction(const ErasureFunction &copy) = delete;
-    ErasureFunction &operator=(const ErasureFunction &copy_assign) = delete;
-
-  private:
-    Fn fn;
-  };
-
-public:
-  template <std::invocable<Env> T> ReaderIO(T &&ff) {
-    ops.push_back(new ErasureFunction<T>(std::forward<T>(ff)));
-  }
-  ~ReaderIO() {
-    destroy();
-  }
-
-  void destroy() {
-    while (!ops.empty()) {
-      delete ops.back();
-      ops.pop_back();
-    }
-  }
-
-  template <typename Env2> ReaderIO(ReaderIO<Env2, V> &&castable) {
-    auto f = [castable = std::move(castable)](Env e) mutable {
-      struct environment_cast<Env, Env2> cast;
-      return castable.build(cast.cast(e));
-    };
-    ops.push_back(new ErasureFunction<decltype(f)>(std::move(f)));
-  }
-
-  ReaderIO(ReaderIO &&other) {
-    ops = std::move(other.ops);
-    other.ops = {};
-  };
-  ReaderIO &operator=(ReaderIO &&move_assign) {
-    ops = std::move(move_assign.ops);
-    move_assign.ops = {};
-    return *this;
-  };
-  ReaderIO(const ReaderIO &copy) = delete;
-  ReaderIO &operator=(const ReaderIO &copy_assign) = delete;
-
-  llvm::Expected<V> build(Env ir) {
-    assert(!ops.empty() && "value has been consumed or moved.");
-    unsigned i = 0;
-    while(i < ops.size()) {
-      llvm::Expected<V> ret = ops[i]->invoke(ir);
-      if (auto err = ret.takeError()) {
-        destroy();
-        return err;
-      }
-      i++;
-      if (i == ops.size()) {
-        destroy();
-        return ret.get();
-      }
-    }
-    __builtin_unreachable();
-  };
-
-  struct promise_type {
-    Env *to_be_filled = nullptr;
-    llvm::Expected<V> *return_value_ = nullptr;
-    ReaderIO<Env, V> get_return_object() {
-      return ReaderIO<Env, V>([this](Env ctx) -> llvm::Expected<V> {
-        this->to_be_filled =
-          &ctx; // I'm sure to_be_filled is only accessed within the scope?
-        auto h = std::coroutine_handle<promise_type>::from_promise(*this);
-        h.resume();
-        assert(return_value_ && "no returnvalue");
-        assert(
-          (h.done() == bool(*return_value_)) &&
-          "unexpected suspension of coroutine"
-        );
-        auto r = std::move(*return_value_);
-        delete return_value_;
-        return_value_ = nullptr;
-        h.destroy(); // should I destroy?
-        return r;
-      });
-    };
-
-    std::suspend_always initial_suspend() { return {}; }
-    std::suspend_always final_suspend() noexcept { return {}; }
-    void unhandled_exception() {}
-
-    template <std::move_constructible ResumeVal> struct trivial_value_awaiter {
-      llvm::Expected<ResumeVal> val;
-      bool await_ready() noexcept { return bool(val); }
-      void await_suspend(const std::coroutine_handle<promise_type> &h) {
-        h.promise().return_value_ = new llvm::Expected<V>(val.takeError());
-      }
-      constexpr ResumeVal await_resume() const noexcept {
-        return val.get(); // this function will be not called again?
-        // thus val is not used anymore
-      }
-    };
-    template <std::move_constructible S>
-    trivial_value_awaiter<S> yield_value(ReaderIO<Env, S> &&s) {
-      assert(to_be_filled);
-      return {s.build(*to_be_filled)};
-    };
-
-    template <typename Env2, std::move_constructible S>
-    trivial_value_awaiter<S> yield_value(ReaderIO<Env2, S> &&s) {
-      assert(to_be_filled);
-      struct environment_cast<Env, Env2> cast;
-      return {s.build(cast.cast(*to_be_filled))};
-    };
-
-    void return_value(V value) {
-      return_value_ = new llvm::Expected<V>(value);
-    };
-  };
-
-  ReaderIO &concat(ReaderIO&& b) {
-    for(auto f :b.ops) {
-      ops.push_back(f);
-    }
-    b.ops = {};
-    return *this;
-  };
-
-private:
-  std::vector<BaseFunction *> ops;
-};
-
-/* bind */
-template <typename Env, typename V, std::invocable<V> Func>
-auto operator>>=(ReaderIO<Env, V> &&src, Func &&fn) {
-  using ret_value_type =
-    typename decltype(fn(std::declval<V>()).build(std::declval<Env>())
-    )::value_type;
-  using R = ReaderIO<Env, ret_value_type>;
-  return R(
-    [src0 = std::move(src), fn = std::forward<Func>(fn)](auto context
-    ) mutable -> llvm::Expected<ret_value_type> {
-      auto ret = src0.build(context);
-      if (auto err = ret.takeError()) {
-        return std::move(err);
-      }
-      return fn(ret.get()).build(context);
-    }
-  );
-}
-
-/* map */
-template <typename Env, typename V, std::invocable<V> Func>
-auto operator|(ReaderIO<Env, V> &&src, Func &&fn) {
-  using ret_value_type = decltype(fn(std::declval<V>()));
-  using R = ReaderIO<Env, ret_value_type>;
-  return R(
-    [src = std::move(src), fn = std::forward<Func>(fn)](auto context
-    ) mutable -> llvm::Expected<ret_value_type> {
-      auto ret = src.build(context);
-      if (auto err = ret.takeError()) {
-        return std::move(err);
-      }
-      return fn(ret.get());
-    }
-  );
-}
-
-/* in-place sequence: a = a then b */
-template <typename Env, typename A>
-ReaderIO<Env, A> &operator<<(ReaderIO<Env, A> &a, ReaderIO<Env, A> &&b) {
-  return a.concat(std::move(b));
-};
-
-/* in-place bind: a = fn(result of a) */
-template <typename Env, typename V, std::invocable<V> Func>
-ReaderIO<Env, V> &operator>>(ReaderIO<Env, V> &a, Func &&fn) {
-  a = std::move(a) >>= std::forward<Func>(fn);
-  return a;
-};
-
-template <typename Env, typename A, typename B, std::move_constructible Func>
-auto lift(ReaderIO<Env, A> &&a, ReaderIO<Env, B> &&b, Func &&func) {
-  return std::move(a) >>= [func = std::forward<Func>(func),
-                           b = std::move(b)](auto a) mutable {
-    return std::move(b) >>=
-           [a, func = std::forward<Func>(func)](auto b) { return func(a, b); };
-  };
-};
-
-template <typename Env, typename A, typename B, typename C, typename Func>
-auto lift(
-  ReaderIO<Env, A> &&a, ReaderIO<Env, B> &&b, ReaderIO<Env, C> &&c, Func &&func
-) {
-  return std::move(a) >>= [func = std::forward<Func>(func), b = std::move(b),
-                           c = std::move(c)](auto a) mutable {
-    return std::move(b) >>= [a, func = std::forward<Func>(func),
-                             c = std::move(c)](auto b) mutable {
-      return std::move(c) >>= [a, b, func = std::forward<Func>(func)](auto c) {
-        return func(a, b, c);
-      };
-    };
-  };
-};
-
-// template <typename Env, typename A, typename B, typename C, typename
-// D,typename Func> auto lift(const ReaderIO<Env,A>& a, const ReaderIO<Env,B>&
-// b, const ReaderIO<Env,C>& c,  const ReaderIO<Env,D>& d, Func&& func) {
-//   return a >>= [=](auto a) {
-//     return b >>= [=](auto b) {
-//       return c >>= [=](auto c) {
-//         return d >>= [=](auto d) {
-//           return func(a,b,c,d);
-//         };
-//       };
-//     };
-//   };
-// };
diff --git a/src/airconv/nt/air_builder.cpp b/src/airconv/nt/air_builder.cpp
index 44e7e9fd9..9e6b59452 100644
--- a/src/airconv/nt/air_builder.cpp
+++ b/src/airconv/nt/air_builder.cpp
@@ -1250,6 +1250,54 @@ AIRBuilder::CreateDotProduct(Value *LHS, Value *RHS) {
   return builder.CreateCall(Fn, Ops);
 }
 
+Value *
+AIRBuilder::CreateGetBuiltinID(llvm::StringRef Name, uint32_t Component) {
+  auto &Context = getContext();
+  auto Attrs = AttributeList::get(
+      Context, {{~0U, Attribute::get(Context, Attribute::AttrKind::ReadNone)},
+                {~0U, Attribute::get(Context, Attribute::AttrKind::NoUnwind)},
+                {~0U, Attribute::get(Context, Attribute::AttrKind::WillReturn)}}
+  );
+
+  std::string FnName = (Name + ".i32").str();
+  auto Fn = getModule()->getOrInsertFunction(
+      FnName, FunctionType::get(getIntTy(), {getIntTy()}, false), Attrs
+  );
+
+  return builder.CreateCall(Fn, {getInt(Component)});
+}
+
+Value *
+AIRBuilder::CreateGetGroupID(uint32_t Component) {
+  return CreateGetBuiltinID("air.get_group_id", Component);
+}
+
+Value *
+AIRBuilder::CreateGetGlobalID(uint32_t Component) {
+  return CreateGetBuiltinID("air.get_global_id", Component);
+}
+
+Value *
+AIRBuilder::CreateGetLocalID(uint32_t Component) {
+  return CreateGetBuiltinID("air.get_local_id", Component);
+}
+
+Value *
+AIRBuilder::CreateGetLocalLinearID() {
+  auto &Context = getContext();
+  auto Attrs = AttributeList::get(
+      Context, {{~0U, Attribute::get(Context, Attribute::AttrKind::ReadNone)},
+                {~0U, Attribute::get(Context, Attribute::AttrKind::NoUnwind)},
+                {~0U, Attribute::get(Context, Attribute::AttrKind::WillReturn)}}
+  );
+
+  auto Fn = getModule()->getOrInsertFunction(
+      "air.get_local_linear_id.i32", FunctionType::get(getIntTy(), {}, false), Attrs
+  );
+
+  return builder.CreateCall(Fn, {});
+}
+
 Value *
 AIRBuilder::CreateCountZero(Value *Val, bool TrailingZero) {
   auto &Context = getContext();
diff --git a/src/airconv/nt/air_builder.hpp b/src/airconv/nt/air_builder.hpp
index 509e3bfab..e552f3ac7 100644
--- a/src/airconv/nt/air_builder.hpp
+++ b/src/airconv/nt/air_builder.hpp
@@ -256,6 +256,16 @@ class AIRBuilder {
 
   /* Compute Shader */
 
+  Value *CreateGetGroupID(uint32_t Component);
+
+  Value *CreateGetGlobalID(uint32_t Component);
+
+  Value *CreateGetLocalID(uint32_t Component);
+
+  Value *CreateGetLocalLinearID();
+
+  Value *CreateGetBuiltinID(llvm::StringRef Name, uint32_t Component);
+
   CallInst *CreateBarrier(MemFlags Flags, bool SimdGroup = false);
 
   CallInst *CreateAtomicFence(MemFlags Flags, ThreadScope Scope, bool Relaxed = false);
diff --git a/src/airconv/nt/air_vertex_puller.cpp b/src/airconv/nt/air_vertex_puller.cpp
new file mode 100644
index 000000000..429c66487
--- /dev/null
+++ b/src/airconv/nt/air_vertex_puller.cpp
@@ -0,0 +1,406 @@
+/*
+ * Copyright 2026 Feifan He for CodeWeavers
+ *
+ * This library is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation; either
+ * version 2.1 of the License, or (at your option) any later version.
+ *
+ * This library is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with this library; if not, write to the Free Software
+ * Foundation, Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301, USA
+ */
+
+#include "air_vertex_puller.hpp"
+#include "../air_signature.hpp"
+#include "llvm/IR/Attributes.h"
+#include "llvm/IR/BasicBlock.h"
+#include "llvm/IR/Constants.h"
+#include "llvm/IR/DerivedTypes.h"
+#include "llvm/Support/Alignment.h"
+
+namespace dxmt::air {
+
+llvm::Value *
+VertexPuller::PullVec4(MTLVertexFormat format, llvm::Value *BaseAddress, llvm::Value *ByteOffset) {
+  using namespace llvm;
+
+  auto current = ir.GetInsertBlock();
+
+  auto pull_vertex = BasicBlock::Create(ir.getContext(), "vertex_pull_start", current->getParent());
+  auto continuation = BasicBlock::Create(ir.getContext(), "vertex_pull_done", current->getParent());
+
+  ir.CreateCondBr(ir.CreateIsNull(BaseAddress), continuation, pull_vertex);
+  ir.SetInsertPoint(pull_vertex);
+
+  auto value = PullVec4Checked(format, BaseAddress, ByteOffset);
+
+  ir.CreateBr(continuation);
+  ir.SetInsertPoint(continuation);
+
+  auto phi = ir.CreatePHI(value->getType(), 2);
+  phi->addIncoming(ConstantAggregateZero::get(value->getType()), current);
+  phi->addIncoming(value, pull_vertex);
+
+  return phi;
+}
+
+llvm::Value *
+VertexPuller::PullVec4Checked(MTLVertexFormat format, llvm::Value *BaseAddress, llvm::Value *ByteOffset) {
+  using namespace llvm;
+
+  llvm::Value *value = nullptr;
+
+  switch (format) {
+  case MTLVertexFormat::Char:
+    value = LoadFromDeviceBuffer(ir.getInt8Ty(), BaseAddress, ByteOffset, 0, 1);
+    value = ir.CreateSExt(value, ir.getInt32Ty());
+    value = ir.CreateInsertElement(PoisonValue::get(FixedVectorType::get(ir.getInt32Ty(), 4)), value, 0ull);
+    value = ir.CreateInsertElement(value, air.getInt(0), 1ull);
+    value = ir.CreateInsertElement(value, air.getInt(0), 2ull);
+    value = ir.CreateInsertElement(value, air.getInt(1), 3ull);
+    break;
+  case MTLVertexFormat::Char2:
+    value = LoadFromDeviceBuffer(FixedVectorType::get(ir.getInt8Ty(), 2), BaseAddress, ByteOffset, 0, 1);
+    value = ir.CreateSExt(value, FixedVectorType::get(ir.getInt32Ty(), 2));
+    value = ir.CreateShuffleVector(value, {0, 1, -1, -1});
+    value = ir.CreateInsertElement(value, air.getInt(0), 2ull);
+    value = ir.CreateInsertElement(value, air.getInt(1), 3ull);
+    break;
+  case MTLVertexFormat::Char4:
+    value = LoadFromDeviceBuffer(FixedVectorType::get(ir.getInt8Ty(), 4), BaseAddress, ByteOffset, 0, 1);
+    value = ir.CreateSExt(value, FixedVectorType::get(ir.getInt32Ty(), 4));
+    break;
+  case MTLVertexFormat::UChar:
+    value = LoadFromDeviceBuffer(ir.getInt8Ty(), BaseAddress, ByteOffset, 0, 1);
+    value = ir.CreateZExt(value, ir.getInt32Ty());
+    value = ir.CreateInsertElement(PoisonValue::get(FixedVectorType::get(ir.getInt32Ty(), 4)), value, 0ull);
+    value = ir.CreateInsertElement(value, air.getInt(0), 1ull);
+    value = ir.CreateInsertElement(value, air.getInt(0), 2ull);
+    value = ir.CreateInsertElement(value, air.getInt(1), 3ull);
+    break;
+  case MTLVertexFormat::UChar2:
+    value = LoadFromDeviceBuffer(FixedVectorType::get(ir.getInt8Ty(), 2), BaseAddress, ByteOffset, 0, 1);
+    value = ir.CreateZExt(value, FixedVectorType::get(ir.getInt32Ty(), 2));
+    value = ir.CreateShuffleVector(value, {0, 1, -1, -1});
+    value = ir.CreateInsertElement(value, air.getInt(0), 2ull);
+    value = ir.CreateInsertElement(value, air.getInt(1), 3ull);
+    break;
+  case MTLVertexFormat::UChar4:
+    value = LoadFromDeviceBuffer(FixedVectorType::get(ir.getInt8Ty(), 4), BaseAddress, ByteOffset, 0, 1);
+    value = ir.CreateZExt(value, FixedVectorType::get(ir.getInt32Ty(), 4));
+    break;
+  case MTLVertexFormat::Short:
+    value = LoadFromDeviceBuffer(ir.getInt16Ty(), BaseAddress, ByteOffset, 0, 2);
+    value = ir.CreateSExt(value, ir.getInt32Ty());
+    value = ir.CreateInsertElement(PoisonValue::get(FixedVectorType::get(ir.getInt32Ty(), 4)), value, 0ull);
+    value = ir.CreateInsertElement(value, air.getInt(0), 1ull);
+    value = ir.CreateInsertElement(value, air.getInt(0), 2ull);
+    value = ir.CreateInsertElement(value, air.getInt(1), 3ull);
+    break;
+  case MTLVertexFormat::Short2:
+    value = LoadFromDeviceBuffer(FixedVectorType::get(ir.getInt16Ty(), 2), BaseAddress, ByteOffset, 0, 2);
+    value = ir.CreateSExt(value, FixedVectorType::get(ir.getInt32Ty(), 2));
+    value = ir.CreateShuffleVector(value, {0, 1, -1, -1});
+    value = ir.CreateInsertElement(value, air.getInt(0), 2ull);
+    value = ir.CreateInsertElement(value, air.getInt(1), 3ull);
+    break;
+  case MTLVertexFormat::Short4:
+    value = LoadFromDeviceBuffer(FixedVectorType::get(ir.getInt16Ty(), 4), BaseAddress, ByteOffset, 0, 2);
+    value = ir.CreateSExt(value, FixedVectorType::get(ir.getInt32Ty(), 4));
+    break;
+  case MTLVertexFormat::UShort:
+    value = LoadFromDeviceBuffer(ir.getInt16Ty(), BaseAddress, ByteOffset, 0, 2);
+    value = ir.CreateZExt(value, ir.getInt32Ty());
+    value = ir.CreateInsertElement(PoisonValue::get(FixedVectorType::get(ir.getInt32Ty(), 4)), value, 0ull);
+    value = ir.CreateInsertElement(value, air.getInt(0), 1ull);
+    value = ir.CreateInsertElement(value, air.getInt(0), 2ull);
+    value = ir.CreateInsertElement(value, air.getInt(1), 3ull);
+    break;
+  case MTLVertexFormat::UShort2:
+    value = LoadFromDeviceBuffer(FixedVectorType::get(ir.getInt16Ty(), 2), BaseAddress, ByteOffset, 0, 2);
+    value = ir.CreateZExt(value, FixedVectorType::get(ir.getInt32Ty(), 2));
+    value = ir.CreateShuffleVector(value, {0, 1, -1, -1});
+    value = ir.CreateInsertElement(value, air.getInt(0), 2ull);
+    value = ir.CreateInsertElement(value, air.getInt(1), 3ull);
+    break;
+  case MTLVertexFormat::UShort4:
+    value = LoadFromDeviceBuffer(FixedVectorType::get(ir.getInt16Ty(), 4), BaseAddress, ByteOffset, 0, 2);
+    value = ir.CreateZExt(value, FixedVectorType::get(ir.getInt32Ty(), 4));
+    break;
+  case MTLVertexFormat::Half:
+    value = LoadFromDeviceBuffer(ir.getHalfTy(), BaseAddress, ByteOffset, 0, 2);
+    value = air.CreateConvertToFloat(value);
+    value = ir.CreateInsertElement(PoisonValue::get(FixedVectorType::get(ir.getFloatTy(), 4)), value, 0ull);
+    value = ir.CreateInsertElement(value, air.getFloat(0.0f), 1ull);
+    value = ir.CreateInsertElement(value, air.getFloat(0.0f), 2ull);
+    value = ir.CreateInsertElement(value, air.getFloat(1.0f), 3ull);
+    break;
+  case MTLVertexFormat::Half2:
+    value = LoadFromDeviceBuffer(FixedVectorType::get(ir.getHalfTy(), 2), BaseAddress, ByteOffset, 0, 2);
+    value = air.CreateConvertToFloat(value);
+    value = ir.CreateShuffleVector(value, {0, 1, -1, -1});
+    value = ir.CreateInsertElement(value, air.getFloat(0.0f), 2ull);
+    value = ir.CreateInsertElement(value, air.getFloat(1.0f), 3ull);
+    break;
+  case MTLVertexFormat::Half4:
+    value = LoadFromDeviceBuffer(FixedVectorType::get(ir.getHalfTy(), 4), BaseAddress, ByteOffset, 0, 2);
+    value = air.CreateConvertToFloat(value);
+    break;
+  case MTLVertexFormat::Float:
+    value = LoadFromDeviceBuffer(ir.getFloatTy(), BaseAddress, ByteOffset, 0, 4);
+    value = ir.CreateInsertElement(PoisonValue::get(FixedVectorType::get(ir.getFloatTy(), 4)), value, 0ull);
+    value = ir.CreateInsertElement(value, air.getFloat(0.0f), 1ull);
+    value = ir.CreateInsertElement(value, air.getFloat(0.0f), 2ull);
+    value = ir.CreateInsertElement(value, air.getFloat(1.0f), 3ull);
+    break;
+  case MTLVertexFormat::Float2:
+    value = LoadFromDeviceBuffer(FixedVectorType::get(ir.getFloatTy(), 2), BaseAddress, ByteOffset, 0, 4);
+    value = ir.CreateShuffleVector(value, {0, 1, -1, -1});
+    value = ir.CreateInsertElement(value, air.getFloat(0.0f), 2ull);
+    value = ir.CreateInsertElement(value, air.getFloat(1.0f), 3ull);
+    break;
+  case MTLVertexFormat::Float3:
+    value = LoadFromDeviceBuffer(FixedVectorType::get(ir.getFloatTy(), 3), BaseAddress, ByteOffset, 0, 4);
+    value = ir.CreateShuffleVector(value, {0, 1, 2, -1});
+    value = ir.CreateInsertElement(value, air.getFloat(1.0f), 3ull);
+    break;
+  case MTLVertexFormat::Float4:
+    value = LoadFromDeviceBuffer(FixedVectorType::get(ir.getFloatTy(), 4), BaseAddress, ByteOffset, 0, 4);
+    break;
+  case MTLVertexFormat::UInt:
+  case MTLVertexFormat::Int:
+    value = LoadFromDeviceBuffer(ir.getInt32Ty(), BaseAddress, ByteOffset, 0, 4);
+    value = ir.CreateInsertElement(PoisonValue::get(FixedVectorType::get(ir.getInt32Ty(), 4)), value, 0ull);
+    value = ir.CreateInsertElement(value, air.getInt(0), 1ull);
+    value = ir.CreateInsertElement(value, air.getInt(0), 2ull);
+    value = ir.CreateInsertElement(value, air.getInt(1), 3ull);
+    break;
+  case MTLVertexFormat::UInt2:
+  case MTLVertexFormat::Int2:
+    value = LoadFromDeviceBuffer(FixedVectorType::get(ir.getInt32Ty(), 2), BaseAddress, ByteOffset, 0, 4);
+    value = ir.CreateShuffleVector(value, {0, 1, -1, -1});
+    value = ir.CreateInsertElement(value, air.getInt(0), 2ull);
+    value = ir.CreateInsertElement(value, air.getInt(1), 3ull);
+    break;
+  case MTLVertexFormat::UInt3:
+  case MTLVertexFormat::Int3:
+    value = LoadFromDeviceBuffer(FixedVectorType::get(ir.getInt32Ty(), 3), BaseAddress, ByteOffset, 0, 4);
+    value = ir.CreateShuffleVector(value, {0, 1, 2, -1});
+    value = ir.CreateInsertElement(value, air.getInt(1), 3ull);
+    break;
+  case MTLVertexFormat::UInt4:
+  case MTLVertexFormat::Int4:
+    value = LoadFromDeviceBuffer(FixedVectorType::get(ir.getInt32Ty(), 4), BaseAddress, ByteOffset, 0, 4);
+    break;
+  case MTLVertexFormat::UCharNormalized:
+  case MTLVertexFormat::UChar2Normalized:
+  case MTLVertexFormat::UChar4Normalized:
+  case MTLVertexFormat::UChar4Normalized_BGRA:
+  case MTLVertexFormat::CharNormalized:
+  case MTLVertexFormat::Char2Normalized:
+  case MTLVertexFormat::Char4Normalized:
+  case MTLVertexFormat::UShortNormalized:
+  case MTLVertexFormat::UShort2Normalized:
+  case MTLVertexFormat::UShort4Normalized:
+  case MTLVertexFormat::ShortNormalized:
+  case MTLVertexFormat::Short2Normalized:
+  case MTLVertexFormat::Short4Normalized:
+  case MTLVertexFormat::FloatRG11B10:
+  case MTLVertexFormat::FloatRGB9E5:
+  case MTLVertexFormat::Int1010102Normalized:
+  case MTLVertexFormat::UInt1010102Normalized:
+    value = UnpackFVec4(format, BaseAddress, ByteOffset);
+    break;
+  case MTLVertexFormat::Invalid:
+    break;
+  }
+  assert(value);
+
+  return value;
+}
+
+llvm::Value *
+VertexPuller::UnpackFVec4(MTLVertexFormat format, llvm::Value *BaseAddress, llvm::Value *AlignedOffset) {
+  using namespace llvm;
+
+  if (format == MTLVertexFormat::UShort4Normalized) {
+    llvm::Value *_xy = CallUnpackImpl(
+        "unorm2x16.v2f32", LoadFromDeviceBuffer(ir.getInt32Ty(), BaseAddress, AlignedOffset, 0, 2), ir.getInt32Ty(),
+        FixedVectorType::get(ir.getFloatTy(), 2)
+    );
+    llvm::Value *_zw = CallUnpackImpl(
+        "unorm2x16.v2f32", LoadFromDeviceBuffer(ir.getInt32Ty(), BaseAddress, AlignedOffset, 1, 2), ir.getInt32Ty(),
+        FixedVectorType::get(ir.getFloatTy(), 2)
+    );
+    return ir.CreateShuffleVector(_xy, _zw, {0, 1, 2, 3});
+  }
+  if (format == MTLVertexFormat::Short4Normalized) {
+    llvm::Value *_xy = CallUnpackImpl(
+        "snorm2x16.v2f32", LoadFromDeviceBuffer(ir.getInt32Ty(), BaseAddress, AlignedOffset, 0, 2), ir.getInt32Ty(),
+        FixedVectorType::get(ir.getFloatTy(), 2)
+    );
+    llvm::Value *_zw = CallUnpackImpl(
+        "snorm2x16.v2f32", LoadFromDeviceBuffer(ir.getInt32Ty(), BaseAddress, AlignedOffset, 1, 2), ir.getInt32Ty(),
+        FixedVectorType::get(ir.getFloatTy(), 2)
+    );
+    return ir.CreateShuffleVector(_xy, _zw, {0, 1, 2, 3});
+  }
+
+  std::string op;
+  llvm::Type *src_type = nullptr;
+  uint32_t align = 0;
+  llvm::Type *dst_type = nullptr;
+  switch (format) {
+  case MTLVertexFormat::UCharNormalized:
+    op = "unorm1x8.f32";
+    src_type = ir.getInt8Ty();
+    align = 1;
+    dst_type = ir.getFloatTy();
+    break;
+  case MTLVertexFormat::UChar2Normalized:
+    op = "unorm2x8.v2f32";
+    src_type = ir.getInt16Ty();
+    align = 1;
+    dst_type = FixedVectorType::get(ir.getFloatTy(), 2);
+    break;
+  case MTLVertexFormat::UChar4Normalized:
+  case MTLVertexFormat::UChar4Normalized_BGRA:
+    op = "unorm4x8.v4f32";
+    src_type = ir.getInt32Ty();
+    align = 1;
+    dst_type = FixedVectorType::get(ir.getFloatTy(), 4);
+    break;
+  case MTLVertexFormat::CharNormalized:
+    op = "snorm1x8.f32";
+    src_type = ir.getInt8Ty();
+    align = 1;
+    dst_type = ir.getFloatTy();
+    break;
+  case MTLVertexFormat::Char2Normalized:
+    op = "snorm2x8.v2f32";
+    src_type = ir.getInt16Ty();
+    align = 1;
+    dst_type = FixedVectorType::get(ir.getFloatTy(), 2);
+    break;
+  case MTLVertexFormat::Char4Normalized:
+    op = "snorm4x8.v4f32";
+    src_type = ir.getInt32Ty();
+    align = 1;
+    dst_type = FixedVectorType::get(ir.getFloatTy(), 4);
+    break;
+  case MTLVertexFormat::UShortNormalized:
+    op = "unorm1x16.f32";
+    src_type = ir.getInt16Ty();
+    align = 2;
+    dst_type = ir.getFloatTy();
+    break;
+  case MTLVertexFormat::UShort2Normalized:
+    op = "unorm2x16.v2f32";
+    src_type = ir.getInt32Ty();
+    align = 2;
+    dst_type = FixedVectorType::get(ir.getFloatTy(), 2);
+    break;
+  case MTLVertexFormat::ShortNormalized:
+    op = "snorm1x16.f32";
+    src_type = ir.getInt16Ty();
+    align = 2;
+    dst_type = ir.getFloatTy();
+    break;
+  case MTLVertexFormat::Short2Normalized:
+    op = "snorm2x16.v2f32";
+    src_type = ir.getInt32Ty();
+    align = 2;
+    dst_type = FixedVectorType::get(ir.getFloatTy(), 2);
+    break;
+  case MTLVertexFormat::UInt1010102Normalized:
+    op = "unorm.rgb10a2.v4f32";
+    src_type = ir.getInt32Ty();
+    align = 4;
+    dst_type = FixedVectorType::get(ir.getFloatTy(), 4);
+    break;
+  case MTLVertexFormat::Int1010102Normalized:
+    assert(0 && "unused & Metal 4 only");
+    op = "snorm.rgb10a2.v4f32";
+    src_type = ir.getInt32Ty();
+    align = 4;
+    dst_type = FixedVectorType::get(ir.getFloatTy(), 4);
+    break;
+  case MTLVertexFormat::FloatRG11B10:
+    op = "unorm.rg11b10f.v3f32";
+    src_type = ir.getInt32Ty();
+    align = 4;
+    dst_type = FixedVectorType::get(ir.getFloatTy(), 3);
+    break;
+  case MTLVertexFormat::FloatRGB9E5:
+    op = "unorm.rgb9e5.v3f32";
+    src_type = ir.getInt32Ty();
+    align = 4;
+    dst_type = FixedVectorType::get(ir.getFloatTy(), 3);
+    break;
+  default:
+    break;
+  }
+  assert(src_type && dst_type);
+  llvm::Value *ret =
+      CallUnpackImpl(op, LoadFromDeviceBuffer(src_type, BaseAddress, AlignedOffset, 0, align), src_type, dst_type);
+  if (format == MTLVertexFormat::UChar4Normalized_BGRA) {
+    ret = ir.CreateShuffleVector(ret, {2, 1, 0, 3});
+  }
+  if (dst_type == ir.getFloatTy()) {
+    ret = ir.CreateInsertElement(PoisonValue::get(FixedVectorType::get(ir.getFloatTy(), 4)), ret, 0ull);
+    ret = ir.CreateInsertElement(ret, air.getFloat(0.0f), 1ull);
+    ret = ir.CreateInsertElement(ret, air.getFloat(0.0f), 2ull);
+    ret = ir.CreateInsertElement(ret, air.getFloat(1.0f), 3ull);
+  } else if (dst_type == FixedVectorType::get(ir.getFloatTy(), 2)) {
+    ret = ir.CreateShuffleVector(ret, {0, 1, -1, -1});
+    ret = ir.CreateInsertElement(ret, air.getFloat(0.0f), 2ull);
+    ret = ir.CreateInsertElement(ret, air.getFloat(1.0f), 3ull);
+  } else if (dst_type == FixedVectorType::get(ir.getFloatTy(), 3)) {
+    ret = ir.CreateShuffleVector(ret, {0, 1, 2, -1});
+    ret = ir.CreateInsertElement(ret, air.getFloat(1.0f), 3ull);
+  }
+  return ret;
+}
+
+llvm::Value *
+VertexPuller::LoadFromDeviceBuffer(
+    llvm::Type *value_type, llvm::Value *BaseAddress, llvm::Value *AlignedOffset, uint32_t ElementOffset, uint32_t align
+) {
+  using namespace llvm;
+  return ir.CreateAlignedLoad(
+      value_type,
+      ir.CreateConstGEP1_32(
+          value_type,
+          ir.CreateBitCast(
+              ir.CreateGEP(ir.getInt8Ty(), BaseAddress, {AlignedOffset}),
+              value_type->getPointerTo((uint32_t)air::AddressSpace::device)
+          ),
+          ElementOffset
+      ),
+      llvm::MaybeAlign(align)
+  );
+}
+
+llvm::Value *
+VertexPuller::CallUnpackImpl(std::string Op, llvm::Value *Src, llvm::Type *TySrc, llvm::Type *TyDst) {
+  using namespace llvm;
+  auto &context = ir.getContext();
+  auto att = AttributeList::get(
+      context, {{~0U, Attribute::get(context, Attribute::AttrKind::NoUnwind)},
+                {~0U, Attribute::get(context, Attribute::AttrKind::WillReturn)},
+                {~0U, Attribute::get(context, Attribute::AttrKind::MustProgress)},
+                {~0U, Attribute::get(context, Attribute::AttrKind::NoFree)},
+                {~0U, Attribute::get(context, Attribute::AttrKind::NoSync)},
+                {~0U, Attribute::get(context, Attribute::AttrKind::ReadNone)}}
+  );
+  auto fn = module.getOrInsertFunction("air.unpack." + Op, FunctionType::get(TyDst, {TySrc}, false), att);
+  return ir.CreateCall(fn, {Src});
+}
+
+} // namespace dxmt::air
diff --git a/src/airconv/nt/air_vertex_puller.hpp b/src/airconv/nt/air_vertex_puller.hpp
new file mode 100644
index 000000000..e65195f2c
--- /dev/null
+++ b/src/airconv/nt/air_vertex_puller.hpp
@@ -0,0 +1,110 @@
+/*
+ * Copyright 2026 Feifan He for CodeWeavers
+ *
+ * This library is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation; either
+ * version 2.1 of the License, or (at your option) any later version.
+ *
+ * This library is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with this library; if not, write to the Free Software
+ * Foundation, Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301, USA
+ */
+
+#pragma once
+
+#include "air_builder.hpp"
+#include "llvm/IR/IRBuilder.h"
+#include "llvm/IR/Module.h"
+#include "llvm/IR/Value.h"
+#include <string>
+
+namespace dxmt::air {
+
+/** a.k.a. MTLAttributeFormat */
+enum class MTLVertexFormat {
+  Invalid = 0,
+  UChar2 = 1,
+  // UChar3 = 2,
+  UChar4 = 3,
+  Char2 = 4,
+  // Char3 = 5,
+  Char4 = 6,
+  UChar2Normalized = 7,
+  // UChar3Normalized = 8,
+  UChar4Normalized = 9,
+  Char2Normalized = 10,
+  // Char3Normalized = 11,
+  Char4Normalized = 12,
+  UShort2 = 13,
+  // UShort3 = 14,
+  UShort4 = 15,
+  Short2 = 16,
+  // Short3 = 17,
+  Short4 = 18,
+  UShort2Normalized = 19,
+  // UShort3Normalized = 20,
+  UShort4Normalized = 21,
+  Short2Normalized = 22,
+  // Short3Normalized = 23,
+  Short4Normalized = 24,
+  Half2 = 25,
+  // Half3 = 26,
+  Half4 = 27,
+  Float = 28,
+  Float2 = 29,
+  Float3 = 30,
+  Float4 = 31,
+  Int = 32,
+  Int2 = 33,
+  Int3 = 34,
+  Int4 = 35,
+  UInt = 36,
+  UInt2 = 37,
+  UInt3 = 38,
+  UInt4 = 39,
+  Int1010102Normalized = 40,
+  UInt1010102Normalized = 41,
+  UChar4Normalized_BGRA = 42,
+  UChar = 45,
+  Char = 46,
+  UCharNormalized = 47,
+  CharNormalized = 48,
+  UShort = 49,
+  Short = 50,
+  UShortNormalized = 51,
+  ShortNormalized = 52,
+  Half = 53,
+  FloatRG11B10 = 54,
+  FloatRGB9E5 = 55,
+};
+
+class VertexPuller {
+public:
+  VertexPuller(llvm::air::AIRBuilder &air, llvm::Module &module) : air(air), ir(air.builder), module(module) {}
+
+  llvm::Value *PullVec4(MTLVertexFormat Format, llvm::Value *BaseAddress, llvm::Value *ByteOffset);
+
+private:
+  llvm::Value *PullVec4Checked(MTLVertexFormat Format, llvm::Value *BaseAddress, llvm::Value *ByteOffset);
+
+  llvm::Value *UnpackFVec4(MTLVertexFormat Format, llvm::Value *BaseAddress, llvm::Value *AlignedOffset);
+
+  llvm::Value *LoadFromDeviceBuffer(
+      llvm::Type *value_type, llvm::Value *BaseAddress, llvm::Value *AlignedOffset, uint32_t ElementOffset,
+      uint32_t align
+  );
+
+  llvm::Value *CallUnpackImpl(std::string Op, llvm::Value *Src, llvm::Type *TySrc, llvm::Type *TyDst);
+
+  llvm::air::AIRBuilder &air;
+  llvm::IRBuilderBase &ir;
+  llvm::Module &module;
+};
+
+} // namespace dxmt::air
diff --git a/src/airconv/nt/dxbc_converter_base.cpp b/src/airconv/nt/dxbc_converter_base.cpp
index 58a101c01..353be5b8e 100644
--- a/src/airconv/nt/dxbc_converter_base.cpp
+++ b/src/airconv/nt/dxbc_converter_base.cpp
@@ -2563,11 +2563,9 @@ Converter::LoadInterpolant(uint32_t Index) {
     return {};
 
   auto interpolant = res.interpolant_map[Index];
-  auto h = interpolant.interpolant(nullptr).build(ctx);
-  if (h.takeError())
-    return {};
+  auto h = ctx.function->getArg(interpolant.arg_index);
 
-  return llvm::Optional<InterpolantHandle>({h.get(), interpolant.perspective});
+  return llvm::Optional<InterpolantHandle>({h, interpolant.perspective});
 }
 
 void
@@ -2639,18 +2637,14 @@ Converter::operator()(const InstMaskedSumOfAbsDiff &msad) {
 
   StoreOperand(msad.dst, Result);
 }
-
 void
 Converter::operator()(const InstEmit &) {
-  if (res.call_emit().build(ctx).takeError()) {
-    // TODO
-  }
+  res.call_emit();
 }
+
 void
 Converter::operator()(const InstCut &) {
-  if (res.call_cut().build(ctx).takeError()) {
-    // TODO
-  }
+  res.call_cut();
 }
 
 llvm::Value *
diff --git a/src/airconv/nt/dxbc_io_helper.cpp b/src/airconv/nt/dxbc_io_helper.cpp
new file mode 100644
index 000000000..a0f501dd7
--- /dev/null
+++ b/src/airconv/nt/dxbc_io_helper.cpp
@@ -0,0 +1,417 @@
+#include "dxbc_io_helper.hpp"
+#include "../airconv_error.hpp"
+#include "../dxbc_converter.hpp"
+#include "air_vertex_puller.hpp"
+
+#include "llvm/ADT/SmallVector.h"
+#include "llvm/IR/Constants.h"
+
+namespace dxmt::dxbc {
+
+IOHelper::IOHelper(
+    llvm::air::AIRBuilder &air, llvm::Function *function, air::AirType &types, io_binding_map &io_legacy
+) :
+    air(air),
+    ir(air.builder),
+    function(function),
+    types(types),
+    io(io_legacy) {}
+
+llvm::Value *
+IOHelper::GetFloat4Splat(float value) {
+  return llvm::ConstantVector::get(
+      {llvm::ConstantFP::get(air.getContext(), llvm::APFloat{value}),
+       llvm::ConstantFP::get(air.getContext(), llvm::APFloat{value}),
+       llvm::ConstantFP::get(air.getContext(), llvm::APFloat{value}),
+       llvm::ConstantFP::get(air.getContext(), llvm::APFloat{value})}
+  );
+};
+
+llvm::Value *
+IOHelper::ExtendToVec4(llvm::Value *value) {
+  auto ty = value->getType();
+  if (ty->isVectorTy()) {
+    auto vecTy = llvm::cast<llvm::FixedVectorType>(ty);
+    if (vecTy->getNumElements() == 4)
+      return value;
+    if (vecTy->getNumElements() == 3)
+      return ir.CreateShuffleVector(value, {0, 1, 2, 2});
+    if (vecTy->getNumElements() == 2)
+      return ir.CreateShuffleVector(value, {0, 1, 1, 1});
+    if (vecTy->getNumElements() == 1)
+      return ir.CreateShuffleVector(value, {0, 0, 0, 0});
+    return ir.CreateShuffleVector(value, {0, 1, 2, 3});
+  } else {
+    return ir.CreateVectorSplat(4, value);
+  }
+}
+
+llvm::Value *
+IOHelper::ToDesiredTypeFromIntVec4(llvm::Value *vec4, llvm::Type *desired, uint32_t mask) {
+  llvm::Type *scalar = desired->getScalarType();
+  assert((scalar == types._int || scalar == types._float) && "unhandled desired element type");
+
+  // reinterpret the integer register as float when a float type is requested
+  llvm::Value *value = scalar == types._float ? ir.CreateBitCast(vec4, types._float4) : vec4;
+
+  if (auto *vector = llvm::dyn_cast<llvm::FixedVectorType>(desired)) {
+    unsigned num_elements = vector->getNumElements();
+    llvm::SmallVector<int, 4> indices;
+    for (unsigned i = 0; i < num_elements; i++)
+      indices.push_back(mask & (1 << i) ? (int)i : llvm::UndefMaskElem);
+    return ir.CreateShuffleVector(value, indices);
+  }
+
+  assert(mask);
+  return ir.CreateExtractElement(value, mask ? (uint64_t)__builtin_ctz(mask) : 0);
+};
+
+llvm::Value *
+IOHelper::LoadFromArrayAt(llvm::Value *array, llvm::Value *index) {
+  auto array_ty = llvm::cast<llvm::ArrayType>( // force line break
+      llvm::cast<llvm::PointerType>(array->getType())->getNonOpaquePointerElementType()
+  );
+  auto ptr = ir.CreateGEP(array_ty, array, {air.getInt(0), index}, "", llvm::isa<llvm::ConstantInt>(index));
+  return ir.CreateLoad(array_ty->getElementType(), ptr);
+};
+
+llvm::Value *
+IOHelper::LoadFromVec4ArrayMasked(llvm::Value *array, llvm::Value *index, uint32_t mask) {
+  if (mask == 0b1111) {
+    return LoadFromArrayAt(array, index);
+  }
+  auto array_type = llvm::cast<llvm::PointerType>(array->getType())->getNonOpaquePointerElementType();
+  auto vec4_type = llvm::cast<llvm::ArrayType>(array_type)->getArrayElementType();
+  auto ele_type = llvm::cast<llvm::VectorType>(vec4_type)->getElementType();
+  llvm::Value *value = llvm::ConstantAggregateZero::get(vec4_type);
+  for (unsigned i = 0; i < 4; i++) {
+    if ((mask & (1 << i)) == 0)
+      continue;
+    auto component_ptr = ir.CreateGEP(array_type, array, {air.getInt(0), index, air.getInt(i)});
+    value = ir.CreateInsertElement(value, ir.CreateLoad(ele_type, component_ptr), uint64_t(i));
+  }
+  return value;
+};
+
+llvm::Error
+IOHelper::StoreAtVec4ArrayMasked(llvm::Value *array, llvm::Value *index, llvm::Value *maybe_vec4, uint32_t mask) {
+  auto vec4 = ExtendToVec4(maybe_vec4);
+  if (mask == 0b1111) {
+    auto ptr = ir.CreateInBoundsGEP(
+        llvm::cast<llvm::PointerType>(array->getType())->getNonOpaquePointerElementType(), array, {air.getInt(0), index}
+    );
+    ir.CreateStore(vec4, ptr);
+    return llvm::Error::success();
+  }
+  for (unsigned i = 0; i < 4; i++) {
+    if ((mask & (1 << i)) == 0)
+      continue;
+    auto component_ptr = ir.CreateGEP(
+        llvm::cast<llvm::PointerType>(array->getType())->getNonOpaquePointerElementType(), array,
+        {air.getInt(0), index, air.getInt(i)}
+    );
+    ir.CreateStore(ir.CreateExtractElement(vec4, i), component_ptr);
+  }
+  return llvm::Error::success();
+};
+
+llvm::Error
+IOHelper::InitializeInput(uint32_t with_fnarg_at, uint32_t to_reg, uint32_t mask, bool fix_w_component) {
+  llvm::Value *arg = function->getArg(with_fnarg_at);
+  if (arg->getType()->getScalarType()->isIntegerTy()) {
+    if (arg->getType()->getScalarType() == types._bool) {
+      // front_facing
+      return StoreAtVec4ArrayMasked(io.input.ptr_int4, air.getInt(to_reg), ir.CreateSExt(arg, types._int), mask);
+    } else if (arg->getType()->getScalarType() != types._int) {
+      return llvm::make_error<UnsupportedFeature>("TODO: handle non 32 bit integer");
+    } else {
+      return StoreAtVec4ArrayMasked(io.input.ptr_int4, air.getInt(to_reg), arg, mask);
+    }
+  } else if (arg->getType()->getScalarType()->isFloatTy()) {
+    if (fix_w_component && isa<llvm::FixedVectorType>(arg->getType()) &&
+        cast<llvm::FixedVectorType>(arg->getType())->getNumElements() == 4) {
+      auto w_component = ir.CreateExtractElement(arg, 3);
+      auto rcp_w = ir.CreateFDiv(air.getFloat(1.0f), w_component);
+      arg = ir.CreateInsertElement(arg, rcp_w, 3);
+    }
+    return StoreAtVec4ArrayMasked(io.input.ptr_float4, air.getInt(to_reg), arg, mask);
+  } else {
+    return llvm::make_error<UnsupportedFeature>("TODO: unhandled input type? might be interpolant...");
+  }
+  // unreachable
+  return llvm::Error::success();
+};
+
+llvm::Error
+IOHelper::InitializeInterpolatedInput(
+    uint32_t with_fnarg_at, uint32_t to_reg, uint32_t mask, air::Interpolation interpolation, uint32_t sampleidx_at
+) {
+  llvm::Value *interpolant = function->getArg(with_fnarg_at);
+  llvm::Value *interpolated_val = nullptr;
+  switch (interpolation) {
+  case air::Interpolation::center_perspective:
+    interpolated_val = air.CreateInterpolateAtCenter(interpolant, true);
+    break;
+  case air::Interpolation::center_no_perspective:
+    interpolated_val = air.CreateInterpolateAtCenter(interpolant, false);
+    break;
+  case air::Interpolation::centroid_perspective:
+    interpolated_val = air.CreateInterpolateAtCentroid(interpolant, true);
+    break;
+  case air::Interpolation::centroid_no_perspective:
+    interpolated_val = air.CreateInterpolateAtCentroid(interpolant, false);
+    break;
+  case air::Interpolation::sample_perspective:
+    assert(~sampleidx_at);
+    interpolated_val = air.CreateInterpolateAtSample(interpolant, function->getArg(sampleidx_at), true);
+    break;
+  case air::Interpolation::sample_no_perspective:
+    assert(~sampleidx_at);
+    interpolated_val = air.CreateInterpolateAtSample(interpolant, function->getArg(sampleidx_at), false);
+    break;
+  case air::Interpolation::flat:
+    return llvm::make_error<UnsupportedFeature>("unexpected interpolant mode: flat");
+  }
+  return StoreAtVec4ArrayMasked(io.input.ptr_float4, air.getInt(to_reg), interpolated_val, mask);
+}
+
+llvm::Error
+IOHelper::PullVertexInput(
+    uint32_t vbuf_table_arg_index, uint32_t to_reg_unused, uint32_t mask, SM50_IA_INPUT_ELEMENT element_info,
+    uint32_t slot_mask
+) {
+  auto &builder = ir;
+  llvm::Value *index;             // uint
+  if (element_info.step_function) // per_instance
+  {
+    if (element_info.step_rate) {
+      index = builder.CreateAdd(
+          io.base_instance_id, builder.CreateUDiv(io.instance_id, builder.getInt32(element_info.step_rate))
+      );
+    } else {
+      // 0 takes special meaning, that the instance data should never be
+      // stepped at all
+      index = io.base_instance_id;
+    }
+  } else {
+    index = io.vertex_id_with_base;
+  }
+  if (!io.vertex_buffer_table) {
+    io.vertex_buffer_table = builder.CreateBitCast(
+        function->getArg(vbuf_table_arg_index),
+        types._dxmt_vertex_buffer_entry->getPointerTo((uint32_t)air::AddressSpace::constant)
+    );
+  }
+  unsigned int shift = 32u - element_info.slot;
+  unsigned int vertex_buffer_entry_index = element_info.slot ? __builtin_popcount((slot_mask << shift) >> shift) : 0;
+  auto vertex_buffer_table = io.vertex_buffer_table;
+  auto vertex_buffer_entry = builder.CreateLoad(
+      types._dxmt_vertex_buffer_entry,
+      builder.CreateConstGEP1_32(types._dxmt_vertex_buffer_entry, vertex_buffer_table, vertex_buffer_entry_index)
+  );
+  auto base_addr = builder.CreateExtractValue(vertex_buffer_entry, {0});
+  auto stride = builder.CreateExtractValue(vertex_buffer_entry, {1});
+  auto byte_offset =
+      builder.CreateAdd(builder.CreateMul(stride, index), builder.getInt32(element_info.aligned_byte_offset));
+  air::VertexPuller puller(air, *function->getParent());
+  auto vec4 = puller.PullVec4((air::MTLVertexFormat)element_info.format, base_addr, byte_offset);
+  if (vec4->getType() == types._float4) {
+    return StoreAtVec4ArrayMasked(io.input.ptr_float4, builder.getInt32(element_info.reg), vec4, mask);
+  } else if (vec4->getType() == types._int4) {
+    return StoreAtVec4ArrayMasked(io.input.ptr_int4, builder.getInt32(element_info.reg), vec4, mask);
+  } else {
+    return llvm::make_error<UnsupportedFeature>("unexpected value type from vertex puller");
+  }
+};
+
+llvm::Error
+IOHelper::InitializeVertexIDInput(uint32_t to_reg, uint32_t mask) {
+  return StoreAtVec4ArrayMasked(io.input.ptr_int4, air.getInt(to_reg), io.vertex_id, mask);
+}
+
+llvm::Error
+IOHelper::InitializeInstanceIDInput(uint32_t to_reg, uint32_t mask) {
+  return StoreAtVec4ArrayMasked(io.input.ptr_int4, air.getInt(to_reg), io.instance_id, mask);
+}
+
+llvm::Error
+IOHelper::InitializeInputAttribute(shader::common::InputAttribute value, uint32_t from_fnarg_at) {
+  auto arg = function->getArg(from_fnarg_at);
+  switch (value) {
+  case shader::common::InputAttribute::ThreadId:
+    io.thread_id_arg = arg;
+    break;
+  case shader::common::InputAttribute::ThreadGroupId:
+    io.thread_group_id_arg = arg;
+    break;
+  case shader::common::InputAttribute::ThreadIdInGroup:
+    io.thread_id_in_group_arg = arg;
+    break;
+  case shader::common::InputAttribute::ThreadIdInGroupFlatten:
+    io.thread_id_in_group_flat_arg = arg;
+    break;
+  case shader::common::InputAttribute::CoverageMask:
+    io.coverage_mask_arg = arg;
+    break;
+  default:
+    return llvm::make_error<UnsupportedFeature>("unhandled input system value");
+  }
+  return llvm::Error::success();
+}
+
+llvm::Error
+IOHelper::CreateDepthOutputRegister() {
+  if (io.depth_output_reg != nullptr)
+    return llvm::make_error<UnsupportedFeature>("oDepth is defined twice");
+  io.depth_output_reg = ir.CreateAlloca(types._float);
+  return llvm::Error::success();
+}
+
+llvm::Error
+IOHelper::CreateStencilRefOutputRegister() {
+  if (io.stencil_ref_reg != nullptr)
+    return llvm::make_error<UnsupportedFeature>("oStencil is defined twice");
+  io.stencil_ref_reg = ir.CreateAlloca(types._int);
+  return llvm::Error::success();
+}
+
+llvm::Error
+IOHelper::CreateCoverageMaskOutputRegister() {
+  if (io.coverage_mask_reg != nullptr)
+    return llvm::make_error<UnsupportedFeature>("oMask is defined twice");
+  io.coverage_mask_reg = ir.CreateAlloca(types._int);
+  return llvm::Error::success();
+}
+
+llvm::Value *
+IOHelper::PopulateOutput(uint32_t from_reg, uint32_t mask, llvm::Value *StructRet, uint32_t to_element) {
+  auto desired_type = function->getReturnType()->getStructElementType(to_element);
+  auto value = ToDesiredTypeFromIntVec4(LoadFromArrayAt(io.output.ptr_int4, air.getInt(from_reg)), desired_type, mask);
+  return ir.CreateInsertValue(StructRet, value, {to_element});
+}
+
+llvm::Value *
+IOHelper::PopulateOutputFixUnorm(uint32_t from_reg, uint32_t mask, llvm::Value *StructRet, uint32_t to_element) {
+  auto fvec4 = LoadFromArrayAt(io.output.ptr_float4, air.getInt(from_reg));
+  auto fixed = ir.CreateFSub(fvec4, GetFloat4Splat(1.0f / 127500.0f) /* magic delta ?! */);
+  return ir.CreateInsertValue(StructRet, fixed, {to_element});
+}
+
+llvm::Value *
+IOHelper::PopulateOutputSanitizePosition(
+    uint32_t from_reg, uint32_t mask, llvm::Value *StructRet, uint32_t to_element
+) {
+  auto fvec4 = LoadFromArrayAt(io.output.ptr_float4, air.getInt(from_reg));
+  auto fixed = air.SanitizePosition(fvec4);
+  return ir.CreateInsertValue(StructRet, fixed, {to_element});
+}
+
+llvm::Value *
+IOHelper::PopulateOutputDepth(llvm::Value *StructRet, uint32_t to_element) {
+  auto value = ir.CreateLoad(types._float, ir.CreateConstInBoundsGEP1_32(types._float, io.depth_output_reg, 0));
+  return ir.CreateInsertValue(StructRet, value, {to_element});
+}
+
+llvm::Value *
+IOHelper::PopulateOutputStencilRef(llvm::Value *StructRet, uint32_t to_element) {
+  auto value = ir.CreateLoad(types._int, ir.CreateConstInBoundsGEP1_32(types._int, io.stencil_ref_reg, 0));
+  return ir.CreateInsertValue(StructRet, value, {to_element});
+}
+
+llvm::Value *
+IOHelper::PopulateOutputCoverageMask(llvm::Value *StructRet, uint32_t to_element, uint32_t sample_mask) {
+  llvm::Value *value = ir.CreateLoad(types._int, ir.CreateConstInBoundsGEP1_32(types._int, io.coverage_mask_reg, 0));
+  if (sample_mask != 0xffffffff)
+    value = ir.CreateAnd(value, sample_mask);
+  return ir.CreateInsertValue(StructRet, value, {to_element});
+}
+
+llvm::Value *
+IOHelper::PopulateDefaultCoverageMask(llvm::Value *StructRet, uint32_t to_element, uint32_t sample_mask) {
+  if (io.coverage_mask_reg)
+    return StructRet;
+  return ir.CreateInsertValue(StructRet, air.getInt(sample_mask), {to_element});
+}
+
+llvm::Value *
+IOHelper::PopulateOutputStreamOutput(
+    llvm::Value *StructRet, const SM50_SHADER_EMULATE_VERTEX_STREAM_OUTPUT_DATA *vertex_so,
+    uint32_t base_vertex_arg_index, uint32_t vertex_id_arg_index, uint32_t stream_output_table_arg_index
+) {
+  auto base_vertex = function->getArg(base_vertex_arg_index);
+  auto vertex_id = function->getArg(vertex_id_arg_index);
+  auto so_entries_type = types._dxmt_stream_output_buffer_entry;
+  auto so_entries = ir.CreateBitCast(
+      function->getArg(stream_output_table_arg_index),
+      so_entries_type->getPointerTo((uint32_t)air::AddressSpace::constant)
+  );
+  auto slot0_entry = ir.CreateLoad(so_entries_type, ir.CreateConstInBoundsGEP1_32(so_entries_type, so_entries, 0));
+  auto slot0 = ir.CreateExtractValue(slot0_entry, {0});
+  auto adjusted_vertex_id = ir.CreateSub(vertex_id, base_vertex);
+  auto output_regs = ir.CreateBitOrPointerCast(io.output.ptr_int4, llvm::PointerType::get(types._int, 0));
+  for (unsigned i = 0; i < vertex_so->num_elements; i++) {
+    auto &element = vertex_so->elements[i];
+    if (element.reg_id == 0xffffffff)
+      continue;
+    auto ptr = ir.CreateConstGEP1_32(types._int, output_regs, (unsigned)(element.reg_id * 4 + element.component));
+    auto target_offset = ir.CreateAdd(
+        ir.CreateMul(adjusted_vertex_id, air.getInt(vertex_so->strides[element.output_slot /* expected to be 0 */])),
+        air.getInt(element.offset)
+    );
+    auto target_ptr = ir.CreateGEP(types._int, slot0, {ir.CreateLShr(target_offset, 2)});
+    ir.CreateStore(ir.CreateLoad(types._int, ptr), target_ptr, true);
+  }
+  return StructRet;
+}
+
+void
+IOHelper::PopulateMeshOutputRenderTargetArrayIndex(uint32_t from_reg, uint32_t mask, llvm::Value *primitive_id) {
+  auto result = ToDesiredTypeFromIntVec4(LoadFromArrayAt(io.output.ptr_int4, air.getInt(from_reg)), types._int, mask);
+  air.CreateSetMeshRenderTargetArrayIndex(primitive_id, result);
+}
+
+void
+IOHelper::PopulateMeshOutputViewportArrayIndex(uint32_t from_reg, uint32_t mask, llvm::Value *primitive_id) {
+  auto result = ToDesiredTypeFromIntVec4(LoadFromArrayAt(io.output.ptr_int4, air.getInt(from_reg)), types._int, mask);
+  air.CreateSetMeshViewportArrayIndex(primitive_id, result);
+}
+
+void
+IOHelper::PopulateMeshOutputPosition(uint32_t from_reg, uint32_t mask, llvm::Value *vertex_id) {
+  auto result = LoadFromArrayAt(io.output.ptr_float4, air.getInt(from_reg));
+  air.CreateSetMeshPosition(vertex_id, result);
+}
+
+void
+IOHelper::PopulateMeshOutputVertexData(
+    uint32_t from_reg, uint32_t mask, uint32_t idx, llvm::Value *vertex_id, air::MSLScalerOrVectorType desired_type
+) {
+  auto result = ToDesiredTypeFromIntVec4(
+      LoadFromVec4ArrayMasked(io.output.ptr_int4, air.getInt(from_reg), mask),
+      air::get_llvm_type(desired_type, air.getContext()), mask
+  );
+  air.CreateSetMeshVertexData(vertex_id, idx, result);
+}
+
+llvm::Error
+Prologues::Run(IOHelper &IO) {
+  for (auto &fn : functions) {
+    if (auto err = fn(IO))
+      return err;
+  }
+  return llvm::Error::success();
+}
+
+llvm::Expected<llvm::Value *>
+Epilogues::Run(llvm::Type *TyRet, IOHelper &IO) {
+  llvm::Value *ret = TyRet->isVoidTy() ? nullptr : llvm::ConstantAggregateZero::get(TyRet);
+  for (auto &fn : functions) {
+    auto next = fn(ret, IO);
+    if (auto err = next.takeError())
+      return err;
+    ret = next.get();
+  }
+  return ret;
+}
+
+} // namespace dxmt::dxbc
diff --git a/src/airconv/nt/dxbc_io_helper.hpp b/src/airconv/nt/dxbc_io_helper.hpp
new file mode 100644
index 000000000..2c8fcda05
--- /dev/null
+++ b/src/airconv/nt/dxbc_io_helper.hpp
@@ -0,0 +1,132 @@
+#pragma once
+#include "llvm/IR/IRBuilder.h"
+#include "llvm/Support/Error.h"
+#include "air_builder.hpp"
+#include "../airconv_public.h"
+#include "../air_signature.hpp"
+#include "../air_type.hpp"
+#include "../shader_common.hpp"
+
+#include <cstdint>
+#include <functional>
+#include <vector>
+
+namespace dxmt::dxbc {
+
+struct io_binding_map;
+
+class IOHelper {
+public:
+  IOHelper(llvm::air::AIRBuilder &air, llvm::Function *function, air::AirType &types, io_binding_map &io_legacy);
+
+  llvm::Error InitializeInput(
+      uint32_t FromFunctionArgumentIndex, uint32_t ToRegister, uint32_t Mask, bool InvertWComponent = false
+  );
+
+  llvm::Error InitializeInterpolatedInput(
+      uint32_t FromFunctionArgumentIndex, uint32_t ToRegister, uint32_t Mask,
+      air::Interpolation Interpolation, uint32_t AtSampleIndex
+  );
+
+  llvm::Error PullVertexInput(
+      uint32_t VertexBufferTableArgumentIndex, uint32_t ToRegister, uint32_t Mask,
+      SM50_IA_INPUT_ELEMENT ElementInfo, uint32_t SlotMask
+  );
+
+  llvm::Error InitializeVertexIDInput(uint32_t ToRegister, uint32_t Mask);
+
+  llvm::Error InitializeInstanceIDInput(uint32_t ToRegister, uint32_t Mask);
+
+  llvm::Error InitializeInputAttribute(shader::common::InputAttribute Value, uint32_t FromFunctionArgumentIndex);
+
+  llvm::Error CreateDepthOutputRegister();
+
+  llvm::Error CreateStencilRefOutputRegister();
+
+  llvm::Error CreateCoverageMaskOutputRegister();
+
+  llvm::Value *
+  PopulateOutput(uint32_t FromRegister, uint32_t Mask, llvm::Value *StructRet, uint32_t AtElement);
+
+  llvm::Value *
+  PopulateOutputFixUnorm(uint32_t FromRegister, uint32_t Mask, llvm::Value *StructRet, uint32_t AtElement);
+
+  llvm::Value *
+  PopulateOutputSanitizePosition(uint32_t FromRegister, uint32_t Mask, llvm::Value *StructRet, uint32_t AtElement);
+
+  llvm::Value *
+  PopulateOutputDepth(llvm::Value *StructRet, uint32_t AtElement);
+
+  llvm::Value *
+  PopulateOutputStencilRef(llvm::Value *StructRet, uint32_t AtElement);
+
+  llvm::Value *
+  PopulateOutputCoverageMask(llvm::Value *StructRet, uint32_t AtElement, uint32_t SampleMask);
+
+  llvm::Value *
+  PopulateDefaultCoverageMask(llvm::Value *StructRet, uint32_t AtElement, uint32_t SampleMask);
+
+  llvm::Value *PopulateOutputStreamOutput(
+      llvm::Value *StructRet, const SM50_SHADER_EMULATE_VERTEX_STREAM_OUTPUT_DATA *VertexSO,
+      uint32_t BaseVertexArgumentIndex, uint32_t VertexIDArgumentIndex, uint32_t StreamOutputTableArgumentIndex
+  );
+
+  void PopulateMeshOutputRenderTargetArrayIndex(uint32_t FromRegister, uint32_t Mask, llvm::Value *PrimitiveId);
+
+  void PopulateMeshOutputViewportArrayIndex(uint32_t FromRegister, uint32_t Mask, llvm::Value *PrimitiveId);
+
+  void PopulateMeshOutputPosition(uint32_t FromRegister, uint32_t Mask, llvm::Value *VertexId);
+
+  void PopulateMeshOutputVertexData(
+      uint32_t FromRegister, uint32_t Mask, uint32_t Index, llvm::Value *VertexId,
+      air::MSLScalerOrVectorType DesiredType
+  );
+
+private:
+  llvm::Value *LoadFromArrayAt(llvm::Value *Array, llvm::Value *Index);
+  llvm::Value *LoadFromVec4ArrayMasked(llvm::Value *Array, llvm::Value *Index, uint32_t Mask);
+  llvm::Error StoreAtVec4ArrayMasked(llvm::Value *Array, llvm::Value *Index, llvm::Value *MaybeVec4, uint32_t Mask);
+  llvm::Value *ExtendToVec4(llvm::Value *Value);
+  llvm::Value *ToDesiredTypeFromIntVec4(llvm::Value *Vec4, llvm::Type *Desired, uint32_t Mask);
+  llvm::Value *GetFloat4Splat(float Value);
+
+  llvm::air::AIRBuilder &air;
+  llvm::IRBuilderBase &ir;
+  llvm::Function *function;
+  air::AirType &types;
+
+  /* Legacy */
+  io_binding_map &io;
+};
+
+using PrologueFn = std::function<llvm::Error(IOHelper &)>;
+
+class Prologues {
+public:
+  void
+  Add(PrologueFn &&Fn) {
+    functions.push_back(std::move(Fn));
+  }
+
+  llvm::Error Run(IOHelper &IO);
+
+private:
+  std::vector<PrologueFn> functions;
+};
+
+using EpilogueFn = std::function<llvm::Expected<llvm::Value *>(llvm::Value *, IOHelper &)>;
+
+class Epilogues {
+public:
+  void
+  Add(EpilogueFn &&Fn) {
+    functions.push_back(std::move(Fn));
+  }
+
+  llvm::Expected<llvm::Value *> Run(llvm::Type *TyRet, IOHelper &IO);
+
+private:
+  std::vector<EpilogueFn> functions;
+};
+
+} // namespace dxmt::dxbc