Add packed 8bit builtin types (#5939)

* Add packed bytes builtin type * fix test
shader-slang · Dec 27, 2024 · 7cecc51 · 7cecc51
1 parent 2ad1f81
commit 7cecc51
Show file tree

Hide file tree

Showing 26 changed files with 273 additions and 0 deletions.
diff --git a/source/slang-core-module/slang-embedded-core-module-source.cpp b/source/slang-core-module/slang-embedded-core-module-source.cpp
@@ -56,6 +56,12 @@ enum BaseTypeConversionRank : uint8_t
     kBaseTypeConversionRank_Int32,
     kBaseTypeConversionRank_IntPtr,
     kBaseTypeConversionRank_Int64,
+
+    // Packed type conversion ranks where the overall rank order does not apply.
+    // They must be explicitly casted to another type.
+    kBaseTypeConversionRank_Int8x4Packed,
+    kBaseTypeConversionRank_UInt8x4Packed,
+
     kBaseTypeConversionRank_Error,
 };
 
@@ -150,6 +156,16 @@ static const BaseTypeConversionInfo kBaseTypes[] = {
      kBaseTypeConversionKind_Unsigned,
      kBaseTypeConversionRank_IntPtr},
 
+    {"int8_t4_packed",
+     BaseType::Int8x4Packed,
+     0,
+     kBaseTypeConversionKind_Unsigned,
+     kBaseTypeConversionRank_Int8x4Packed},
+    {"uint8_t4_packed",
+     BaseType::UInt8x4Packed,
+     0,
+     kBaseTypeConversionKind_Unsigned,
+     kBaseTypeConversionRank_UInt8x4Packed},
 };
 
 void Session::finalizeSharedASTBuilder()
@@ -176,6 +192,11 @@ void Session::finalizeSharedASTBuilder()
         globalAstBuilder->getBuiltinType(baseType.tag);
 }
 
+static bool isConversionRankPackedType(BaseTypeConversionRank rank)
+{
+    return (rank == BaseTypeConversionRank::kBaseTypeConversionRank_Int8x4Packed) ||
+           (rank == BaseTypeConversionRank::kBaseTypeConversionRank_UInt8x4Packed);
+}
 
 // Given two base types, we need to be able to compute the cost of converting between them.
 ConversionCost getBaseTypeConversionCost(
@@ -189,6 +210,14 @@ ConversionCost getBaseTypeConversionCost(
         return kConversionCost_None;
     }
 
+    // Handle special case for packed types, where they must be explicitly casted to another type.
+    bool isToPackedType = isConversionRankPackedType(toInfo.conversionRank);
+    bool isFromPackedType = isConversionRankPackedType(fromInfo.conversionRank);
+    if (isToPackedType || isFromPackedType)
+    {
+        return kConversionCost_GeneralConversion;
+    }
+
     // Conversions within the same kind are easist to handle
     if (toInfo.conversionKind == fromInfo.conversionKind)
     {

diff --git 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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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: