summaryrefslogtreecommitdiffstats
path: root/source
diff options
context:
space:
mode:
Diffstat (limited to 'source')
-rw-r--r--source/slang/hlsl.meta.slang639
-rw-r--r--source/slang/slang-capabilities.capdef8
-rw-r--r--source/slang/slang-emit-metal.cpp1
-rw-r--r--source/slang/slang-stdlib-textures.cpp59
-rw-r--r--source/slang/slang-stdlib-textures.h7
5 files changed, 650 insertions, 64 deletions
diff --git a/source/slang/hlsl.meta.slang b/source/slang/hlsl.meta.slang
index 95ca03beb..92b68c3e6 100644
--- a/source/slang/hlsl.meta.slang
+++ b/source/slang/hlsl.meta.slang
@@ -964,14 +964,16 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,access,isShadow,0,forma
[__readNone]
[ForceInline]
- [require(glsl_hlsl_spirv, texture_querylod)]
+ [require(glsl_hlsl_metal_spirv, texture_querylod)]
float CalculateLevelOfDetail(SamplerState s, TextureCoord location)
{
__requireComputeDerivative();
__target_switch
{
case hlsl:
- __intrinsic_asm "CalculateLevelOfDetail";
+ __intrinsic_asm ".CalculateLevelOfDetail";
+ case metal:
+ __intrinsic_asm ".calculate_clamped_lod";
case glsl:
__intrinsic_asm "textureQueryLod($p, $2).x";
case spirv:
@@ -984,14 +986,16 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,access,isShadow,0,forma
[__readNone]
[ForceInline]
- [require(glsl_hlsl_spirv, texture_querylod)]
+ [require(glsl_hlsl_metal_spirv, texture_querylod)]
float CalculateLevelOfDetailUnclamped(SamplerState s, TextureCoord location)
{
__requireComputeDerivative();
__target_switch
{
case hlsl:
- __intrinsic_asm "CalculateLevelOfDetailUnclamped";
+ __intrinsic_asm ".CalculateLevelOfDetailUnclamped";
+ case metal:
+ __intrinsic_asm ".calculate_unclamped_lod";
case glsl:
__intrinsic_asm "textureQueryLod($p, $2).y";
case spirv:
@@ -1008,7 +1012,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format>
{
[__readNone]
[ForceInline]
- [require(cpp_cuda_glsl_hlsl_spirv, texture_sm_4_1_fragment)]
+ [require(cpp_cuda_glsl_hlsl_metal_spirv, texture_sm_4_1_fragment)]
T Sample(SamplerState s, vector<float, Shape.dimensions+isArray> location)
{
__requireComputeDerivative();
@@ -1017,6 +1021,32 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format>
case cpp:
case hlsl:
__intrinsic_asm ".Sample";
+ case metal:
+ if (isArray == 1)
+ {
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_1D):
+ __intrinsic_asm "$0.sample($1, ($2).x, uint(($2).y))";
+ case $(SLANG_TEXTURE_2D):
+ __intrinsic_asm "$0.sample($1, ($2).xy, uint(($2).z))";
+ case $(SLANG_TEXTURE_CUBE):
+ __intrinsic_asm "$0.sample($1, ($2).xyz, uint(($2).w))";
+ }
+ }
+ else
+ {
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_1D):
+ case $(SLANG_TEXTURE_2D):
+ case $(SLANG_TEXTURE_3D):
+ case $(SLANG_TEXTURE_CUBE):
+ __intrinsic_asm ".sample";
+ }
+ }
+ // TODO: This needs to be handled by the capability system
+ __intrinsic_asm "<invalid intrinsic>";
case glsl:
__intrinsic_asm "$ctexture($p, $2)$z";
case cuda:
@@ -1062,7 +1092,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format>
[__readNone]
[ForceInline]
- [require(cpp_glsl_hlsl_spirv, texture_sm_4_1_fragment)]
+ [require(cpp_glsl_hlsl_metal_spirv, texture_sm_4_1_fragment)]
T Sample(SamplerState s, vector<float, Shape.dimensions+isArray> location, constexpr vector<int, Shape.planeDimensions> offset)
{
__requireComputeDerivative();
@@ -1071,6 +1101,26 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format>
case cpp:
case hlsl:
__intrinsic_asm ".Sample";
+ case metal:
+ if (isArray == 1)
+ {
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_2D):
+ __intrinsic_asm "$0.sample($1, ($2).xy, uint(($2).z), $3)";
+ }
+ }
+ else
+ {
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_2D):
+ case $(SLANG_TEXTURE_3D):
+ __intrinsic_asm ".sample";
+ }
+ }
+ // TODO: This needs to be handled by the capability system
+ __intrinsic_asm "<invalid intrinsic>";
case glsl:
__intrinsic_asm "$ctextureOffset($p, $2, $3)$z";
case spirv:
@@ -1086,7 +1136,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format>
[__readNone]
[ForceInline]
__glsl_extension(GL_ARB_sparse_texture_clamp)
- [require(cpp_glsl_hlsl_spirv, texture_sm_4_1_fragment)]
+ [require(cpp_glsl_hlsl_metal_spirv, texture_sm_4_1_fragment)]
T Sample(SamplerState s, vector<float, Shape.dimensions+isArray> location, constexpr vector<int, Shape.planeDimensions> offset, float clamp)
{
__requireComputeDerivative();
@@ -1095,6 +1145,26 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format>
case cpp:
case hlsl:
__intrinsic_asm ".Sample";
+ case metal:
+ if (isArray == 1)
+ {
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_2D):
+ __intrinsic_asm "$0.sample($1, ($2).xy, uint(($2).z), min_lod_clamp($4), $3)";
+ }
+ }
+ else
+ {
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_2D):
+ case $(SLANG_TEXTURE_3D):
+ __intrinsic_asm "$0.sample($1, $2, min_lod_clamp($4), $3)";
+ }
+ }
+ // TODO: This needs to be handled by the capability system
+ __intrinsic_asm "<invalid intrinsic>";
case glsl:
__intrinsic_asm "$ctextureOffsetClampARB($p, $2, $3, $4)$z";
case spirv:
@@ -1110,7 +1180,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format>
[__readNone]
[ForceInline]
- [require(cpp_glsl_hlsl_spirv, texture_sm_4_1_fragment)]
+ [require(cpp_glsl_hlsl_metal_spirv, texture_sm_4_1_fragment)]
T Sample(SamplerState s, vector<float, Shape.dimensions+isArray> location, constexpr vector<int, Shape.planeDimensions> offset, float clamp, out uint status)
{
__target_switch
@@ -1124,7 +1194,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format>
[__readNone]
[ForceInline]
- [require(cpp_glsl_hlsl_spirv, texture_sm_4_1_fragment)]
+ [require(cpp_glsl_hlsl_metal_spirv, texture_sm_4_1_fragment)]
T SampleBias(SamplerState s, vector<float, Shape.dimensions+isArray> location, float bias)
{
__requireComputeDerivative();
@@ -1133,6 +1203,29 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format>
case cpp:
case hlsl:
__intrinsic_asm ".SampleBias";
+ case metal:
+ if (isArray == 1)
+ {
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_2D):
+ __intrinsic_asm "$0.sample($1, ($2).xy, uint(($2).z), bias($3))";
+ case $(SLANG_TEXTURE_CUBE):
+ __intrinsic_asm "$0.sample($1, ($2).xyz, uint(($2).w), bias($3))";
+ }
+ }
+ else
+ {
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_2D):
+ case $(SLANG_TEXTURE_3D):
+ case $(SLANG_TEXTURE_CUBE):
+ __intrinsic_asm "$0.sample($1, $2, bias($3))";
+ }
+ }
+ // TODO: This needs to be handled by the capability system
+ __intrinsic_asm "<invalid intrinsic>";
case glsl:
__intrinsic_asm "$ctexture($p, $2, $3)$z";
case spirv:
@@ -1147,7 +1240,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format>
[__readNone]
[ForceInline]
- [require(cpp_glsl_hlsl_spirv, texture_sm_4_1_fragment)]
+ [require(cpp_glsl_hlsl_metal_spirv, texture_sm_4_1_fragment)]
T SampleBias(SamplerState s, vector<float, Shape.dimensions+isArray> location, float bias, constexpr vector<int, Shape.planeDimensions> offset)
{
__requireComputeDerivative();
@@ -1156,6 +1249,26 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format>
case cpp:
case hlsl:
__intrinsic_asm ".SampleBias";
+ case metal:
+ if (isArray == 1)
+ {
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_2D):
+ __intrinsic_asm "$0.sample($1, ($2).xy, uint(($2).z), bias($3), $4)";
+ }
+ }
+ else
+ {
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_2D):
+ case $(SLANG_TEXTURE_3D):
+ __intrinsic_asm "$0.sample($1, $2, bias($3), $4)";
+ }
+ }
+ // TODO: This needs to be handled by the capability system
+ __intrinsic_asm "<invalid intrinsic>";
case glsl:
__intrinsic_asm "$ctextureOffset($p, $2, $4, $3)$z";
case spirv:
@@ -1170,7 +1283,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format>
[__readNone]
[ForceInline]
- [require(glsl_hlsl_spirv, texture_shadowlod)]
+ [require(glsl_hlsl_metal_spirv, texture_shadowlod)]
float SampleCmp(SamplerComparisonState s, vector<float, Shape.dimensions+isArray> location, float compareValue)
{
__target_switch
@@ -1190,6 +1303,27 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format>
}
case hlsl:
__intrinsic_asm ".SampleCmp";
+ case metal:
+ if (isArray == 1)
+ {
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_2D):
+ __intrinsic_asm "$0.sample_compare($1, ($2).xy, uint(($2).z), $3)";
+ case $(SLANG_TEXTURE_CUBE):
+ __intrinsic_asm "$0.sample_compare($1, ($2).xyz, uint(($2).w), $3)";
+ }
+ }
+ else
+ {
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_2D):
+ case $(SLANG_TEXTURE_CUBE):
+ __intrinsic_asm ".sample_compare";
+ }
+ }
+ __intrinsic_asm "<invalid intrinsic>";
case spirv:
return spirv_asm
{
@@ -1201,7 +1335,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format>
[__readNone]
[ForceInline]
- [require(glsl_hlsl_spirv, texture_shadowlod)]
+ [require(glsl_hlsl_metal_spirv, texture_shadowlod)]
float SampleCmpLevelZero(SamplerComparisonState s, vector<float, Shape.dimensions+isArray> location, float compareValue)
{
__target_switch
@@ -1217,6 +1351,27 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format>
}
case hlsl:
__intrinsic_asm ".SampleCmpLevelZero";
+ case metal:
+ if (isArray == 1)
+ {
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_2D):
+ __intrinsic_asm "$0.sample_compare($1, ($2).xy, uint(($2).z), $3, level(0))";
+ case $(SLANG_TEXTURE_CUBE):
+ __intrinsic_asm "$0.sample_compare($1, ($2).xyz, uint(($2).w), $3, level(0))";
+ }
+ }
+ else
+ {
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_2D):
+ case $(SLANG_TEXTURE_CUBE):
+ __intrinsic_asm "$0.sample_compare($1, $2, $3, level(0))";
+ }
+ }
+ __intrinsic_asm "<invalid intrinsic>";
case spirv:
const float zeroFloat = 0.0f;
return spirv_asm
@@ -1229,7 +1384,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format>
[__readNone]
[ForceInline]
- [require(glsl_hlsl_spirv, texture_shadowlod)]
+ [require(glsl_hlsl_metal_spirv, texture_shadowlod)]
float SampleCmp(SamplerComparisonState s, vector<float, Shape.dimensions+isArray> location, float compareValue, constexpr vector<int, Shape.planeDimensions> offset)
{
__target_switch
@@ -1245,6 +1400,24 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format>
}
case hlsl:
__intrinsic_asm ".SampleCmp";
+ case metal:
+ if (isArray == 1)
+ {
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_2D):
+ __intrinsic_asm "$0.sample_compare($1, ($2).xy, uint(($2).z), $3, $4)";
+ }
+ }
+ else
+ {
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_2D):
+ __intrinsic_asm ".sample_compare";
+ }
+ }
+ __intrinsic_asm "<invalid intrinsic>";
case spirv:
return spirv_asm
{
@@ -1256,7 +1429,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format>
[__readNone]
[ForceInline]
- [require(glsl_hlsl_spirv, texture_shadowlod)]
+ [require(glsl_hlsl_metal_spirv, texture_shadowlod)]
float SampleCmpLevelZero(SamplerComparisonState s, vector<float, Shape.dimensions+isArray> location, float compareValue, constexpr vector<int, Shape.planeDimensions> offset)
{
__target_switch
@@ -1272,6 +1445,26 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format>
}
case hlsl:
__intrinsic_asm ".SampleCmpLevelZero";
+ case metal:
+ if (isShadow == 1)
+ {
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_2D):
+ if (isArray == 1)
+ {
+ // T sample_compare(sampler s, float2 coord, uint array, float compare_value, lod_options options, int2 offset = int2(0)) const
+ __intrinsic_asm "$0.sample_compare($1, ($2).xy, uint(($2).z), $3, level(0), $4)";
+ }
+ else
+ {
+ // T sample_compare(sampler s, float2 coord, float compare_value, lod_options options, int2 offset = int2(0)) const
+ __intrinsic_asm "$0.sample_compare($1, $2, $3, level(0), $4)";
+ }
+ break;
+ }
+ }
+ __intrinsic_asm "<invalid intrinsic>";
case spirv:
const float zeroFloat = 0.0f;
return spirv_asm
@@ -1284,14 +1477,39 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format>
[__readNone]
[ForceInline]
- [require(cpp_glsl_hlsl_spirv, texture_sm_4_1)]
+ [require(cpp_glsl_hlsl_metal_spirv, texture_sm_4_1)]
T SampleGrad(SamplerState s, vector<float, Shape.dimensions+isArray> location, vector<float, Shape.dimensions> gradX, vector<float, Shape.dimensions> gradY)
{
__target_switch
{
case cpp:
case hlsl:
- __intrinsic_asm ".SampleGrad";
+ __intrinsic_asm ".SampleGrad";
+ case metal:
+ if (isArray == 1)
+ {
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_2D):
+ __intrinsic_asm "$0.sample($1, ($2).xy, uint(($2).z), gradient2d($3, $4))";
+ case $(SLANG_TEXTURE_CUBE):
+ __intrinsic_asm "$0.sample($1, ($2).xyz, uint(($2).w), gradientcube($3, $4))";
+ }
+ }
+ else
+ {
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_2D):
+ __intrinsic_asm "$0.sample($1, $2, gradient2d($3, $4))";
+ case $(SLANG_TEXTURE_3D):
+ __intrinsic_asm "$0.sample($1, $2, gradient3d($3, $4))";
+ case $(SLANG_TEXTURE_CUBE):
+ __intrinsic_asm "$0.sample($1, $2, gradientcube($3, $4))";
+ }
+ }
+ // TODO: This needs to be handled by the capability system
+ __intrinsic_asm "<invalid intrinsic>";
case glsl:
__intrinsic_asm "$ctextureGrad($p, $2, $3, $4)$z";
case spirv:
@@ -1306,14 +1524,35 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format>
[__readNone]
[ForceInline]
- [require(cpp_glsl_hlsl_spirv, texture_sm_4_1)]
+ [require(cpp_glsl_hlsl_metal_spirv, texture_sm_4_1)]
T SampleGrad(SamplerState s, vector<float, Shape.dimensions+isArray> location, vector<float, Shape.dimensions> gradX, vector<float, Shape.dimensions> gradY, constexpr vector<int, Shape.dimensions> offset)
{
__target_switch
{
case cpp:
case hlsl:
- __intrinsic_asm ".SampleGrad";
+ __intrinsic_asm ".SampleGrad";
+ case metal:
+ if (isArray == 1)
+ {
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_2D):
+ __intrinsic_asm "$0.sample($1, ($2).xy, uint(($2).z), gradient2d($3, $4), $5)";
+ }
+ }
+ else
+ {
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_2D):
+ __intrinsic_asm "$0.sample($1, $2, gradient2d($3, $4), $5)";
+ case $(SLANG_TEXTURE_3D):
+ __intrinsic_asm "$0.sample($1, $2, gradient3d($3, $4), $5)";
+ }
+ }
+ // TODO: This needs to be handled by the capability system
+ __intrinsic_asm "<invalid intrinsic>";
case glsl:
__intrinsic_asm "$ctextureGradOffset($p, $2, $3, $4, $5)$z";
case spirv:
@@ -1330,14 +1569,35 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format>
[__readNone]
[ForceInline]
__glsl_extension(GL_ARB_sparse_texture_clamp)
- [require(cpp_glsl_hlsl_spirv, texture_sm_4_1)]
+ [require(cpp_glsl_hlsl_metal_spirv, texture_sm_4_1)]
T SampleGrad(SamplerState s, vector<float, Shape.dimensions+isArray> location, vector<float, Shape.dimensions> gradX, vector<float, Shape.dimensions> gradY, constexpr vector<int, Shape.dimensions> offset, float lodClamp)
{
__target_switch
{
case cpp:
case hlsl:
- __intrinsic_asm ".SampleGrad";
+ __intrinsic_asm ".SampleGrad";
+ case metal:
+ if (isArray == 1)
+ {
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_2D):
+ __intrinsic_asm "$0.sample($1, ($2).xy, uint(($2).z), gradient2d($3, $4), min_lod_clamp($6), $5)";
+ }
+ }
+ else
+ {
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_2D):
+ __intrinsic_asm "$0.sample($1, $2, gradient2d($3, $4), min_lod_clamp($6), $5)";
+ case $(SLANG_TEXTURE_3D):
+ __intrinsic_asm "$0.sample($1, $2, gradient3d($3, $4), min_lod_clamp($6), $5)";
+ }
+ }
+ // TODO: This needs to be handled by the capability system
+ __intrinsic_asm "<invalid intrinsic>";
case glsl:
__intrinsic_asm "$ctextureGradOffsetClampARB($p, $2, $3, $4, $5, $6)$z";
case spirv:
@@ -1353,7 +1613,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format>
[__readNone]
[ForceInline]
- [require(cpp_cuda_glsl_hlsl_spirv, texture_sm_4_1)]
+ [require(cpp_cuda_glsl_hlsl_metal_spirv, texture_sm_4_1)]
T SampleLevel(SamplerState s, vector<float, Shape.dimensions+isArray> location, float level)
{
__target_switch
@@ -1361,6 +1621,29 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format>
case cpp:
case hlsl:
__intrinsic_asm ".SampleLevel";
+ case metal:
+ if (isArray == 1)
+ {
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_2D):
+ __intrinsic_asm "$0.sample($1, ($2).xy, uint(($2).z), level($3))";
+ case $(SLANG_TEXTURE_CUBE):
+ __intrinsic_asm "$0.sample($1, ($2).xyz, uint(($2).w), level($3))";
+ }
+ }
+ else
+ {
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_2D):
+ case $(SLANG_TEXTURE_3D):
+ case $(SLANG_TEXTURE_CUBE):
+ __intrinsic_asm "$0.sample($1, $2, level($3))";
+ }
+ }
+ // TODO: This needs to be handled by the capability system
+ __intrinsic_asm "<invalid intrinsic>";
case glsl:
__intrinsic_asm "$ctextureLod($p, $2, $3)$z";
case cuda:
@@ -1406,7 +1689,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format>
[__readNone]
[ForceInline]
- [require(cpp_glsl_hlsl_spirv, texture_sm_4_1)]
+ [require(cpp_glsl_hlsl_metal_spirv, texture_sm_4_1)]
T SampleLevel(SamplerState s, vector<float, Shape.dimensions+isArray> location, float level, constexpr vector<int, Shape.planeDimensions> offset)
{
__target_switch
@@ -1414,6 +1697,28 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format>
case cpp:
case hlsl:
__intrinsic_asm ".SampleLevel";
+ case metal:
+ if (isArray == 1)
+ {
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_2D):
+ __intrinsic_asm "$0.sample($1, ($2).xy, uint(($2).z), level($3), $4)";
+ case $(SLANG_TEXTURE_CUBE):
+ __intrinsic_asm "$0.sample($1, ($2).xyz, uint(($2).w), level($3), $4)";
+ }
+ }
+ else
+ {
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_2D):
+ case $(SLANG_TEXTURE_3D):
+ case $(SLANG_TEXTURE_CUBE):
+ __intrinsic_asm "$0.sample($1, $2, level($3), $4)";
+ }
+ }
+ __intrinsic_asm "<invalid intrinsic>";
case glsl:
__intrinsic_asm "$ctextureLodOffset($p, $2, $3, $4)$z";
case spirv:
@@ -1474,13 +1779,46 @@ Array<T,4> __makeArray<T>(T v0, T v1, T v2, T v3);
// Gather for scalar textures.
__generic<TElement, T, Shape: __ITextureShape, let isArray:int, let sampleCount:int, let access:int, let isShadow:int, let format:int>
[ForceInline]
-[require(glsl_spirv, GLSL_400)]
-vector<TElement,4> __glsl_gather(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 0, format> texture, SamplerState s, vector<float, Shape.dimensions+isArray> location, int component)
+[require(glsl_metal_spirv, GLSL_400)]
+vector<TElement,4> __texture_gather(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 0, format> texture, SamplerState s, vector<float, Shape.dimensions+isArray> location, int component)
{
__target_switch
{
case glsl:
__intrinsic_asm "textureGather($p, $2, $3)";
+ case metal:
+ if (isShadow == 0)
+ {
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_2D):
+ if (isArray == 1)
+ {
+ // Tv gather(sampler s, float2 coord, uint array, int2 offset = int2(0), component c = component::x) const
+ __intrinsic_asm "$0.gather($1, ($2).xy, uint(($2).z), int2(0), metal::component($3))";
+ }
+ else
+ {
+ // Tv gather(sampler s, float2 coord, int2 offset = int2(0), component c = component::x) const
+ __intrinsic_asm "$0.gather($1, $2, int2(0), metal::component($3))";
+ }
+ break;
+ case $(SLANG_TEXTURE_CUBE):
+ if (isArray == 1)
+ {
+ // Tv gather(sampler s, float3 coord, uint array, component c = component::x) const
+ __intrinsic_asm "$0.gather($1, ($2).xyz, uint(($2).w), metal::component($3))";
+ }
+ else
+ {
+ // Tv gather(sampler s, float3 coord, component c = component::x) const
+ __intrinsic_asm "$0.gather($1, $2, metal::component($3))";
+ }
+ break;
+ }
+ }
+ // TODO: This needs to be handled by the capability system
+ __intrinsic_asm "<invalid intrinsic>";
case spirv:
return spirv_asm {
%sampledImage : __sampledImageType(texture) = OpSampledImage $texture $s;
@@ -1491,7 +1829,7 @@ vector<TElement,4> __glsl_gather(__TextureImpl<T, Shape, isArray, 0, sampleCount
__generic<TElement, T, Shape: __ITextureShape, let isArray:int, let sampleCount:int, let access:int, let isShadow:int, let format:int>
[ForceInline]
[require(glsl_spirv, GLSL_400)]
-vector<TElement,4> __glsl_gather(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 1, format> sampler, vector<float, Shape.dimensions+isArray> location, int component)
+vector<TElement,4> __texture_gather(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 1, format> sampler, vector<float, Shape.dimensions+isArray> location, int component)
{
__target_switch
{
@@ -1505,13 +1843,32 @@ vector<TElement,4> __glsl_gather(__TextureImpl<T, Shape, isArray, 0, sampleCount
}
__generic<TElement, T, Shape: __ITextureShape, let isArray:int, let sampleCount:int, let access:int, let isShadow:int, let format:int>
[ForceInline]
-[require(glsl_spirv, GLSL_400)]
-vector<TElement,4> __glsl_gather_offset(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 0, format> texture, SamplerState s, constexpr vector<float, Shape.dimensions+isArray> location, constexpr vector<int, Shape.planeDimensions> offset, int component)
+[require(glsl_metal_spirv, GLSL_400)]
+vector<TElement,4> __texture_gather_offset(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 0, format> texture, SamplerState s, constexpr vector<float, Shape.dimensions+isArray> location, constexpr vector<int, Shape.planeDimensions> offset, int component)
{
__target_switch
{
case glsl:
__intrinsic_asm "textureGatherOffset($p, $2, $3, $4)";
+ case metal:
+ if (Shape.flavor == $(SLANG_TEXTURE_2D))
+ {
+ if (isShadow == 0)
+ {
+ if (isArray == 1)
+ {
+ // Tv gather(sampler s, float2 coord, uint array, int2 offset = int2(0), component c = component::x) const
+ __intrinsic_asm "$0.gather($1, ($2).xy, uint(($2).z), $3, metal::component($4))";
+ }
+ else
+ {
+ // Tv gather(sampler s, float2 coord, int2 offset = int2(0), component c = component::x) const
+ __intrinsic_asm "$0.gather($1, $2, $3, metal::component($4))";
+ }
+ }
+ }
+ // TODO: This needs to be handled by the capability system
+ __intrinsic_asm "<Metal support gather with offset only for 2D>";
case spirv:
return spirv_asm {
%sampledImage : __sampledImageType(texture) = OpSampledImage $texture $s;
@@ -1522,7 +1879,7 @@ vector<TElement,4> __glsl_gather_offset(__TextureImpl<T, Shape, isArray, 0, samp
__generic<TElement, T, Shape: __ITextureShape, let isArray:int, let sampleCount:int, let access:int, let isShadow:int, let format:int>
[ForceInline]
[require(glsl_spirv, GLSL_400)]
-vector<TElement,4> __glsl_gather_offset(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 1, format> sampler, vector<float, Shape.dimensions+isArray> location, constexpr vector<int, Shape.planeDimensions> offset, int component)
+vector<TElement,4> __texture_gather_offset(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 1, format> sampler, vector<float, Shape.dimensions+isArray> location, constexpr vector<int, Shape.planeDimensions> offset, int component)
{
__target_switch
{
@@ -1537,7 +1894,7 @@ vector<TElement,4> __glsl_gather_offset(__TextureImpl<T, Shape, isArray, 0, samp
__generic<TElement, T, Shape: __ITextureShape, let isArray:int, let sampleCount:int, let access:int, let isShadow:int, let format:int>
[ForceInline]
[require(glsl_spirv, GLSL_400)]
-vector<TElement,4> __glsl_gather_offsets(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 0, format> texture, SamplerState s, vector<float, Shape.dimensions+isArray> location,
+vector<TElement,4> __texture_gather_offsets(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 0, format> texture, SamplerState s, vector<float, Shape.dimensions+isArray> location,
constexpr vector<int, Shape.planeDimensions> offset1,
constexpr vector<int, Shape.planeDimensions> offset2,
constexpr vector<int, Shape.planeDimensions> offset3,
@@ -1560,7 +1917,7 @@ vector<TElement,4> __glsl_gather_offsets(__TextureImpl<T, Shape, isArray, 0, sam
__generic<TElement, T, Shape: __ITextureShape, let isArray:int, let sampleCount:int, let access:int, let isShadow:int, let format:int>
[ForceInline]
[require(glsl_spirv, GLSL_400)]
-vector<TElement,4> __glsl_gather_offsets(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 1, format> sampler, vector<float, Shape.dimensions+isArray> location,
+vector<TElement,4> __texture_gather_offsets(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 1, format> sampler, vector<float, Shape.dimensions+isArray> location,
constexpr vector<int, Shape.planeDimensions> offset1,
constexpr vector<int, Shape.planeDimensions> offset2,
constexpr vector<int, Shape.planeDimensions> offset3,
@@ -1581,13 +1938,45 @@ vector<TElement,4> __glsl_gather_offsets(__TextureImpl<T, Shape, isArray, 0, sam
}
__generic<TElement, T, Shape: __ITextureShape, let isArray:int, let sampleCount:int, let access:int, let isShadow:int, let format:int>
[ForceInline]
-[require(glsl_spirv, GLSL_400)]
-vector<TElement,4> __glsl_gatherCmp(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 0, format> texture, SamplerComparisonState s, vector<float, Shape.dimensions+isArray> location, TElement compareValue)
+[require(glsl_metal_spirv, GLSL_400)]
+vector<TElement,4> __texture_gatherCmp(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 0, format> texture, SamplerComparisonState s, vector<float, Shape.dimensions+isArray> location, TElement compareValue)
{
__target_switch
{
case glsl:
__intrinsic_asm "textureGather($p, $2, $3)";
+ case metal:
+ if (isShadow == 1)
+ {
+ if (Shape.flavor == $(SLANG_TEXTURE_2D))
+ {
+ if (isArray == 1)
+ {
+ // Tv gather_compare(sampler s, float2 coord, uint array, float compare_value, int2 offset = int2(0)) const
+ __intrinsic_asm "$0.gather_compare($1, ($2).xy, uint(($2).z), $3)";
+ }
+ else
+ {
+ // Tv gather_compare(sampler s, float2 coord, float compare_value, int2 offset = int2(0)) const
+ __intrinsic_asm "$0.gather_compare($1, $2, $3)";
+ }
+ }
+ else if (Shape.flavor == $(SLANG_TEXTURE_CUBE))
+ {
+ if (isArray == 1)
+ {
+ // Tv gather_compare(sampler s, float3 coord, uint array, float compare_value) const
+ __intrinsic_asm "$0.gather_compare($1, ($2).xyz, uint(($2).w), $3)";
+ }
+ else
+ {
+ // Tv gather_compare(sampler s, float3 coord, float compare_value) const
+ __intrinsic_asm "$0.gather_compare($1, $2, $3)";
+ }
+ }
+ }
+ // TODO: This needs to be handled by the capability system
+ __intrinsic_asm "<invalid intrinsics>";
case spirv:
return spirv_asm {
%sampledImage : __sampledImageType(texture) = OpSampledImage $texture $s;
@@ -1598,7 +1987,7 @@ vector<TElement,4> __glsl_gatherCmp(__TextureImpl<T, Shape, isArray, 0, sampleCo
__generic<TElement, T, Shape: __ITextureShape, let isArray:int, let sampleCount:int, let access:int, let isShadow:int, let format:int>
[ForceInline]
[require(glsl_spirv, GLSL_400)]
-vector<TElement,4> __glsl_gatherCmp(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 1, format> sampler, vector<float, Shape.dimensions+isArray> location, TElement compareValue)
+vector<TElement,4> __texture_gatherCmp(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 1, format> sampler, vector<float, Shape.dimensions+isArray> location, TElement compareValue)
{
__target_switch
{
@@ -1612,13 +2001,32 @@ vector<TElement,4> __glsl_gatherCmp(__TextureImpl<T, Shape, isArray, 0, sampleCo
}
__generic<TElement, T, Shape: __ITextureShape, let isArray:int, let sampleCount:int, let access:int, let isShadow:int, let format:int>
[ForceInline]
-[require(glsl_spirv, GLSL_400)]
-vector<TElement,4> __glsl_gatherCmp_offset(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 0, format> texture, SamplerComparisonState s, vector<float, Shape.dimensions+isArray> location, TElement compareValue, constexpr vector<int, Shape.planeDimensions> offset)
+[require(glsl_metal_spirv, GLSL_400)]
+vector<TElement,4> __texture_gatherCmp_offset(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 0, format> texture, SamplerComparisonState s, vector<float, Shape.dimensions+isArray> location, TElement compareValue, constexpr vector<int, Shape.planeDimensions> offset)
{
__target_switch
{
case glsl:
__intrinsic_asm "textureGatherOffset($p, $2, $3, $4)";
+ case metal:
+ if (isShadow == 1)
+ {
+ if (Shape.flavor == $(SLANG_TEXTURE_2D))
+ {
+ if (isArray == 1)
+ {
+ // Tv gather_compare(sampler s, float2 coord, uint array, float compare_value, int2 offset = int2(0)) const
+ __intrinsic_asm "$0.gather_compare($1, ($2).xy, uint(($2).z), $3, $4)";
+ }
+ else
+ {
+ // Tv gather_compare(sampler s, float2 coord, float compare_value, int2 offset = int2(0)) const
+ __intrinsic_asm "$0.gather_compare($1, $2, $3, $4)";
+ }
+ }
+ }
+ // TODO: This needs to be handled by the capability system
+ __intrinsic_asm "<invalid intrinsics>";
case spirv:
return spirv_asm {
%sampledImage : __sampledImageType(texture) = OpSampledImage $texture $s;
@@ -1629,7 +2037,7 @@ vector<TElement,4> __glsl_gatherCmp_offset(__TextureImpl<T, Shape, isArray, 0, s
__generic<TElement, T, Shape: __ITextureShape, let isArray:int, let sampleCount:int, let access:int, let isShadow:int, let format:int>
[ForceInline]
[require(glsl_spirv, GLSL_400)]
-vector<TElement,4> __glsl_gatherCmp_offset(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 1, format> sampler, vector<float, Shape.dimensions+isArray> location, TElement compareValue, constexpr vector<int, Shape.planeDimensions> offset)
+vector<TElement,4> __texture_gatherCmp_offset(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 1, format> sampler, vector<float, Shape.dimensions+isArray> location, TElement compareValue, constexpr vector<int, Shape.planeDimensions> offset)
{
__target_switch
{
@@ -1644,7 +2052,7 @@ vector<TElement,4> __glsl_gatherCmp_offset(__TextureImpl<T, Shape, isArray, 0, s
__generic<TElement, T, Shape: __ITextureShape, let isArray:int, let sampleCount:int, let access:int, let isShadow:int, let format:int>
[ForceInline]
[require(glsl_spirv, GLSL_400)]
-vector<TElement,4> __glsl_gatherCmp_offsets(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 0, format> texture, SamplerComparisonState s, vector<float, Shape.dimensions+isArray> location, TElement compareValue,
+vector<TElement,4> __texture_gatherCmp_offsets(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 0, format> texture, SamplerComparisonState s, vector<float, Shape.dimensions+isArray> location, TElement compareValue,
vector<int, Shape.planeDimensions> offset1,
vector<int, Shape.planeDimensions> offset2,
vector<int, Shape.planeDimensions> offset3,
@@ -1666,7 +2074,7 @@ vector<TElement,4> __glsl_gatherCmp_offsets(__TextureImpl<T, Shape, isArray, 0,
__generic<TElement, T, Shape: __ITextureShape, let isArray:int, let sampleCount:int, let access:int, let isShadow:int, let format:int>
[ForceInline]
[require(glsl_spirv, GLSL_400)]
-vector<TElement,4> __glsl_gatherCmp_offsets(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 1, format> sampler, vector<float, Shape.dimensions+isArray> location, TElement compareValue,
+vector<TElement,4> __texture_gatherCmp_offsets(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 1, format> sampler, vector<float, Shape.dimensions+isArray> location, TElement compareValue,
vector<int, Shape.planeDimensions> offset1,
vector<int, Shape.planeDimensions> offset2,
vector<int, Shape.planeDimensions> offset3,
@@ -1703,7 +2111,9 @@ for (int isScalarTexture = 0; isScalarTexture <= 1; isScalarTexture++) {
${{{{
// Gather component
const char* samplerStateParam = isCombined ? "" : " s,";
- for (int isCmp = 0; isCmp <= 1; ++isCmp) {
+ const char* metalSupport = isCombined ? "" : "metal_";
+ const char* caseMetal = isCombined ? "" : "case metal:";
+ for (int isCmp = 0; isCmp < 2; ++isCmp) {
const char* cmp = isCmp ? "Cmp" : "";
const char* cmpParam = isCmp ? ", T compareValue" : "";
const char* compareArg = isCmp ? ", compareValue" : "";
@@ -1711,32 +2121,34 @@ ${{{{
const char* componentNames[] = { "", "Red", "Green", "Blue", "Alpha"};
const char* glslComponentNames[] = { ", 0", ", 1", ", 2", ", 3" };
- for (auto componentId = 0; componentId <= 4; componentId++) {
+ for (auto componentId = 0; componentId < 5; componentId++) {
auto componentName = componentNames[componentId];
auto glslComponent = (isCmp ? "" :glslComponentNames[componentId == 0 ? 0 : componentId - 1]);
}}}}
[ForceInline]
- [require(glsl_hlsl_spirv, texture_gather)]
+ [require(glsl_hlsl_$(metalSupport)spirv, texture_gather)]
vector<T,4> Gather$(cmp)$(componentName)($(samplerStateType)$(samplerStateParam) vector<float, Shape.dimensions+isArray> location $(cmpParam))
{
__target_switch
{
case hlsl: __intrinsic_asm ".Gather$(cmp)$(componentName)";
+ $(caseMetal)
case glsl:
case spirv:
- return __glsl_gather$(cmp)<T>(this,$(samplerStateParam) location $(compareArg) $(glslComponent));
+ return __texture_gather$(cmp)<T>(this,$(samplerStateParam) location $(compareArg) $(glslComponent));
}
}
[ForceInline]
- [require(glsl_hlsl_spirv, texture_gather)]
+ [require(glsl_hlsl_$(metalSupport)spirv, texture_gather)]
vector<T,4> Gather$(cmp)$(componentName)($(samplerStateType)$(samplerStateParam) vector<float, Shape.dimensions+isArray> location $(cmpParam), constexpr vector<int, Shape.planeDimensions> offset)
{
__target_switch
{
case hlsl: __intrinsic_asm ".Gather$(cmp)$(componentName)";
+ $(caseMetal)
case glsl:
case spirv:
- return __glsl_gather$(cmp)_offset<T>(this,$(samplerStateParam) location $(compareArg), offset $(glslComponent));
+ return __texture_gather$(cmp)_offset<T>(this,$(samplerStateParam) location $(compareArg), offset $(glslComponent));
}
}
[ForceInline]
@@ -1752,7 +2164,7 @@ ${{{{
case hlsl: __intrinsic_asm ".Gather$(cmp)$(componentName)";
case glsl:
case spirv:
- return __glsl_gather$(cmp)_offsets<T>(this,$(samplerStateParam) location $(compareArg), offset1,offset2,offset3,offset4 $(glslComponent));
+ return __texture_gather$(cmp)_offsets<T>(this,$(samplerStateParam) location $(compareArg), offset1,offset2,offset3,offset4 $(glslComponent));
}
}
${{{{
@@ -1785,7 +2197,7 @@ extension __TextureImpl<T,Shape,isArray,0,sampleCount,0,isShadow,isCombined,form
__glsl_extension(GL_EXT_samplerless_texture_functions)
[__readNone]
[ForceInline]
- [require(cpp_glsl_hlsl_spirv, texture_sm_4_1_samplerless)]
+ [require(cpp_glsl_hlsl_metal_spirv, texture_sm_4_1_samplerless)]
T Load(vector<int, Shape.dimensions+isArray+1> location)
{
__target_switch
@@ -1793,6 +2205,66 @@ extension __TextureImpl<T,Shape,isArray,0,sampleCount,0,isShadow,isCombined,form
case cpp:
case hlsl:
__intrinsic_asm ".Load";
+ case metal:
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_1D):
+ // lod is not supported for 1D texture
+ if (isArray == 1)
+ // Tv read(uint coord, uint array, uint lod = 0) const
+ __intrinsic_asm "$0.read(uint(($1).x), uint(($1).y))";
+ else
+ // Tv read(uint coord, uint lod = 0) const
+ __intrinsic_asm "$0.read(uint(($1).x))";
+ break;
+ case $(SLANG_TEXTURE_2D):
+ if (isShadow == 1)
+ {
+ if (isArray == 1)
+ // T read(uint2 coord, uint array, uint lod = 0) const
+ __intrinsic_asm "$0.read(vec<uint,2>(($1).xy), uint(($1).z), uint(($1).w))";
+ else
+ // T read(uint2 coord, uint lod = 0) const
+ __intrinsic_asm "$0.read(vec<uint,2>(($1).xy), uint(($1).z))";
+ }
+ else
+ {
+ if (isArray == 1)
+ // Tv read(uint2 coord, uint array, uint lod = 0) const
+ __intrinsic_asm "$0.read(vec<uint,2>(($1).xy), uint(($1).z), uint(($1).w))";
+ else
+ // Tv read(uint2 coord, uint lod = 0) const
+ __intrinsic_asm "$0.read(vec<uint,2>(($1).xy), uint(($1).z))";
+ }
+ break;
+ case $(SLANG_TEXTURE_3D):
+ if (isShadow == 0 && isArray == 0)
+ // Tv read(uint3 coord, uint lod = 0) const
+ __intrinsic_asm "$0.read(vec<uint,3>(($1).xyz), uint(($1).w))";
+ break;
+ case $(SLANG_TEXTURE_CUBE):
+ if (isShadow == 1)
+ {
+ if (isArray == 1)
+ // T read(uint2 coord, uint face, uint array, uint lod = 0) const
+ __intrinsic_asm "$0.read(vec<uint,2>(($1).xy), uint(($1).z), uint(($1).w))";
+ else
+ // T read(uint2 coord, uint face, uint lod = 0) const
+ __intrinsic_asm "$0.read(vec<uint,2>(($1).xy), uint(($1).z), uint(($1).w))";
+ }
+ else
+ {
+ if (isArray == 1)
+ // Tv read(uint2 coord, uint face, uint array, uint lod = 0) const
+ __intrinsic_asm "$0.read(vec<uint,2>(($1).xy), uint(($1).z), uint(($1).w))";
+ else
+ // Tv read(uint2 coord, uint face, uint lod = 0) const
+ __intrinsic_asm "$0.read(vec<uint,2>(($1).xy), uint(($1).z), uint(($1).w))";
+ }
+ break;
+ }
+ // TODO: This needs to be handled by the capability system
+ __intrinsic_asm "<invalid intrinsics>";
case glsl:
__intrinsic_asm "$ctexelFetch($0, ($1).$w1b, ($1).$w1e)$z";
case spirv:
@@ -1919,14 +2391,42 @@ extension __TextureImpl<T,Shape,isArray,1,sampleCount,0,isShadow,isCombined,form
__glsl_extension(GL_EXT_samplerless_texture_functions)
[__readNone]
[ForceInline]
- [require(cpp_glsl_hlsl_spirv, texture_sm_4_1_samplerless)]
+ [require(cpp_glsl_hlsl_metal_spirv, texture_sm_4_1_samplerless)]
T Load(vector<int, Shape.dimensions+isArray> location, int sampleIndex)
{
__target_switch
{
case cpp:
case hlsl:
- __intrinsic_asm ".Load";
+ __intrinsic_asm ".Load";
+ case metal:
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_2D):
+ if (isShadow == 1)
+ {
+ if (isArray == 1)
+ // Document seems to have a typo. `lod` must be `sample`.
+ // Tv read(uint2 coord, uint array, uint lod = 0) const
+ __intrinsic_asm "$0.read(($1).xy, ($1).z, uint($2))";
+ else
+ // T read(uint2 coord, uint sample) const
+ __intrinsic_asm "$0.read($1, uint($2))";
+ }
+ else
+ {
+ if (isArray == 1)
+ // Document seems to have a typo. `lod` must be `sample`.
+ // Tv read(uint2 coord, uint array, uint lod = 0) const
+ __intrinsic_asm "$0.read(($1).xy, ($1).z, uint($2))";
+ else
+ // Tv read(uint2 coord, uint sample) const
+ __intrinsic_asm "$0.read($1, uint($2))";
+ }
+ break;
+ }
+ // TODO: This needs to be handled by the capability system
+ __intrinsic_asm "<Not supported>";
case glsl:
__intrinsic_asm "$ctexelFetch($0, $1, ($2))$z";
case spirv:
@@ -2212,7 +2712,7 @@ extension __TextureImpl<T,Shape,isArray,1,sampleCount,$(access),isShadow, 0,form
{
[__readNone]
[ForceInline]
- [require(cpp_glsl_hlsl_spirv, texture_sm_4_1_compute_fragment)]
+ [require(cpp_glsl_hlsl_metal_spirv, texture_sm_4_1_compute_fragment)]
T Load(vector<int, Shape.dimensions+isArray> location, int sampleIndex)
{
__target_switch
@@ -2220,6 +2720,34 @@ extension __TextureImpl<T,Shape,isArray,1,sampleCount,$(access),isShadow, 0,form
case cpp:
case hlsl:
__intrinsic_asm ".Load";
+ case metal:
+ switch (Shape.flavor)
+ {
+ case $(SLANG_TEXTURE_2D):
+ if (isShadow == 1)
+ {
+ if (isArray == 1)
+ // The document seems to have a typo. `lod` must mean `sample`.
+ // Tv read(uint2 coord, uint array, uint lod = 0) const
+ __intrinsic_asm "$0.read(vec<uint,2>(($1).xy), uint(($1).z), $2)";
+ else
+ // T read(uint2 coord, uint sample) const
+ __intrinsic_asm "$0.read(vec<uint,2>(($1).xy), $2)";
+ }
+ else
+ {
+ if (isArray == 1)
+ // The document seems to have a typo. `lod` must mean `sample`.
+ // Tv read(uint2 coord, uint array, uint lod = 0) const
+ __intrinsic_asm "$0.read(vec<uint,2>(($1).xy), uint(($1).z), $2)";
+ else
+ // Tv read(uint2 coord, uint sample) const
+ __intrinsic_asm "$0.read(vec<uint,2>(($1).xy), $2)";
+ }
+ break;
+ }
+ // TODO: This needs to be handled by the capability system
+ __intrinsic_asm "<Not supported>";
case glsl:
__intrinsic_asm "$(glslIntrinsicMS)";
case spirv:
@@ -10135,13 +10663,14 @@ T NonUniformResourceIndex<T>(T value) { return value; }
// Normalize a vector
__generic<T : __BuiltinFloatingPointType, let N : int>
[__readNone]
-[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)]
+[require(cpp_cuda_glsl_hlsl_metal_spirv, sm_2_0_GLSL_140)]
vector<T,N> normalize(vector<T,N> x)
{
__target_switch
{
case glsl: __intrinsic_asm "normalize";
case hlsl: __intrinsic_asm "normalize";
+ case metal: __intrinsic_asm "normalize";
case spirv: return spirv_asm {
OpExtInst $$vector<T,N> result glsl450 Normalize $x
};
@@ -10152,13 +10681,14 @@ vector<T,N> normalize(vector<T,N> x)
__generic<T : __BuiltinFloatingPointType>
[__readNone]
-[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)]
+[require(cpp_cuda_glsl_hlsl_metal_spirv, sm_2_0_GLSL_140)]
T normalize(T x)
{
__target_switch
{
case glsl: __intrinsic_asm "normalize";
case hlsl: __intrinsic_asm "normalize";
+ case metal: __intrinsic_asm "normalize";
case spirv: return spirv_asm {
OpExtInst $$T result glsl450 Normalize $x
};
@@ -14024,8 +14554,8 @@ for (int aa = 0; aa < kBaseBufferAccessLevelCount; ++aa)
char const* glslTextureSizeFunc = (isReadOnly) ? "textureSize" : "imageSize";
char const* glslLoadFuncName = (isReadOnly) ? "texelFetch" : "imageLoad";
char const* spvLoadInstName = (isReadOnly) ? "OpImageFetch" : "OpImageRead";
- char const* requireToSetQuery = (isReadOnly) ? "[require(glsl_hlsl_spirv, texture_size)]" : "[require(glsl_hlsl_spirv, image_size)]";
- char const* requireToSet = (isReadOnly) ? "[require(glsl_hlsl_spirv, texture_sm_4_1)]" : "[require(glsl_hlsl_spirv, texture_sm_4_1_compute_fragment)]";
+ char const* requireToSetQuery = (isReadOnly) ? "[require(glsl_hlsl_metal_spirv, texture_size)]" : "[require(glsl_hlsl_metal_spirv, image_size)]";
+ char const* requireToSet = (isReadOnly) ? "[require(glsl_hlsl_metal_spirv, texture_sm_4_1)]" : "[require(glsl_hlsl_metal_spirv, texture_sm_4_1_compute_fragment)]";
}}}}
__generic<T, let format:int>
@@ -14039,6 +14569,7 @@ extension __TextureImpl<T, __ShapeBuffer, 0, 0, 0, $(aa), 0, 0, format>
{
case hlsl: __intrinsic_asm ".GetDimensions";
case glsl: __intrinsic_asm "($1 = $(glslTextureSizeFunc)($0))";
+ case metal: __intrinsic_asm "(*($1) = $0.get_width())";
case spirv:
dim = spirv_asm {
OpCapability ImageQuery;
@@ -14055,6 +14586,7 @@ extension __TextureImpl<T, __ShapeBuffer, 0, 0, 0, $(aa), 0, 0, format>
__target_switch
{
case hlsl: __intrinsic_asm ".Load";
+ case metal: __intrinsic_asm "$0.read(uint($1))";
case glsl: __intrinsic_asm "$(glslLoadFuncName)($0, $1)$z";
case spirv: return spirv_asm {
%sampled:__sampledType(T) = $(spvLoadInstName) $this $location;
@@ -14084,6 +14616,7 @@ ${{{{
{
case hlsl: __intrinsic_asm "($0)[$1] = $2";
case glsl: __intrinsic_asm "imageStore($0, int($1), $V2)";
+ case metal: __intrinsic_asm "$0.write($2, $1)";
case spirv: spirv_asm {
OpImageWrite $this $index __convertTexel(newValue);
};
diff --git a/source/slang/slang-capabilities.capdef b/source/slang/slang-capabilities.capdef
index 4d4fc7ccf..5c672d398 100644
--- a/source/slang/slang-capabilities.capdef
+++ b/source/slang/slang-capabilities.capdef
@@ -62,9 +62,9 @@ def spirv_1_6 : spirv_1_5;
alias spirv = spirv_1_0;
alias spirv_latest = spirv_1_6;
-alias any_target = hlsl | glsl | c | cpp | cuda | spirv;
-alias any_textual_target = hlsl | glsl | c | cpp | cuda;
-alias any_gfx_target = hlsl | glsl | spirv;
+alias any_target = hlsl | metal | glsl | c | cpp | cuda | spirv;
+alias any_textual_target = hlsl | metal | glsl | c | cpp | cuda;
+alias any_gfx_target = hlsl | metal | glsl | spirv;
alias any_cpp_target = cpp | cuda;
alias cpp_cuda = cpp | cuda;
@@ -75,6 +75,7 @@ alias cpp_cuda_glsl_hlsl_metal_spirv = cpp | cuda | glsl | hlsl | metal | spirv_
alias cpp_cuda_hlsl = cpp | cuda | hlsl;
alias cpp_glsl = cpp | glsl;
alias cpp_glsl_hlsl_spirv = cpp | glsl | hlsl | spirv_1_0;
+alias cpp_glsl_hlsl_metal_spirv = cpp | glsl | hlsl | metal | spirv_1_0;
alias cpp_hlsl = cpp | hlsl;
alias cuda_glsl_hlsl = cuda | glsl | hlsl;
alias cuda_glsl_hlsl_spirv = cuda | glsl | hlsl | spirv_1_0;
@@ -84,6 +85,7 @@ alias cuda_hlsl = cuda | hlsl;
alias cuda_hlsl_spirv = cuda | hlsl | spirv;
alias glsl_hlsl_spirv = glsl | hlsl | spirv;
alias glsl_hlsl_metal_spirv = glsl | hlsl | metal | spirv;
+alias glsl_metal_spirv = glsl | metal | spirv;
alias glsl_spirv = glsl | spirv;
alias hlsl_spirv = hlsl | spirv;
diff --git a/source/slang/slang-emit-metal.cpp b/source/slang/slang-emit-metal.cpp
index 1eb4b9abe..ff497a20f 100644
--- a/source/slang/slang-emit-metal.cpp
+++ b/source/slang/slang-emit-metal.cpp
@@ -963,6 +963,7 @@ void MetalSourceEmitter::emitFrontMatterImpl(TargetRequest*)
{
m_writer->emit("#include <metal_stdlib>\n");
m_writer->emit("#include <metal_math>\n");
+ m_writer->emit("#include <metal_texture>\n");
m_writer->emit("using namespace metal;\n");
}
diff --git a/source/slang/slang-stdlib-textures.cpp b/source/slang/slang-stdlib-textures.cpp
index f874c5da9..380ed1677 100644
--- a/source/slang/slang-stdlib-textures.cpp
+++ b/source/slang/slang-stdlib-textures.cpp
@@ -68,7 +68,8 @@ void TextureTypeInfo::writeFuncBody(
const String& cuda,
const String& spirvDefault,
const String& spirvRWDefault,
- const String& spirvCombined)
+ const String& spirvCombined,
+ const String& metal)
{
BraceScope funcScope{i, sb};
{
@@ -90,6 +91,11 @@ void TextureTypeInfo::writeFuncBody(
sb << i << "case cuda:\n";
sb << i << "__intrinsic_asm \"" << cuda << "\";\n";
}
+ if (metal.getLength())
+ {
+ sb << i << "case metal:\n";
+ sb << i << "__intrinsic_asm \"" << metal << "\";\n";
+ }
if(spirvDefault.getLength() && spirvCombined.getLength())
{
sb << i << "case spirv:\n";
@@ -127,14 +133,14 @@ void TextureTypeInfo::writeFuncWithSig(
const String& spirvRWDefault,
const String& spirvCombined,
const String& cuda,
+ const String& metal,
const ReadNoneMode readNoneMode)
{
if (readNoneMode == ReadNoneMode::Always)
sb << i << "[__readNone]\n";
- sb << i << "[__readNone]\n";
sb << i << "[ForceInline]\n";
sb << i << sig << "\n";
- writeFuncBody(funcName, glsl, cuda, spirvDefault, spirvRWDefault, spirvCombined);
+ writeFuncBody(funcName, glsl, cuda, spirvDefault, spirvRWDefault, spirvCombined, metal);
sb << "\n";
}
@@ -147,6 +153,7 @@ void TextureTypeInfo::writeFunc(
const String& spirvRWDefault,
const String& spirvCombined,
const String& cuda,
+ const String& metal,
const ReadNoneMode readNoneMode)
{
writeFuncWithSig(
@@ -157,6 +164,7 @@ void TextureTypeInfo::writeFunc(
spirvRWDefault,
spirvCombined,
cuda,
+ metal,
readNoneMode
);
}
@@ -184,27 +192,56 @@ void TextureTypeInfo::writeGetDimensionFunctions()
int sizeDimCount = 0;
StringBuilder params;
+ int paramCount = 0;
+
+ StringBuilder metal;
+ const char* metalMipLevel = "0";
+
if (includeMipInfo)
- params << "uint mipLevel, ";
+ {
+ ++paramCount;
+ params << "uint mipLevel,";
+
+ if (baseShape != SLANG_TEXTURE_1D)
+ metalMipLevel = "$1";
+ }
switch (baseShape)
{
case SLANG_TEXTURE_1D:
+ ++paramCount;
params << t << "width";
+ metal << "(*($" << String(paramCount) << ") = $0.get_width(" << String(metalMipLevel) << ")),";
+
sizeDimCount = 1;
break;
case SLANG_TEXTURE_2D:
case SLANG_TEXTURE_CUBE:
+ ++paramCount;
params << t << "width,";
+ metal << "(*($" << String(paramCount) << ") = $0.get_width(" << String(metalMipLevel) << ")),";
+
+ ++paramCount;
params << t << "height";
+ metal << "(*($" << String(paramCount) << ") = $0.get_height(" << String(metalMipLevel) << ")),";
+
sizeDimCount = 2;
break;
case SLANG_TEXTURE_3D:
+ ++paramCount;
params << t << "width,";
+ metal << "(*($" << String(paramCount) << ") = $0.get_width(" << String(metalMipLevel) << ")),";
+
+ ++paramCount;
params << t << "height,";
+ metal << "(*($" << String(paramCount) << ") = $0.get_height(" << String(metalMipLevel) << ")),";
+
+ ++paramCount;
params << t << "depth";
+ metal << "(*($" << String(paramCount) << ") = $0.get_depth(" << String(metalMipLevel) << ")),";
+
sizeDimCount = 3;
break;
@@ -215,18 +252,27 @@ void TextureTypeInfo::writeGetDimensionFunctions()
if (isArray)
{
+ ++sizeDimCount;
+ ++paramCount;
params << ", " << t << "elements";
- sizeDimCount++;
+ metal << "(*($" << String(paramCount) << ") = $0.get_array_size()),";
}
if (isMultisample)
{
+ ++paramCount;
params << ", " << t << "sampleCount";
+ metal << "(*($" << String(paramCount) << ") = $0.get_num_samples()),";
}
if (includeMipInfo)
+ {
+ ++paramCount;
params << ", " << t << "numberOfLevels";
+ metal << "(*($" << String(paramCount) << ") = $0.get_num_mip_levels()),";
+ }
+ metal.reduceLength(metal.getLength() - 1); // drop the last comma
StringBuilder glsl;
{
@@ -407,7 +453,7 @@ void TextureTypeInfo::writeGetDimensionFunctions()
sb << " __glsl_version(450)\n";
sb << " __glsl_extension(GL_EXT_samplerless_texture_functions)\n";
- sb << " [require(glsl_spirv, texture_sm_4_1)]\n";
+ sb << " [require(glsl_hlsl_metal_spirv, texture_sm_4_1)]\n";
writeFunc(
"void",
"GetDimensions",
@@ -417,6 +463,7 @@ void TextureTypeInfo::writeGetDimensionFunctions()
spirvRWDefault,
spirvCombined,
"",
+ metal,
ReadNoneMode::Always);
}
}
diff --git a/source/slang/slang-stdlib-textures.h b/source/slang/slang-stdlib-textures.h
index 6c07b1a46..a6166cc2d 100644
--- a/source/slang/slang-stdlib-textures.h
+++ b/source/slang/slang-stdlib-textures.h
@@ -68,8 +68,9 @@ public:
const String& cuda,
const String& spirvDefault,
const String& spirvRWDefault,
- const String& spirvCombined
- );
+ const String& spirvCombined,
+ const String& metal
+ );
void writeFuncWithSig(
const char* funcName,
const String& sig,
@@ -78,6 +79,7 @@ public:
const String& spirvRWDefault = String{},
const String& spirvCombined = String{},
const String& cuda = String{},
+ const String& metal = String{},
const ReadNoneMode readNoneMode = ReadNoneMode::Never
);
void writeFunc(
@@ -89,6 +91,7 @@ public:
const String& spirvRWDefault = String{},
const String& spirvCombined = String{},
const String& cuda = String{},
+ const String& metal = String{},
const ReadNoneMode readNoneMode = ReadNoneMode::Never
);
generated by cgit v1.2.3 (git 2.39.1) at 2026-06-01 19:13:09 +0000