Skip to content

Commit f933014

Browse files
fairywreathfairywreath
committed
remove builtin packed types, use typealias instead
1 parent 55c127c commit f933014

27 files changed

+18
-291
lines changed

‎source/slang-core-module/slang-embedded-core-module-source.cpp

-31
Original file line numberDiff line numberDiff line change
@@ -56,12 +56,6 @@ enum BaseTypeConversionRank : uint8_t
5656
kBaseTypeConversionRank_Int32,
5757
kBaseTypeConversionRank_IntPtr,
5858
kBaseTypeConversionRank_Int64,
59-
60-
// Packed type conversion ranks where the overall rank order does not apply.
61-
// They must be explicitly casted to another type.
62-
kBaseTypeConversionRank_Int8x4Packed,
63-
kBaseTypeConversionRank_UInt8x4Packed,
64-
6559
kBaseTypeConversionRank_Error,
6660
};
6761

@@ -155,17 +149,6 @@ static const BaseTypeConversionInfo kBaseTypes[] = {
155149
UINT_MASK,
156150
kBaseTypeConversionKind_Unsigned,
157151
kBaseTypeConversionRank_IntPtr},
158-
159-
{"int8_t4_packed",
160-
BaseType::Int8x4Packed,
161-
0,
162-
kBaseTypeConversionKind_Unsigned,
163-
kBaseTypeConversionRank_Int8x4Packed},
164-
{"uint8_t4_packed",
165-
BaseType::UInt8x4Packed,
166-
0,
167-
kBaseTypeConversionKind_Unsigned,
168-
kBaseTypeConversionRank_UInt8x4Packed},
169152
};
170153

171154
void Session::finalizeSharedASTBuilder()
@@ -192,12 +175,6 @@ void Session::finalizeSharedASTBuilder()
192175
globalAstBuilder->getBuiltinType(baseType.tag);
193176
}
194177

195-
static bool isConversionRankPackedType(BaseTypeConversionRank rank)
196-
{
197-
return (rank == BaseTypeConversionRank::kBaseTypeConversionRank_Int8x4Packed) ||
198-
(rank == BaseTypeConversionRank::kBaseTypeConversionRank_UInt8x4Packed);
199-
}
200-
201178
// Given two base types, we need to be able to compute the cost of converting between them.
202179
ConversionCost getBaseTypeConversionCost(
203180
BaseTypeConversionInfo const& toInfo,
@@ -210,14 +187,6 @@ ConversionCost getBaseTypeConversionCost(
210187
return kConversionCost_None;
211188
}
212189

213-
// Handle special case for packed types, where they must be explicitly casted to another type.
214-
bool isToPackedType = isConversionRankPackedType(toInfo.conversionRank);
215-
bool isFromPackedType = isConversionRankPackedType(fromInfo.conversionRank);
216-
if (isToPackedType || isFromPackedType)
217-
{
218-
return kConversionCost_GeneralConversion;
219-
}
220-
221190
// Conversions within the same kind are easist to handle
222191
if (toInfo.conversionKind == fromInfo.conversionKind)
223192
{

‎source/slang/hlsl.meta.slang

+16-13
Original file line numberDiff line numberDiff line change
@@ -24017,21 +24017,24 @@ T workgroupUniformLoad<T>(__ref T v)
2401724017
}
2401824018

2401924019
//
24020-
// Pack/Unpack Math Intrinsics
24020+
// HLSL Pack/Unpack Math Intrinsics
2402124021
//
2402224022
// These were introduced in SM 6.6 but requirements are dropped to SM 5.0 here
2402324023
// to expose these intrinsics on targets that do not have SM 6.6 features.
2402424024
//
2402524025

2402624026
//@public:
2402724027

24028+
typealias uint8_t4_packed = uint;
24029+
typealias int8_t4_packed = uint;
24030+
2402824031
/// Unpack 4 signed 8-bit values into a vector of 16 bit integers.
2402924032
[__readNone]
2403024033
[ForceInline]
2403124034
[require(cpp_cuda_glsl_hlsl_metal_spirv_wgsl, shader5_sm_5_0)]
2403224035
int16_t4 unpack_s8s16(int8_t4_packed packed)
2403324036
{
24034-
return unpackInt4x8ToInt16(uint(packed));
24037+
return unpackInt4x8ToInt16(packed);
2403524038
}
2403624039

2403724040
/// Unpack 4 unsigned 8-bit values into a vector of 16 bit integers.
@@ -24040,7 +24043,7 @@ int16_t4 unpack_s8s16(int8_t4_packed packed)
2404024043
[require(cpp_cuda_glsl_hlsl_metal_spirv_wgsl, shader5_sm_5_0)]
2404124044
uint16_t4 unpack_u8u16(uint8_t4_packed packed)
2404224045
{
24043-
return unpackUint4x8ToUint16(uint(packed));
24046+
return unpackUint4x8ToUint16(packed);
2404424047
}
2404524048

2404624049
/// Unpack 4 signed 8-bit values into a vector of 32 bit integers.
@@ -24049,7 +24052,7 @@ uint16_t4 unpack_u8u16(uint8_t4_packed packed)
2404924052
[require(cpp_cuda_glsl_hlsl_metal_spirv_wgsl, shader5_sm_5_0)]
2405024053
int32_t4 unpack_s8s32(int8_t4_packed packed)
2405124054
{
24052-
return unpackInt4x8ToInt32(uint(packed));
24055+
return unpackInt4x8ToInt32(packed);
2405324056
}
2405424057

2405524058
/// Unpack 4 unsigned 8-bit values into a vector of 32 bit integers.
@@ -24058,7 +24061,7 @@ int32_t4 unpack_s8s32(int8_t4_packed packed)
2405824061
[require(cpp_cuda_glsl_hlsl_metal_spirv_wgsl, shader5_sm_5_0)]
2405924062
uint32_t4 unpack_u8u32(uint8_t4_packed packed)
2406024063
{
24061-
return unpackUint4x8ToUint32(uint(packed));
24064+
return unpackUint4x8ToUint32(packed);
2406224065
}
2406324066

2406424067
/// Pack a vector of 4 unsigned 32 bit integers into a packed value of 4 8-bit integers, dropping unused bits.
@@ -24067,7 +24070,7 @@ uint32_t4 unpack_u8u32(uint8_t4_packed packed)
2406724070
[require(cpp_cuda_glsl_hlsl_metal_spirv_wgsl, shader5_sm_5_0)]
2406824071
uint8_t4_packed pack_u8(uint32_t4 unpackedValue)
2406924072
{
24070-
return uint8_t4_packed(packUint4x8(unpackedValue));
24073+
return packUint4x8(unpackedValue);
2407124074
}
2407224075

2407324076
/// Pack a vector of 4 signed 32 bit integers into a packed value of 4 8-bit integers, dropping unused bits.
@@ -24076,7 +24079,7 @@ uint8_t4_packed pack_u8(uint32_t4 unpackedValue)
2407624079
[require(cpp_cuda_glsl_hlsl_metal_spirv_wgsl, shader5_sm_5_0)]
2407724080
int8_t4_packed pack_s8(int32_t4 unpackedValue)
2407824081
{
24079-
return int8_t4_packed(packInt4x8(unpackedValue));
24082+
return packInt4x8(unpackedValue);
2408024083
}
2408124084

2408224085
/// Pack a vector of 4 unsigned 16 bit integers into a packed value of 4 8-bit integers, dropping unused bits.
@@ -24085,7 +24088,7 @@ int8_t4_packed pack_s8(int32_t4 unpackedValue)
2408524088
[require(cpp_cuda_glsl_hlsl_metal_spirv_wgsl, shader5_sm_5_0)]
2408624089
uint8_t4_packed pack_u8(uint16_t4 unpackedValue)
2408724090
{
24088-
return uint8_t4_packed(packUint4x8(unpackedValue));
24091+
return packUint4x8(unpackedValue);
2408924092
}
2409024093

2409124094
/// Pack a vector of 4 signed 16 bit integers into a packed value of 4 8-bit integers, dropping unused bits.
@@ -24094,7 +24097,7 @@ uint8_t4_packed pack_u8(uint16_t4 unpackedValue)
2409424097
[require(cpp_cuda_glsl_hlsl_metal_spirv_wgsl, shader5_sm_5_0)]
2409524098
int8_t4_packed pack_s8(int16_t4 unpackedValue)
2409624099
{
24097-
return int8_t4_packed(packInt4x8(unpackedValue));
24100+
return packInt4x8(unpackedValue);
2409824101
}
2409924102

2410024103
/// Pack a vector of 4 unsigned 32 bit integers into a packed value of 4 8-bit integers,
@@ -24104,7 +24107,7 @@ int8_t4_packed pack_s8(int16_t4 unpackedValue)
2410424107
[require(cpp_cuda_glsl_hlsl_metal_spirv_wgsl, shader5_sm_5_0)]
2410524108
uint8_t4_packed pack_clamp_u8(int32_t4 unpackedValue)
2410624109
{
24107-
return uint8_t4_packed(packUint4x8Clamp(unpackedValue));
24110+
return packUint4x8Clamp(unpackedValue);
2410824111
}
2410924112

2411024113
/// Pack a vector of 4 signed 32 bit integers into a packed value of 4 8-bit integers,
@@ -24114,7 +24117,7 @@ uint8_t4_packed pack_clamp_u8(int32_t4 unpackedValue)
2411424117
[require(cpp_cuda_glsl_hlsl_metal_spirv_wgsl, shader5_sm_5_0)]
2411524118
int8_t4_packed pack_clamp_s8(int32_t4 unpackedValue)
2411624119
{
24117-
return int8_t4_packed(packInt4x8Clamp(unpackedValue));
24120+
return packInt4x8Clamp(unpackedValue);
2411824121
}
2411924122

2412024123
/// Pack a vector of 4 unsigned 16 bit integers into a packed value of 4 8-bit integers,
@@ -24124,7 +24127,7 @@ int8_t4_packed pack_clamp_s8(int32_t4 unpackedValue)
2412424127
[require(cpp_cuda_glsl_hlsl_metal_spirv_wgsl, shader5_sm_5_0)]
2412524128
uint8_t4_packed pack_clamp_u8(int16_t4 unpackedValue)
2412624129
{
24127-
return uint8_t4_packed(packUint4x8Clamp(unpackedValue));
24130+
return packUint4x8Clamp(unpackedValue);
2412824131
}
2412924132

2413024133
/// Pack a vector of 4 signed 16 bit integers into a packed value of 4 8-bit integers,
@@ -24134,7 +24137,7 @@ uint8_t4_packed pack_clamp_u8(int16_t4 unpackedValue)
2413424137
[require(cpp_cuda_glsl_hlsl_metal_spirv_wgsl, shader5_sm_5_0)]
2413524138
int8_t4_packed pack_clamp_s8(int16_t4 unpackedValue)
2413624139
{
24137-
return int8_t4_packed(packInt4x8Clamp(unpackedValue));
24140+
return packInt4x8Clamp(unpackedValue);
2413824141
}
2413924142

2414024143
// Work-graphs

‎source/slang/slang-check-conversion.cpp

-2
Original file line numberDiff line numberDiff line change
@@ -1026,8 +1026,6 @@ int getTypeBitSize(Type* t)
10261026
return 16;
10271027
case BaseType::Int:
10281028
case BaseType::UInt:
1029-
case BaseType::Int8x4Packed:
1030-
case BaseType::UInt8x4Packed:
10311029
return 32;
10321030
case BaseType::Int64:
10331031
case BaseType::UInt64:

‎source/slang/slang-check-decl.cpp

-2
Original file line numberDiff line numberDiff line change
@@ -1978,8 +1978,6 @@ void SemanticsDeclHeaderVisitor::checkVarDeclCommon(VarDeclBase* varDecl)
19781978
case BaseType::UInt:
19791979
case BaseType::UInt64:
19801980
case BaseType::UIntPtr:
1981-
case BaseType::Int8x4Packed:
1982-
case BaseType::UInt8x4Packed:
19831981
break;
19841982
default:
19851983
getSink()->diagnose(varDecl, Diagnostics::staticConstRequirementMustBeIntOrBool);

‎source/slang/slang-emit-c-like.cpp

-9
Original file line numberDiff line numberDiff line change
@@ -266,11 +266,6 @@ void CLikeSourceEmitter::emitSimpleType(IRType* type)
266266
case kIROp_UIntPtrType:
267267
return UnownedStringSlice("uintptr_t");
268268

269-
case kIROp_Int8x4PackedType:
270-
return UnownedStringSlice("int8_t4_packed");
271-
case kIROp_UInt8x4PackedType:
272-
return UnownedStringSlice("uint8_t4_packed");
273-
274269
case kIROp_HalfType:
275270
return UnownedStringSlice("half");
276271

@@ -1315,8 +1310,6 @@ void CLikeSourceEmitter::emitSimpleValueImpl(IRInst* inst)
13151310
return;
13161311
}
13171312
case BaseType::UInt:
1318-
case BaseType::Int8x4Packed:
1319-
case BaseType::UInt8x4Packed:
13201313
{
13211314
m_writer->emit(UInt(uint32_t(litInst->value.intVal)));
13221315
m_writer->emit("U");
@@ -4025,8 +4018,6 @@ void CLikeSourceEmitter::emitVecNOrScalar(
40254018
m_writer->emit("ushort");
40264019
break;
40274020
case kIROp_UIntType:
4028-
case kIROp_Int8x4PackedType:
4029-
case kIROp_UInt8x4PackedType:
40304021
m_writer->emit("uint");
40314022
break;
40324023
case kIROp_UInt64Type:

‎source/slang/slang-emit-cpp.cpp

-5
Original file line numberDiff line numberDiff line change
@@ -101,13 +101,8 @@ static const char s_xyzwNames[] = "xyzw";
101101
case kIROp_UIntPtrType:
102102
return UnownedStringSlice("uintptr_t");
103103

104-
case kIROp_Int8x4PackedType:
105-
case kIROp_UInt8x4PackedType:
106-
return UnownedStringSlice("uint32_t");
107-
108104
// Not clear just yet how we should handle half... we want all processing as float
109105
// probly, but when reading/writing to memory converting
110-
111106
case kIROp_HalfType:
112107
return UnownedStringSlice("half");
113108

‎source/slang/slang-emit-cuda.cpp

-3
Original file line numberDiff line numberDiff line change
@@ -84,9 +84,6 @@ UnownedStringSlice CUDASourceEmitter::getBuiltinTypeName(IROp op)
8484
case kIROp_UIntPtrType:
8585
return UnownedStringSlice("uint");
8686
#endif
87-
case kIROp_Int8x4PackedType:
88-
case kIROp_UInt8x4PackedType:
89-
return UnownedStringSlice("uint");
9087

9188
case kIROp_HalfType:
9289
return UnownedStringSlice("__half");

‎source/slang/slang-emit-glsl.cpp

-16
Original file line numberDiff line numberDiff line change
@@ -1330,8 +1330,6 @@ void GLSLSourceEmitter::emitSimpleValueImpl(IRInst* inst)
13301330
return;
13311331
}
13321332
case BaseType::UInt:
1333-
case BaseType::Int8x4Packed:
1334-
case BaseType::UInt8x4Packed:
13351333
{
13361334
m_writer->emit(UInt(uint32_t(litInst->value.intVal)));
13371335
m_writer->emit("U");
@@ -2169,8 +2167,6 @@ bool GLSLSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOu
21692167
break;
21702168

21712169
case BaseType::UInt:
2172-
case BaseType::Int8x4Packed:
2173-
case BaseType::UInt8x4Packed:
21742170
if (fromType == BaseType::Float)
21752171
{
21762172
m_writer->emit("floatBitsToUint");
@@ -3180,18 +3176,6 @@ void GLSLSourceEmitter::emitSimpleTypeImpl(IRType* type)
31803176
#endif
31813177
return;
31823178
}
3183-
case kIROp_Int8x4PackedType:
3184-
{
3185-
_requireBaseType(BaseType::Int8x4Packed);
3186-
m_writer->emit("uint");
3187-
return;
3188-
}
3189-
case kIROp_UInt8x4PackedType:
3190-
{
3191-
_requireBaseType(BaseType::UInt8x4Packed);
3192-
m_writer->emit("uint");
3193-
return;
3194-
}
31953179
case kIROp_VoidType:
31963180
case kIROp_BoolType:
31973181
case kIROp_Int8Type:

‎source/slang/slang-emit-hlsl.cpp

-6
Original file line numberDiff line numberDiff line change
@@ -902,8 +902,6 @@ bool HLSLSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOu
902902
case BaseType::UInt64:
903903
case BaseType::UIntPtr:
904904
case BaseType::Bool:
905-
case BaseType::Int8x4Packed:
906-
case BaseType::UInt8x4Packed:
907905
// Because the intermediate type will always
908906
// be an integer type, we can convert to
909907
// another integer type of the same size
@@ -943,8 +941,6 @@ bool HLSLSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOu
943941
case BaseType::UInt:
944942
case BaseType::Int:
945943
case BaseType::Bool:
946-
case BaseType::Int8x4Packed:
947-
case BaseType::UInt8x4Packed:
948944
break;
949945
case BaseType::UInt16:
950946
case BaseType::Int16:
@@ -1330,8 +1326,6 @@ void HLSLSourceEmitter::emitSimpleTypeImpl(IRType* type)
13301326
case kIROp_Int16Type:
13311327
case kIROp_UInt16Type:
13321328
case kIROp_HalfType:
1333-
case kIROp_Int8x4PackedType:
1334-
case kIROp_UInt8x4PackedType:
13351329
{
13361330
m_writer->emit(getDefaultBuiltinTypeName(type->getOp()));
13371331
return;

‎source/slang/slang-emit-metal.cpp

-4
Original file line numberDiff line numberDiff line change
@@ -1099,10 +1099,6 @@ void MetalSourceEmitter::emitSimpleTypeImpl(IRType* type)
10991099
case kIROp_UIntPtrType:
11001100
m_writer->emit("ulong");
11011101
return;
1102-
case kIROp_Int8x4PackedType:
1103-
case kIROp_UInt8x4PackedType:
1104-
m_writer->emit("uint");
1105-
return;
11061102
case kIROp_StructType:
11071103
m_writer->emit(getName(type));
11081104
return;

‎source/slang/slang-emit-spirv.cpp

-4
Original file line numberDiff line numberDiff line change
@@ -1466,8 +1466,6 @@ struct SPIRVEmitContext : public SourceEmitterBase, public SPIRVEmitSharedContex
14661466
case kIROp_Int8Type:
14671467
case kIROp_IntType:
14681468
case kIROp_Int64Type:
1469-
case kIROp_Int8x4PackedType:
1470-
case kIROp_UInt8x4PackedType:
14711469
{
14721470
const IntInfo i = getIntTypeInfo(as<IRType>(inst));
14731471
if (i.width == 16)
@@ -7551,8 +7549,6 @@ struct SPIRVEmitContext : public SourceEmitterBase, public SPIRVEmitSharedContex
75517549
case kIROp_UInt64Type:
75527550
case kIROp_UInt8Type:
75537551
case kIROp_UIntPtrType:
7554-
case kIROp_Int8x4PackedType:
7555-
case kIROp_UInt8x4PackedType:
75567552
spvEncoding = 6; // Unsigned
75577553
break;
75587554
case kIROp_FloatType:

‎source/slang/slang-emit-wgsl.cpp

-6
Original file line numberDiff line numberDiff line change
@@ -511,10 +511,6 @@ void WGSLSourceEmitter::emitSimpleTypeImpl(IRType* type)
511511
case kIROp_UIntPtrType:
512512
m_writer->emit("u64");
513513
return;
514-
case kIROp_Int8x4PackedType:
515-
case kIROp_UInt8x4PackedType:
516-
m_writer->emit("u32");
517-
return;
518514
case kIROp_StructType:
519515
m_writer->emit(getName(type));
520516
return;
@@ -967,8 +963,6 @@ void WGSLSourceEmitter::emitSimpleValueImpl(IRInst* inst)
967963
return;
968964
}
969965
case BaseType::UInt:
970-
case BaseType::Int8x4Packed:
971-
case BaseType::UInt8x4Packed:
972966
{
973967
m_writer->emit("u32(");
974968
m_writer->emit(UInt(uint32_t(litInst->value.intVal)));

0 commit comments

Comments
 (0)