Add packed 8bit builtin types (#5939)

* Add packed bytes builtin type

* fix test

21 files changed, 113 insertions, 0 deletions

diff --git a/source/slang/slang-check-conversion.cpp b/source/slang/slang-check-conversion.cpp
index 74151a4bb..7cb71a1bd 100644
--- a/source/slang/slang-check-conversion.cpp
+++ b/source/slang/slang-check-conversion.cpp
@@ -699,6 +699,8 @@ int getTypeBitSize(Type* t)
         return 16;
     case BaseType::Int:
     case BaseType::UInt:
+    case BaseType::Int8x4Packed:
+    case BaseType::UInt8x4Packed:
         return 32;
     case BaseType::Int64:
     case BaseType::UInt64:
diff --git a/source/slang/slang-check-decl.cpp b/source/slang/slang-check-decl.cpp
index 31e7b70bf..c9497ce54 100644
--- a/source/slang/slang-check-decl.cpp
+++ b/source/slang/slang-check-decl.cpp
@@ -1939,6 +1939,8 @@ void SemanticsDeclHeaderVisitor::checkVarDeclCommon(VarDeclBase* varDecl)
             case BaseType::UInt:
             case BaseType::UInt64:
             case BaseType::UIntPtr:
+            case BaseType::Int8x4Packed:
+            case BaseType::UInt8x4Packed:
                 break;
             default:
                 getSink()->diagnose(varDecl, Diagnostics::staticConstRequirementMustBeIntOrBool);
diff --git a/source/slang/slang-emit-c-like.cpp b/source/slang/slang-emit-c-like.cpp
index b70069d42..3175f1b07 100644
--- a/source/slang/slang-emit-c-like.cpp
+++ b/source/slang/slang-emit-c-like.cpp
@@ -274,6 +274,11 @@ void CLikeSourceEmitter::emitSimpleType(IRType* type)
     case kIROp_UIntPtrType:
         return UnownedStringSlice("uintptr_t");
 
+    case kIROp_Int8x4PackedType:
+        return UnownedStringSlice("int8_t4_packed");
+    case kIROp_UInt8x4PackedType:
+        return UnownedStringSlice("uint8_t4_packed");
+
     case kIROp_HalfType:
         return UnownedStringSlice("half");
 
@@ -1272,6 +1277,8 @@ void CLikeSourceEmitter::emitSimpleValueImpl(IRInst* inst)
                         return;
                     }
                 case BaseType::UInt:
+                case BaseType::Int8x4Packed:
+                case BaseType::UInt8x4Packed:
                     {
                         m_writer->emit(UInt(uint32_t(litInst->value.intVal)));
                         m_writer->emit("U");
@@ -3896,6 +3903,8 @@ void CLikeSourceEmitter::emitVecNOrScalar(
                 m_writer->emit("ushort");
                 break;
             case kIROp_UIntType:
+            case kIROp_Int8x4PackedType:
+            case kIROp_UInt8x4PackedType:
                 m_writer->emit("uint");
                 break;
             case kIROp_UInt64Type:
diff --git a/source/slang/slang-emit-cpp.cpp b/source/slang/slang-emit-cpp.cpp
index ed51f7825..f91c4d06e 100644
--- a/source/slang/slang-emit-cpp.cpp
+++ b/source/slang/slang-emit-cpp.cpp
@@ -101,6 +101,10 @@ static const char s_xyzwNames[] = "xyzw";
     case kIROp_UIntPtrType:
         return UnownedStringSlice("uintptr_t");
 
+    case kIROp_Int8x4PackedType:
+    case kIROp_UInt8x4PackedType:
+        return UnownedStringSlice("uint32_t");
+
         // Not clear just yet how we should handle half... we want all processing as float
         // probly, but when reading/writing to memory converting
 
diff --git a/source/slang/slang-emit-cuda.cpp b/source/slang/slang-emit-cuda.cpp
index c51d0c20a..58ac377bf 100644
--- a/source/slang/slang-emit-cuda.cpp
+++ b/source/slang/slang-emit-cuda.cpp
@@ -84,6 +84,10 @@ UnownedStringSlice CUDASourceEmitter::getBuiltinTypeName(IROp op)
     case kIROp_UIntPtrType:
         return UnownedStringSlice("uint");
 #endif
+    case kIROp_Int8x4PackedType:
+    case kIROp_UInt8x4PackedType:
+        return UnownedStringSlice("uint");
+
     case kIROp_HalfType:
         return UnownedStringSlice("__half");
 
diff --git a/source/slang/slang-emit-glsl.cpp b/source/slang/slang-emit-glsl.cpp
index d6f3795e3..a19c9c01f 100644
--- a/source/slang/slang-emit-glsl.cpp
+++ b/source/slang/slang-emit-glsl.cpp
@@ -1236,6 +1236,8 @@ void GLSLSourceEmitter::emitSimpleValueImpl(IRInst* inst)
                         return;
                     }
                 case BaseType::UInt:
+                case BaseType::Int8x4Packed:
+                case BaseType::UInt8x4Packed:
                     {
                         m_writer->emit(UInt(uint32_t(litInst->value.intVal)));
                         m_writer->emit("U");
@@ -1984,6 +1986,8 @@ bool GLSLSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOu
                 break;
 
             case BaseType::UInt:
+            case BaseType::Int8x4Packed:
+            case BaseType::UInt8x4Packed:
                 if (fromType == BaseType::Float)
                 {
                     m_writer->emit("floatBitsToUint");
@@ -3050,6 +3054,18 @@ void GLSLSourceEmitter::emitSimpleTypeImpl(IRType* type)
 #endif
             return;
         }
+    case kIROp_Int8x4PackedType:
+        {
+            _requireBaseType(BaseType::Int8x4Packed);
+            m_writer->emit("uint");
+            return;
+        }
+    case kIROp_UInt8x4PackedType:
+        {
+            _requireBaseType(BaseType::UInt8x4Packed);
+            m_writer->emit("uint");
+            return;
+        }
     case kIROp_VoidType:
     case kIROp_BoolType:
     case kIROp_Int8Type:
diff --git a/source/slang/slang-emit-hlsl.cpp b/source/slang/slang-emit-hlsl.cpp
index c43f075b2..40d6f75d9 100644
--- a/source/slang/slang-emit-hlsl.cpp
+++ b/source/slang/slang-emit-hlsl.cpp
@@ -822,6 +822,8 @@ bool HLSLSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOu
             case BaseType::UInt64:
             case BaseType::UIntPtr:
             case BaseType::Bool:
+            case BaseType::Int8x4Packed:
+            case BaseType::UInt8x4Packed:
                 // Because the intermediate type will always
                 // be an integer type, we can convert to
                 // another integer type of the same size
@@ -861,6 +863,8 @@ bool HLSLSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOu
             case BaseType::UInt:
             case BaseType::Int:
             case BaseType::Bool:
+            case BaseType::Int8x4Packed:
+            case BaseType::UInt8x4Packed:
                 break;
             case BaseType::UInt16:
             case BaseType::Int16:
@@ -1193,6 +1197,8 @@ void HLSLSourceEmitter::emitSimpleTypeImpl(IRType* type)
     case kIROp_Int16Type:
     case kIROp_UInt16Type:
     case kIROp_HalfType:
+    case kIROp_Int8x4PackedType:
+    case kIROp_UInt8x4PackedType:
         {
             m_writer->emit(getDefaultBuiltinTypeName(type->getOp()));
             return;
diff --git a/source/slang/slang-emit-metal.cpp b/source/slang/slang-emit-metal.cpp
index 282d8f95c..5e4381d21 100644
--- a/source/slang/slang-emit-metal.cpp
+++ b/source/slang/slang-emit-metal.cpp
@@ -1076,6 +1076,10 @@ void MetalSourceEmitter::emitSimpleTypeImpl(IRType* type)
     case kIROp_UIntPtrType:
         m_writer->emit("ulong");
         return;
+    case kIROp_Int8x4PackedType:
+    case kIROp_UInt8x4PackedType:
+        m_writer->emit("uint");
+        return;
     case kIROp_StructType:
         m_writer->emit(getName(type));
         return;
diff --git a/source/slang/slang-emit-spirv.cpp b/source/slang/slang-emit-spirv.cpp
index 51ff9066c..676a16228 100644
--- a/source/slang/slang-emit-spirv.cpp
+++ b/source/slang/slang-emit-spirv.cpp
@@ -1465,6 +1465,8 @@ struct SPIRVEmitContext : public SourceEmitterBase, public SPIRVEmitSharedContex
         case kIROp_Int8Type:
         case kIROp_IntType:
         case kIROp_Int64Type:
+        case kIROp_Int8x4PackedType:
+        case kIROp_UInt8x4PackedType:
             {
                 const IntInfo i = getIntTypeInfo(as<IRType>(inst));
                 if (i.width == 16)
@@ -7366,6 +7368,8 @@ struct SPIRVEmitContext : public SourceEmitterBase, public SPIRVEmitSharedContex
             case kIROp_UInt64Type:
             case kIROp_UInt8Type:
             case kIROp_UIntPtrType:
+            case kIROp_Int8x4PackedType:
+            case kIROp_UInt8x4PackedType:
                 spvEncoding = 6; // Unsigned
                 break;
             case kIROp_FloatType:
diff --git a/source/slang/slang-emit-wgsl.cpp b/source/slang/slang-emit-wgsl.cpp
index da5e38327..f9181a50d 100644
--- a/source/slang/slang-emit-wgsl.cpp
+++ b/source/slang/slang-emit-wgsl.cpp
@@ -495,6 +495,10 @@ void WGSLSourceEmitter::emitSimpleTypeImpl(IRType* type)
     case kIROp_UIntPtrType:
         m_writer->emit("u64");
         return;
+    case kIROp_Int8x4PackedType:
+    case kIROp_UInt8x4PackedType:
+        m_writer->emit("u32");
+        return;
     case kIROp_StructType:
         m_writer->emit(getName(type));
         return;
@@ -940,6 +944,8 @@ void WGSLSourceEmitter::emitSimpleValueImpl(IRInst* inst)
                         return;
                     }
                 case BaseType::UInt:
+                case BaseType::Int8x4Packed:
+                case BaseType::UInt8x4Packed:
                     {
                         m_writer->emit("u32(");
                         m_writer->emit(UInt(uint32_t(litInst->value.intVal)));
diff --git a/source/slang/slang-ir-any-value-marshalling.cpp b/source/slang/slang-ir-any-value-marshalling.cpp
index 6dc01d495..673a3d168 100644
--- a/source/slang/slang-ir-any-value-marshalling.cpp
+++ b/source/slang/slang-ir-any-value-marshalling.cpp
@@ -153,6 +153,8 @@ struct AnyValueMarshallingContext
         case kIROp_IntPtrType:
         case kIROp_UIntPtrType:
         case kIROp_PtrType:
+        case kIROp_Int8x4PackedType:
+        case kIROp_UInt8x4PackedType:
             context->marshalBasicType(builder, dataType, concreteTypedVar);
             break;
         case kIROp_VectorType:
@@ -309,6 +311,8 @@ struct AnyValueMarshallingContext
                     break;
                 }
             case kIROp_UIntType:
+            case kIROp_Int8x4PackedType:
+            case kIROp_UInt8x4PackedType:
 #if SLANG_PTR_IS_32
             case kIROp_UIntPtrType:
 #endif
@@ -537,6 +541,8 @@ struct AnyValueMarshallingContext
                     break;
                 }
             case kIROp_UIntType:
+            case kIROp_Int8x4PackedType:
+            case kIROp_UInt8x4PackedType:
                 {
                     ensureOffsetAt4ByteBoundary();
                     if (fieldOffset < static_cast<uint32_t>(anyValInfo->fieldKeys.getCount()))
@@ -812,6 +818,8 @@ SlangInt _getAnyValueSizeRaw(IRType* type, SlangInt offset)
     case kIROp_FloatType:
     case kIROp_UIntType:
     case kIROp_BoolType:
+    case kIROp_Int8x4PackedType:
+    case kIROp_UInt8x4PackedType:
         return alignUp(offset, 4) + 4;
     case kIROp_UInt64Type:
     case kIROp_Int64Type:
diff --git a/source/slang/slang-ir-byte-address-legalize.cpp b/source/slang/slang-ir-byte-address-legalize.cpp
index 617c8c7c4..9207e6a2f 100644
--- a/source/slang/slang-ir-byte-address-legalize.cpp
+++ b/source/slang/slang-ir-byte-address-legalize.cpp
@@ -840,6 +840,8 @@ struct ByteAddressBufferLegalizationContext
         case kIROp_IntType:
         case kIROp_FloatType:
         case kIROp_BoolType:
+        case kIROp_Int8x4PackedType:
+        case kIROp_UInt8x4PackedType:
             // The basic 32-bit types (and `bool`) can be handled by
             // loading `uint` values and then bit-casting.
             //
diff --git a/source/slang/slang-ir-layout.cpp b/source/slang/slang-ir-layout.cpp
index 52810cb2e..8180ea6aa 100644
--- a/source/slang/slang-ir-layout.cpp
+++ b/source/slang/slang-ir-layout.cpp
@@ -128,6 +128,9 @@ static Result _calcSizeAndAlignment(
         BASE(UIntPtr, kPointerSize);
         BASE(Double, 8);
 
+        BASE(Int8x4Packed, 4);
+        BASE(UInt8x4Packed, 4);
+
         // We are currently handling `bool` following the HLSL
         // precednet of storing it in 4 bytes.
         //
diff --git a/source/slang/slang-ir-lower-bit-cast.cpp b/source/slang/slang-ir-lower-bit-cast.cpp
index 450ad8ac9..0a7783639 100644
--- a/source/slang/slang-ir-lower-bit-cast.cpp
+++ b/source/slang/slang-ir-lower-bit-cast.cpp
@@ -178,6 +178,8 @@ struct BitCastLoweringContext
         case kIROp_UIntType:
         case kIROp_FloatType:
         case kIROp_BoolType:
+        case kIROp_Int8x4PackedType:
+        case kIROp_UInt8x4PackedType:
 #if SLANG_PTR_IS_32
         case kIROp_IntPtrType:
         case kIROp_UIntPtrType:
diff --git a/source/slang/slang-ir-util.cpp b/source/slang/slang-ir-util.cpp
index 3d2a539a6..7788a50d5 100644
--- a/source/slang/slang-ir-util.cpp
+++ b/source/slang/slang-ir-util.cpp
@@ -105,6 +105,8 @@ IROp getTypeStyle(IROp op)
     case kIROp_UInt64Type:
     case kIROp_IntPtrType:
     case kIROp_UIntPtrType:
+    case kIROp_Int8x4PackedType:
+    case kIROp_UInt8x4PackedType:
         {
             // All int like
             return kIROp_IntType;
@@ -140,6 +142,8 @@ IROp getTypeStyle(BaseType op)
     case BaseType::UInt:
     case BaseType::UInt64:
     case BaseType::UIntPtr:
+    case BaseType::Int8x4Packed:
+    case BaseType::UInt8x4Packed:
         return kIROp_IntType;
     case BaseType::Half:
     case BaseType::Float:
@@ -445,6 +449,12 @@ void getTypeNameHint(StringBuilder& sb, IRInst* type)
     case kIROp_UIntPtrType:
         sb << "uintptr";
         break;
+    case kIROp_Int8x4PackedType:
+        sb << "int8_t4_packed";
+        break;
+    case kIROp_UInt8x4PackedType:
+        sb << "uint8_t4_packed";
+        break;
     case kIROp_CharType:
         sb << "char";
         break;
@@ -1735,6 +1745,10 @@ UnownedStringSlice getBasicTypeNameHint(IRType* basicType)
         return UnownedStringSlice::fromLiteral("uint64");
     case kIROp_UIntPtrType:
         return UnownedStringSlice::fromLiteral("uintptr");
+    case kIROp_Int8x4PackedType:
+        return UnownedStringSlice::fromLiteral("int8_t4_packed");
+    case kIROp_UInt8x4PackedType:
+        return UnownedStringSlice::fromLiteral("uint8_t4_packed");
     case kIROp_FloatType:
         return UnownedStringSlice::fromLiteral("float");
     case kIROp_HalfType:
diff --git a/source/slang/slang-ir.cpp b/source/slang/slang-ir.cpp
index 5e5d94b14..ff1cd49ea 100644
--- a/source/slang/slang-ir.cpp
+++ b/source/slang/slang-ir.cpp
@@ -3798,6 +3798,8 @@ IRInst* IRBuilder::emitDefaultConstruct(IRType* type, bool fallback)
     case kIROp_UIntType:
     case kIROp_UIntPtrType:
     case kIROp_UInt64Type:
+    case kIROp_Int8x4PackedType:
+    case kIROp_UInt8x4PackedType:
     case kIROp_CharType:
         return getIntValue(type, 0);
     case kIROp_BoolType:
@@ -7421,6 +7423,8 @@ bool isIntegralType(IRType* t)
         case BaseType::UInt64:
         case BaseType::IntPtr:
         case BaseType::UIntPtr:
+        case BaseType::Int8x4Packed:
+        case BaseType::UInt8x4Packed:
             return true;
         default:
             return false;
@@ -7467,6 +7471,10 @@ IntInfo getIntTypeInfo(const IRType* intType)
     case kIROp_Int64Type:
         return {64, true};
 
+    case kIROp_Int8x4PackedType:
+    case kIROp_UInt8x4PackedType:
+        return {32, false};
+
     case kIROp_IntPtrType:  // target platform dependent
     case kIROp_UIntPtrType: // target platform dependent
     default:
diff --git a/source/slang/slang-lower-to-ir.cpp b/source/slang/slang-lower-to-ir.cpp
index 6ff270dc6..5bbe44e9b 100644
--- a/source/slang/slang-lower-to-ir.cpp
+++ b/source/slang/slang-lower-to-ir.cpp
@@ -4551,6 +4551,8 @@ struct ExprLoweringVisitorBase : public ExprVisitor<Derived, LoweredValInfo>
             case BaseType::UInt64:
             case BaseType::UIntPtr:
             case BaseType::IntPtr:
+            case BaseType::Int8x4Packed:
+            case BaseType::UInt8x4Packed:
                 return LoweredValInfo::simple(getBuilder()->getIntValue(type, 0));
 
             case BaseType::Half:
diff --git a/source/slang/slang-mangle.cpp b/source/slang/slang-mangle.cpp
index 7d4587480..d955b7bd9 100644
--- a/source/slang/slang-mangle.cpp
+++ b/source/slang/slang-mangle.cpp
@@ -186,6 +186,11 @@ void emitBaseType(ManglingContext* context, BaseType baseType)
     case BaseType::IntPtr:
         emitRaw(context, "ip");
         break;
+    case BaseType::Int8x4Packed:
+        emitRaw(context, "c4p");
+        break;
+    case BaseType::UInt8x4Packed:
+        emitRaw(context, "C4p");
         break;
 
     default:
diff --git a/source/slang/slang-type-layout.cpp b/source/slang/slang-type-layout.cpp
index 17f1ae677..c6acf5250 100644
--- a/source/slang/slang-type-layout.cpp
+++ b/source/slang/slang-type-layout.cpp
@@ -110,6 +110,10 @@ struct DefaultLayoutRulesImpl : SimpleLayoutRulesImpl
                 sizeof(intptr_t),
                 sizeof(intptr_t));
 
+        case BaseType::Int8x4Packed:
+        case BaseType::UInt8x4Packed:
+            return SimpleLayoutInfo(LayoutResourceKind::Uniform, 4, 4);
+
         case BaseType::Half:
             return SimpleLayoutInfo(LayoutResourceKind::Uniform, 2, 2);
         case BaseType::Float:
diff --git a/source/slang/slang-type-system-shared.h b/source/slang/slang-type-system-shared.h
index f80267d2b..d11bc8cd3 100644
--- a/source/slang/slang-type-system-shared.h
+++ b/source/slang/slang-type-system-shared.h
@@ -22,6 +22,8 @@ namespace Slang
     X(Char)                  \
     X(IntPtr)                \
     X(UIntPtr)               \
+    X(Int8x4Packed)          \
+    X(UInt8x4Packed)         \
     /* end */
 
 enum class BaseType
diff --git a/source/slang/slang.cpp b/source/slang/slang.cpp
index 7d18cb50b..9d8e3a68b 100644
--- a/source/slang/slang.cpp
+++ b/source/slang/slang.cpp
@@ -81,6 +81,8 @@ namespace Slang
      BaseTypeInfo::Flag::Signed | BaseTypeInfo::Flag::Integer,
      uint8_t(BaseType::IntPtr)},
     {uint8_t(sizeof(uintptr_t)), BaseTypeInfo::Flag::Integer, uint8_t(BaseType::UIntPtr)},
+    {uint8_t(sizeof(uint32_t)), BaseTypeInfo::Flag::Integer, uint8_t(BaseType::Int8x4Packed)},
+    {uint8_t(sizeof(uint32_t)), BaseTypeInfo::Flag::Integer, uint8_t(BaseType::UInt8x4Packed)},
 };
 
 /* static */ bool BaseTypeInfo::check()
@@ -132,6 +134,10 @@ namespace Slang
         return UnownedStringSlice::fromLiteral("intptr_t");
     case BaseType::UIntPtr:
         return UnownedStringSlice::fromLiteral("uintptr_t");
+    case BaseType::Int8x4Packed:
+        return UnownedStringSlice::fromLiteral("int8_t4_packed");
+    case BaseType::UInt8x4Packed:
+        return UnownedStringSlice::fromLiteral("uint8_t4_packed");
     default:
         {
             SLANG_ASSERT(!"Unknown basic type");