summaryrefslogtreecommitdiffstats
diff options
context:
space:
mode:
authorJay Kwak <82421531+jkwak-work@users.noreply.github.com>2024-05-14 15:42:12 -0700
committerGitHub <noreply@github.com>2024-05-14 15:42:12 -0700
commitd76bed6c1b03e5d7ef19c947fdd5fcaf33b595f7 (patch)
treea5709a08298ead8f5fe2fb51f1b0f30c61505a3d
parent5ceb8569b1ac7898c437b0c47ad29a5d8a9f7d90 (diff)
Implement texture functions for Metal target (#4158)
* Impl texture APIs for Metal target This commit is to implement texture functions for Metal target. The following functions are implemented and tested. - GetDimensions() - CalculateLevelOfDetail() - CalculateLevelOfDetailUnclamped() - Sample() - SampleBias() - SampleLevel() - SampleCmp() - SampleCmpLevelZero() - Gather() - SampleGrad() - Load() Metal has limited support for the texture functions compared to HLSL. - LOD is not supported for 1D texture, - Depth textures are limited to 2D, 2DArray, Cube and CubeArray textures. - "Offset" variants are limited to 2D, 2DArray, 2D-Depth, 2DArray-Depth and 3D textures. The functions that cannot be implemented for Metal should properly be handled by the capability system later. * Fix the failing test, multi-file.hlsl I am not sure why this change is needed. * Fix compile errors on macOS 2nd try * Remove a typo character to fix the compile error * Trivial clean up * Remove `as_type` where it was intended as static_cast * Use a simpler sytax for __intrinsic_asm * Trivial clean up * Remove TEST_AFTER_FIXING_CAPABILITY_PROBLEM after fixing normalize * Fix the failing test properly * Fix an incorrect setup of Depth-cube texture --------- Co-authored-by: Yong He <yonghe@outlook.com>
Diffstat
-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
-rw-r--r--tests/metal/texture.slang642
6 files changed, 1292 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
);
diff --git a/tests/metal/texture.slang b/tests/metal/texture.slang
new file mode 100644
index 000000000..fc9347bea
--- /dev/null
+++ b/tests/metal/texture.slang
@@ -0,0 +1,642 @@
+//TEST:SIMPLE(filecheck=METAL): -stage compute -entry computeMain -target metal -DEMIT_SOURCE
+//TEST:SIMPLE(filecheck=METALLIB): -stage compute -entry computeMain -target metallib -DEMIT_SOURCE -DMETALLIB
+//TEST:SIMPLE(filecheck=HLSL): -stage compute -entry computeMain -target hlsl -DEMIT_SOURCE
+//TEST:SIMPLE(filecheck=SPIR): -stage compute -entry computeMain -target spirv -emit-spirv-directly -DEMIT_SOURCE
+
+//TEST(compute):COMPARE_COMPUTE_EX(filecheck-buffer=FUNCTIONAL):-slang -compute -dx12 -profile cs_6_6 -use-dxil -shaderobj -output-using-type
+//TEST(compute, vulkan):COMPARE_COMPUTE_EX(filecheck-buffer=FUNCTIONAL):-vk -emit-spirv-directly -compute -shaderobj -output-using-type -render-feature hardware-device
+
+#if !defined(METALLIB)
+ // Metal doesn't support some features, and we need a new test that
+ // can check if the capability system can error them out properly.
+ #define NEED_TO_TEST_FOR_METAL_CAPABILITY 1
+#endif
+
+#if defined(EMIT_SOURCE)
+ // It appears that Slang-test doesn't initialize the depth cube texture
+ // properly.
+ #define TEST_WHEN_DEPTH_CUBE_WORKS
+#endif
+
+//TEST_INPUT: ubuffer(data=[0], stride=4):out,name outputBuffer
+RWStructuredBuffer<int> outputBuffer;
+
+//TEST_INPUT: Texture1D(size=4, content = one):name t1D
+Texture1D<float> t1D;
+//TEST_INPUT: Texture2D(size=4, content = one):name t2D
+Texture2D<float> t2D;
+//TEST_INPUT: Texture3D(size=4, content = one):name t3D
+Texture3D<float> t3D;
+//TEST_INPUT: TextureCube(size=4, content = one):name tCube
+TextureCube<float> tCube;
+
+//TEST_INPUT: Texture1D(size=4, content = one, arrayLength=2):name t1DArray
+Texture1DArray<float> t1DArray;
+//TEST_INPUT: Texture2D(size=4, content = one, arrayLength=2):name t2DArray
+Texture2DArray<float> t2DArray;
+//TEST_INPUT: TextureCube(size=4, content = one, arrayLength=2):name tCubeArray
+TextureCubeArray<float> tCubeArray;
+
+// Metal doc says "For depth texture types, T must be float."
+__generic<T : __BuiltinType, let sampleCount:int=0, let format:int=0>
+typealias depth2d = __TextureImpl<
+ T,
+ __Shape2D,
+ 0, // isArray
+ 0, // isMS
+ sampleCount,
+ 0, // access
+ 1, // isShadow
+ 0, // isCombined
+ format
+>;
+
+__generic<T : __BuiltinType, let sampleCount:int=0, let format:int=0>
+typealias depth2d_array = __TextureImpl<
+ T,
+ __Shape2D,
+ 1, // isArray
+ 0, // isMS
+ sampleCount,
+ 0, // access
+ 1, // isShadow
+ 0, // isCombined
+ format
+>;
+
+__generic<T : __BuiltinType, let sampleCount:int=0, let format:int=0>
+typealias depthcube = __TextureImpl<
+ T,
+ __ShapeCube,
+ 0, // isArray
+ 0, // isMS
+ sampleCount,
+ 0, // access
+ 1, // isShadow
+ 0, // isCombined
+ format
+>;
+
+__generic<T : __BuiltinType, let sampleCount:int=0, let format:int=0>
+typealias depthcube_array = __TextureImpl<
+ T,
+ __ShapeCube,
+ 1, // isArray
+ 0, // isMS
+ sampleCount,
+ 0, // access
+ 1, // isShadow
+ 0, // isCombined
+ format
+>;
+
+//TEST_INPUT: Texture2D(size=4, content = one):name d2D
+depth2d<float> d2D;
+//TEST_INPUT: Texture2D(size=4, content = one):name dCube
+depthcube<float> dCube;
+//TEST_INPUT: Texture2D(size=4, content = one, arrayLength=2):name d2DArray
+depth2d_array<float> d2DArray;
+//TEST_INPUT: TextureCube(size=4, content = one, arrayLength=2):name dCubeArray
+depthcube_array<float> dCubeArray;
+
+//TEST_INPUT: Sampler:name samplerState
+SamplerState samplerState;
+//TEST_INPUT: Sampler:name shadowSampler
+SamplerComparisonState shadowSampler;
+
+
+bool TEST_texture_float()
+{
+ // Metal textures support `Tv` types, which "denotes a 4-component vector
+ // type based on the templated type <T> for declaring the texture type:
+ // - If T is float, Tv is float4.
+ // - If T is half, Tv is half4.
+ // - If T is int, Tv is int4.
+ // - If T is uint, Tv is uint4.
+ // - If T is short, Tv is short4.
+ // - If T is ushort, Tv is ushort4."
+ typealias Tv = vector<float,4>;
+
+ float u = 0;
+ float u2 = 0.5;
+ constexpr const float ddx = 0.0f;
+ constexpr const float ddy = 0.0f;
+
+ uint width = 0, height = 0, depth = 0;
+ float fwidth = 0.0f, fheight = 0.0f, fdepth = 0.0f;
+ uint numLevels = 0, elements = 0, sampleCount = 0;
+ float fnumLevels = 0.0f, felements = 0.0f;
+
+ bool voidResult = true;
+
+ // ======================
+ // void GetDimensions()
+ // ======================
+
+ // METAL: .get_width(
+ // METALLIB: call {{.*}}.get_width_texture_1d(
+ t1D.GetDimensions(width);
+ voidResult = voidResult && (uint(4) == width);
+
+ // METAL: .get_width({{.*}}.get_num_mip_levels()
+ // METALLIB: call {{.*}}.get_num_mip_levels_texture_1d(
+ t1D.GetDimensions(0, width, numLevels);
+ voidResult = voidResult && (uint(4) == width);
+ voidResult = voidResult && (uint(3) == numLevels);
+
+ // METAL: .get_width({{.*}}.get_height(
+ // METALLIB: call {{.*}}.get_height_texture_2d(
+ t2D.GetDimensions(width, height);
+ voidResult = voidResult && (uint(4) == width);
+ voidResult = voidResult && (uint(4) == height);
+
+ // METAL: .get_width({{.*}}.get_height({{.*}}.get_num_mip_levels()
+ // METALLIB: call {{.*}}.get_num_mip_levels_texture_2d(
+ t2D.GetDimensions(0, width, height, numLevels);
+ voidResult = voidResult && (uint(4) == width);
+ voidResult = voidResult && (uint(4) == height);
+ voidResult = voidResult && (uint(3) == numLevels);
+
+ // METAL: .get_width({{.*}}.get_height({{.*}}.get_depth(
+ // METALLIB: call {{.*}}.get_depth_texture_3d(
+ t3D.GetDimensions(width, height, depth);
+ voidResult = voidResult && (uint(4) == width);
+ voidResult = voidResult && (uint(4) == height);
+ voidResult = voidResult && (uint(4) == depth);
+
+ // METAL: .get_width({{.*}}.get_height({{.*}}.get_depth({{.*}}.get_num_mip_levels()
+ // METALLIB: call {{.*}}.get_num_mip_levels_texture_3d(
+ t3D.GetDimensions(0, width, height, depth, numLevels);
+ voidResult = voidResult && (uint(4) == width);
+ voidResult = voidResult && (uint(4) == height);
+ voidResult = voidResult && (uint(4) == depth);
+ voidResult = voidResult && (uint(3) == numLevels);
+
+ // METAL: .get_width({{.*}}.get_height({{.*}}
+ // METALLIB: call {{.*}}.get_height_texture_cube(
+ tCube.GetDimensions(width, height);
+ voidResult = voidResult && (uint(4) == width);
+ voidResult = voidResult && (uint(4) == height);
+
+ // METAL: .get_width({{.*}}.get_height({{.*}}.get_num_mip_levels()
+ // METALLIB: call {{.*}}.get_num_mip_levels_texture_cube(
+ tCube.GetDimensions(0, width, height, numLevels);
+ voidResult = voidResult && (uint(4) == width);
+ voidResult = voidResult && (uint(4) == height);
+ voidResult = voidResult && (uint(3) == numLevels);
+
+ // METAL: .get_width({{.*}}.get_array_size(
+ // METALLIB: call {{.*}}.get_array_size_texture_1d_array(
+ t1DArray.GetDimensions(width, elements);
+ voidResult = voidResult && (uint(4) == width);
+ voidResult = voidResult && (uint(2) == elements);
+
+ // METAL: .get_width({{.*}}.get_num_mip_levels(
+ // METALLIB: call {{.*}}.get_num_mip_levels_texture_1d_array(
+ t1DArray.GetDimensions(0, width, elements, numLevels);
+ voidResult = voidResult && (uint(4) == width);
+ voidResult = voidResult && (uint(2) == elements);
+ voidResult = voidResult && (uint(3) == numLevels);
+
+ // METAL: .get_width({{.*}}.get_height({{.*}}.get_array_size(
+ // METALLIB: call {{.*}}.get_array_size_texture_2d_array(
+ t2DArray.GetDimensions(width, height, elements);
+ voidResult = voidResult && (uint(4) == width);
+ voidResult = voidResult && (uint(4) == height);
+ voidResult = voidResult && (uint(2) == elements);
+
+ // METAL: .get_width({{.*}}.get_height({{.*}}.get_num_mip_levels(
+ // METALLIB: call {{.*}}.get_num_mip_levels_texture_2d_array(
+ t2DArray.GetDimensions(0, width, height, elements, numLevels);
+ voidResult = voidResult && (uint(4) == width);
+ voidResult = voidResult && (uint(4) == height);
+ voidResult = voidResult && (uint(2) == elements);
+ voidResult = voidResult && (uint(3) == numLevels);
+
+ // METAL: .get_width({{.*}}.get_height({{.*}}.get_array_size(
+ // METALLIB: call {{.*}}.get_array_size_texture_cube_array(
+ tCubeArray.GetDimensions(width, height, elements);
+ voidResult = voidResult && (uint(4) == width);
+ voidResult = voidResult && (uint(4) == height);
+ voidResult = voidResult && (uint(2) == elements);
+
+ // METAL: .get_width({{.*}}.get_height({{.*}}.get_num_mip_levels(
+ // METALLIB: call {{.*}}.get_num_mip_levels_texture_cube_array(
+ tCubeArray.GetDimensions(0, width, height, elements, numLevels);
+ voidResult = voidResult && (uint(4) == width);
+ voidResult = voidResult && (uint(4) == height);
+ voidResult = voidResult && (uint(2) == elements);
+ voidResult = voidResult && (uint(3) == numLevels);
+
+ bool result = voidResult
+ // ===============================
+ // float CalculateLevelOfDetail()
+ // ===============================
+
+#if defined(NEED_TO_TEST_FOR_METAL_CAPABILITY)
+ // Metal doesn't support LOD for 1D texture
+ && float(0) == t1D.CalculateLevelOfDetail(samplerState, u)
+#endif
+
+ // METAL: t2D{{.*}}.calculate_clamped_lod({{.*}}
+ // METALLIB: call {{.*}}.calculate_clamped_lod_texture_2d(
+ && float(0) == t2D.CalculateLevelOfDetail(samplerState, float2(u, u))
+
+ // METAL: t3D{{.*}}.calculate_clamped_lod({{.*}}
+ // METALLIB: call {{.*}}.calculate_clamped_lod_texture_3d(
+ && float(0) == t3D.CalculateLevelOfDetail(samplerState, float3(u, u, u))
+
+ // METAL: tCube{{.*}}.calculate_clamped_lod({{.*}}
+ // METALLIB: call {{.*}}.calculate_clamped_lod_texture_cube(
+ && float(0) == tCube.CalculateLevelOfDetail(samplerState, float3(u, u, u))
+
+ // METAL: t2DArray{{.*}}.calculate_clamped_lod({{.*}}
+ // METALLIB: call {{.*}}.calculate_clamped_lod_texture_2d_array(
+ && float(0) == t2DArray.CalculateLevelOfDetail(samplerState, float2(u, u))
+
+ // METAL: tCubeArray{{.*}}.calculate_clamped_lod({{.*}}
+ // METALLIB: call {{.*}}.calculate_clamped_lod_texture_cube_array(
+ && float(0) == tCubeArray.CalculateLevelOfDetail(samplerState, float3(u, u, u))
+
+ // ========================================
+ // float CalculateLevelOfDetailUnclamped()
+ // ========================================
+
+#if defined(NEED_TO_TEST_FOR_METAL_CAPABILITY)
+ // Metal doesn't support LOD for 1D texture
+ && float(0) >= t1D.CalculateLevelOfDetailUnclamped(samplerState, u)
+#endif
+
+ // METAL: t2D{{.*}}.calculate_unclamped_lod({{.*}}
+ // METALLIB: call {{.*}}.calculate_unclamped_lod_texture_2d(
+ && float(0) >= t2D.CalculateLevelOfDetailUnclamped(samplerState, float2(u, u))
+
+ // METAL: t3D{{.*}}.calculate_unclamped_lod({{.*}}
+ // METALLIB: call {{.*}}.calculate_unclamped_lod_texture_3d(
+ && float(0) >= t3D.CalculateLevelOfDetailUnclamped(samplerState, float3(u, u, u))
+
+ // METAL: tCube{{.*}}.calculate_unclamped_lod({{.*}}
+ // METALLIB: call {{.*}}.calculate_unclamped_lod_texture_cube(
+ && float(0) >= tCube.CalculateLevelOfDetailUnclamped(samplerState, normalize(float3(u, 1 - u, u)))
+
+ // METAL: t2DArray{{.*}}.calculate_unclamped_lod({{.*}}
+ // METALLIB: call {{.*}}.calculate_unclamped_lod_texture_2d_array(
+ && float(0) >= t2DArray.CalculateLevelOfDetailUnclamped(samplerState, float2(u, u))
+
+ // METAL: tCubeArray{{.*}}.calculate_unclamped_lod({{.*}}
+ // METALLIB: call {{.*}}.calculate_unclamped_lod_texture_cube_array(
+ && float(0) >= tCubeArray.CalculateLevelOfDetailUnclamped(samplerState, normalize(float3(u, 1 - u, u)))
+
+ // ===========
+ // T Sample()
+ // ===========
+
+ // METAL: t1D{{.*}}.sample(
+ // METALLIB: call {{.*}}.sample_texture_1d.v4f32(
+ && all(Tv(1) == t1D.Sample(samplerState, u))
+
+ // METAL: t2D{{.*}}.sample({{.*}}
+ // METALLIB: call {{.*}}.sample_texture_2d.v4f32(
+ && all(Tv(1) == t2D.Sample(samplerState, float2(u, u)))
+
+ // METAL: t3D{{.*}}.sample({{.*}}
+ // METALLIB: call {{.*}}.sample_texture_3d.v4f32(
+ && all(Tv(1) == t3D.Sample(samplerState, float3(u, u, u)))
+
+ // METAL: tCube{{.*}}.sample({{.*}}
+ // METALLIB: call {{.*}}.sample_texture_cube.v4f32(
+ && all(Tv(1) == tCube.Sample(samplerState, normalize(float3(u, 1 - u, u))))
+
+ // METAL: t1DArray{{.*}}.sample(
+ // METALLIB: call {{.*}}.sample_texture_1d_array.v4f32(
+ && all(Tv(1) == t1DArray.Sample(samplerState, float2(u, 0)))
+
+ // METAL: t2DArray{{.*}}.sample({{.*}}
+ // METALLIB: call {{.*}}.sample_texture_2d_array.v4f32(
+ && all(Tv(1) == t2DArray.Sample(samplerState, float3(u, u, 0)))
+
+ // METAL: tCubeArray{{.*}}.sample({{.*}}
+ // METALLIB: call {{.*}}.sample_texture_cube_array.v4f32(
+ && all(Tv(1) == tCubeArray.Sample(samplerState, normalize(float4(u, 1 - u, u, 0))))
+
+ // Offset variant
+
+#if defined(NEED_TO_TEST_FOR_METAL_CAPABILITY)
+ // Metal doesn't support Offset for 1D and Cube texture
+ && all(Tv(1) == t1D.Sample(samplerState, u2, 1))
+ && all(Tv(1) == t1DArray.Sample(samplerState, float2(u2, 0), 1))
+#endif
+
+ // METAL: t2D{{.*}}.sample({{.*}}
+ // METALLIB: call {{.*}}.sample_texture_2d.v4f32(
+ && all(Tv(1) == t2D.Sample(samplerState, float2(u, u), int2(1, 1)))
+
+ // METAL: t3D{{.*}}.sample({{.*}}
+ // METALLIB: call {{.*}}.sample_texture_3d.v4f32(
+ && all(Tv(1) == t3D.Sample(samplerState, float3(u, u, u), int3(1, 1, 1)))
+
+ // METAL: t2DArray{{.*}}.sample({{.*}}
+ // METALLIB: call {{.*}}.sample_texture_2d_array.v4f32(
+ && all(Tv(1) == t2DArray.Sample(samplerState, float3(u, u, 0), int2(1, 1)))
+
+ // Clamp variant
+
+#if defined(NEED_TO_TEST_FOR_METAL_CAPABILITY)
+ // Metal doesn't support Offset for 1D and Cube texture
+ && all(Tv(1) == t1D.Sample(samplerState, u2, 1, float(1)))
+ && all(Tv(1) == t1DArray.Sample(samplerState, float2(u2, 0), 1, float(1)))
+#endif
+
+ // METAL: t2D{{.*}}.sample({{.*}} min_lod_clamp(
+ // METALLIB: call {{.*}}.sample_texture_2d.v4f32(
+ && all(Tv(1) == t2D.Sample(samplerState, float2(u, u), int2(1, 1), float(1)))
+
+ // METAL: t3D{{.*}}.sample({{.*}} min_lod_clamp(
+ // METALLIB: call {{.*}}.sample_texture_3d.v4f32(
+ && all(Tv(1) == t3D.Sample(samplerState, float3(u, u, u), int3(1, 1, 1), float(1)))
+
+ // METAL: t2DArray{{.*}}.sample({{.*}} min_lod_clamp(
+ // METALLIB: call {{.*}}.sample_texture_2d_array.v4f32(
+ && all(Tv(1) == t2DArray.Sample(samplerState, float3(u, u, 0), int2(1, 1), float(1)))
+
+ // ===============
+ // T SampleBias()
+ // ===============
+
+#if defined(NEED_TO_TEST_FOR_METAL_CAPABILITY)
+ // Metal doesn't support Bias for 1D texture
+ && all(Tv(1) == t1D.SampleBias(samplerState, u, float(-1)))
+ && all(Tv(1) == t1DArray.SampleBias(samplerState, float2(u, 0), float(-1)))
+#endif
+
+ // METAL: t2D{{.*}}.sample({{.*}}
+ // METALLIB: call {{.*}}.sample_texture_2d.v4f32(
+ && all(Tv(1) == t2D.SampleBias(samplerState, float2(u, u), float(-1)))
+
+ // METAL: t3D{{.*}}.sample({{.*}}
+ // METALLIB: call {{.*}}.sample_texture_3d.v4f32(
+ && all(Tv(1) == t3D.SampleBias(samplerState, float3(u, u, u), float(-1)))
+
+ // METAL: tCube{{.*}}.sample({{.*}}
+ // METALLIB: call {{.*}}.sample_texture_cube.v4f32(
+ && all(Tv(1) == tCube.SampleBias(samplerState, normalize(float3(u, 1 - u, u)), float(-1)))
+
+ // METAL: t2DArray{{.*}}.sample({{.*}}
+ // METALLIB: call {{.*}}.sample_texture_2d_array.v4f32(
+ && all(Tv(1) == t2DArray.SampleBias(samplerState, float3(u, u, 0), float(-1)))
+
+ // METAL: tCubeArray{{.*}}.sample({{.*}}
+ // METALLIB: call {{.*}}.sample_texture_cube_array.v4f32(
+ && all(Tv(1) == tCubeArray.SampleBias(samplerState, normalize(float4(u, 1 - u, u, 0)), float(-1)))
+
+ // Offset variant
+
+#if defined(NEED_TO_TEST_FOR_METAL_CAPABILITY)
+ // Metal doesn't support Offset for 1D and Cube texture
+ && all(Tv(1) == t1D.SampleBias(samplerState, u2, float(-1), 1))
+ && all(Tv(1) == t1DArray.SampleBias(samplerState, float2(u2, 0), float(-1), 1))
+#endif
+
+ // METAL: t2D{{.*}}.sample({{.*}}
+ // METALLIB: call {{.*}}.sample_texture_2d.v4f32(
+ && all(Tv(1) == t2D.SampleBias(samplerState, float2(u, u), float(-1), int2(1, 1)))
+
+ // METAL: t3D{{.*}}.sample({{.*}}
+ // METALLIB: call {{.*}}.sample_texture_3d.v4f32(
+ && all(Tv(1) == t3D.SampleBias(samplerState, float3(u, u, u), float(-1), int3(1, 1, 1)))
+
+ // METAL: t2DArray{{.*}}.sample({{.*}}
+ // METALLIB: call {{.*}}.sample_texture_2d_array.v4f32(
+ && all(Tv(1) == t2DArray.SampleBias(samplerState, float3(u, u, 0), float(-1), int2(1, 1)))
+
+ // ===================================
+ // T SampleLevel()
+ // ===================================
+
+#if defined(NEED_TO_TEST_FOR_METAL_CAPABILITY)
+ // Metal doesn't support LOD for 1D texture
+ && all(Tv(1) == t1D.SampleLevel(samplerState, u, 0))
+ && all(Tv(1) == t1DArray.SampleLevel(samplerState, float2(u, 0), 0))
+#endif
+
+ // METAL: t2D{{.*}}.sample({{.*}} level(
+ // METALLIB: call {{.*}}.sample_texture_2d.v4f32(
+ && all(Tv(1) == t2D.SampleLevel(samplerState, float2(u, u), 0))
+
+ // METAL: t3D{{.*}}.sample({{.*}} level(
+ // METALLIB: call {{.*}}.sample_texture_3d.v4f32(
+ && all(Tv(1) == t3D.SampleLevel(samplerState, float3(u, u, u), 0))
+
+ // METAL: tCube{{.*}}.sample({{.*}} level(
+ // METALLIB: call {{.*}}.sample_texture_cube.v4f32(
+ && all(Tv(1) == tCube.SampleLevel(samplerState, normalize(float3(u, 1 - u, u)), 0))
+
+ // METAL: t2DArray{{.*}}.sample({{.*}} level(
+ // METALLIB: call {{.*}}.sample_texture_2d_array.v4f32(
+ && all(Tv(1) == t2DArray.SampleLevel(samplerState, float3(u, u, 0), 0))
+
+ // METAL: tCubeArray{{.*}}.sample({{.*}} level(
+ // METALLIB: call {{.*}}.sample_texture_cube_array.v4f32(
+ && all(Tv(1) == tCubeArray.SampleLevel(samplerState, float4(normalize(float3(u, 1 - u, u)), 0), 0))
+
+ // Offset variant
+
+#if defined(NEED_TO_TEST_FOR_METAL_CAPABILITY)
+ // Metal doesn't support LOD for 1D texture
+ && all(Tv(1) == t1D.SampleLevel(samplerState, u2, 0, 1))
+ && all(Tv(1) == t1DArray.SampleLevel(samplerState, float2(u2, 0), 0, 1))
+#endif
+
+ // METAL: t2D{{.*}}.sample({{.*}} level(
+ // METALLIB: call {{.*}}.sample_texture_2d.v4f32(
+ && all(Tv(1) == t2D.SampleLevel(samplerState, float2(u, u), 0, int2(1, 1)))
+
+ // METAL: t3D{{.*}}.sample({{.*}} level(
+ // METALLIB: call {{.*}}.sample_texture_3d.v4f32(
+ && all(Tv(1) == t3D.SampleLevel(samplerState, float3(u, u, u), 0, int3(1, 1, 1)))
+
+ // METAL: t2DArray{{.*}}.sample({{.*}} level(
+ // METALLIB: call {{.*}}.sample_texture_2d_array.v4f32(
+ && all(Tv(1) == t2DArray.SampleLevel(samplerState, float3(u, u, 0), 0, int2(1, 1)))
+
+ // ==================
+ // float SampleCmp()
+ // ==================
+
+ // METAL: d2D{{.*}}.sample_compare(
+ // METALLIB: call {{.*}}.sample_compare_depth_2d.f32(
+ && float(1) == d2D.SampleCmp(shadowSampler, float2(u, u), 0)
+
+ // METAL: d2DArray{{.*}}.sample_compare(
+ // METALLIB: call {{.*}}.sample_compare_depth_2d_array.f32(
+ && float(1) == d2DArray.SampleCmp(shadowSampler, normalize(float3(u, 1 - u, u)), 0)
+
+#if defined(TEST_WHEN_DEPTH_CUBE_WORKS)
+ // METAL: dCube{{.*}}.sample_compare(
+ // METALLIB: call {{.*}}.sample_compare_depth_cube.f32(
+ && float(1) == dCube.SampleCmp(shadowSampler, normalize(float3(u, 1 - u, u)), 0)
+
+ // METAL: dCubeArray{{.*}}.sample_compare(
+ // METALLIB: call {{.*}}.sample_compare_depth_cube_array.f32(
+ && float(1) == dCubeArray.SampleCmp(shadowSampler, normalize(float4(u, 1-u, u, u)), 0)
+#endif
+
+ // Offset variant
+
+ // METAL: d2D{{.*}}.sample_compare(
+ // METALLIB: call {{.*}}.sample_compare_depth_2d.f32(
+ && float(1) == d2D.SampleCmp(shadowSampler, float2(u2, u), 0, int2(0, 0))
+
+ // ===================================
+ // float SampleCmpLevelZero()
+ // ===================================
+
+ // METAL: d2D{{.*}}.sample_compare(
+ // METALLIB: call {{.*}}.sample_compare_depth_2d.f32(
+ && float(1) == d2D.SampleCmpLevelZero(shadowSampler, float2(u, u), 0)
+
+ // METAL: d2DArray{{.*}}.sample_compare(
+ // METALLIB: call {{.*}}.sample_compare_depth_2d_array.f32(
+ && float(1) == d2DArray.SampleCmpLevelZero(shadowSampler, normalize(float3(u, 1 - u, u)), 0)
+
+#if defined(TEST_WHEN_DEPTH_CUBE_WORKS)
+ // METAL: dCube{{.*}}.sample_compare(
+ // METALLIB: call {{.*}}.sample_compare_depth_cube.f32(
+ && float(1) == dCube.SampleCmpLevelZero(shadowSampler, normalize(float3(u, 1 - u, u)), 0)
+
+ // METAL: dCubeArray{{.*}}.sample_compare(
+ // METALLIB: call {{.*}}.sample_compare_depth_cube_array.f32(
+ && float(1) == dCubeArray.SampleCmpLevelZero(shadowSampler, normalize(float4(u, 1-u, u, u)), 0)
+#endif
+
+ // Offset variant
+
+ // METAL: d2D{{.*}}.sample_compare(
+ // METALLIB: call {{.*}}.sample_compare_depth_2d.f32(
+ && float(1) == d2D.SampleCmpLevelZero(shadowSampler, float2(u2, u), 0, int2(0, 0))
+
+ // ==================================
+ // vector<T,4> Gather()
+ // ==================================
+
+ // METAL: t2D{{.*}}.gather(
+ // METALLIB: call {{.*}}.gather_texture_2d.v4f32(
+ && all(Tv(1) == t2D.Gather(samplerState, float2(u, u)))
+
+ // METAL: tCube{{.*}}.gather(
+ // METALLIB: call {{.*}}.gather_texture_cube.v4f32(
+ && all(Tv(1) == tCube.Gather(samplerState, normalize(float3(u, 1 - u, u))))
+
+ // METAL: t2DArray{{.*}}.gather(
+ // METALLIB: call {{.*}}.gather_texture_2d_array.v4f32(
+ && all(Tv(1) == t2DArray.Gather(samplerState, float3(u, u, 0)))
+
+ // METAL: tCubeArray{{.*}}.gather(
+ // METALLIB: call {{.*}}.gather_texture_cube_array.v4f32(
+ && all(Tv(1) == tCubeArray.Gather(samplerState, float4(normalize(float3(u, 1 - u, u)), 0)))
+
+ // Offset variant
+
+ // METAL: t2D{{.*}}.gather(
+ // METALLIB: call {{.*}}.gather_texture_2d.v4f32(
+ && all(Tv(1) == t2D.Gather(samplerState, float2(u2, u), int2(0, 0)))
+
+ // METAL: t2DArray{{.*}}.gather(
+ // METALLIB: call {{.*}}.gather_texture_2d_array.v4f32(
+ && all(Tv(1) == t2DArray.Gather(samplerState, float3(u2, u, 0), int2(0, 0)))
+
+ // =====================================
+ // T SampleGrad()
+ // =====================================
+
+#if defined(NEED_TO_TEST_FOR_METAL_CAPABILITY)
+ // Metal doesn't support LOD for 1D texture
+ && all(Tv(1) == t1D.SampleGrad(samplerState, 0.0f, ddx, ddy))
+ && all(Tv(1) == t1DArray.SampleGrad(samplerState, float2(0.0f, 0.0f), ddx, ddy))
+ && all(Tv(1) == t1D.SampleGrad(samplerState, 0.0f, ddx, ddy, 0))
+ && all(Tv(1) == t1DArray.SampleGrad(samplerState, float2(0.0f, 0.0f), ddx, ddy, 0))
+#endif
+
+ // METAL: t2D{{.*}}.sample(
+ // METALLIB: call {{.*}}.sample_texture_2d_grad.v4f32(
+ && all(Tv(1) == t2D.SampleGrad(samplerState, float2(u, u), float2(ddx, ddx), float2(ddy, ddy)))
+
+ // METAL: t3D{{.*}}.sample(
+ // METALLIB: call {{.*}}.sample_texture_3d_grad.v4f32(
+ && all(Tv(1) == t3D.SampleGrad(samplerState, float3(u, u, u), float3(ddx, ddx, ddx), float3(ddy, ddy, ddy)))
+
+ // METAL: tCube{{.*}}.sample(
+ // METALLIB: call {{.*}}.sample_texture_cube_grad.v4f32(
+ && all(Tv(1) == tCube.SampleGrad(samplerState, normalize(float3(u, 1 - u, u)), float3(ddx, ddx, ddx), float3(ddy, ddy, ddy)))
+
+ // METAL: t2DArray{{.*}}.sample(
+ // METALLIB: call {{.*}}.sample_texture_2d_array_grad.v4f32(
+ && all(Tv(1) == t2DArray.SampleGrad(samplerState, float3(u, u, 0.0f), float2(ddx, ddx), float2(ddy, ddy)))
+
+ // Offset variant
+
+ // METAL: t2D{{.*}}.sample(
+ // METALLIB: call {{.*}}.sample_texture_2d_grad.v4f32(
+ && all(Tv(1) == t2D.SampleGrad(samplerState, float2(u2, u), float2(ddx, ddx), float2(ddy, ddy), int2(0, 0)))
+
+ // METAL: t3D{{.*}}.sample(
+ // METALLIB: call {{.*}}.sample_texture_3d_grad.v4f32(
+ && all(Tv(1) == t3D.SampleGrad(samplerState, float3(u2, u, u), float3(ddx, ddx, ddx), float3(ddy, ddy, ddy), int3(0, 0, 0)))
+
+ // METAL: t2DArray{{.*}}.sample(
+ // METALLIB: call {{.*}}.sample_texture_2d_array_grad.v4f32(
+ && all(Tv(1) == t2DArray.SampleGrad(samplerState, float3(u2, u, 0.0f), float2(ddx, ddx), float2(ddy, ddy), int2(0, 0)))
+
+ // ===================
+ // T Load()
+ // ===================
+
+ // METAL: t1D{{.*}}.read(
+ // METALLIB: call {{.*}}.read_texture_1d.v4f32(
+ && all(Tv(1) == t1D.Load(int2(0, 0)))
+
+ // METAL: t2D{{.*}}.read(
+ // METALLIB: call {{.*}}.read_texture_2d.v4f32(
+ && all(Tv(1) == t2D.Load(int3(0, 0, 0)))
+
+ // METAL: t3D{{.*}}.read(
+ // METALLIB: call {{.*}}.read_texture_3d.v4f32(
+ && all(Tv(1) == t3D.Load(int4(0, 0, 0, 0)))
+
+ // METAL: t1DArray{{.*}}.read(
+ // METALLIB: call {{.*}}.read_texture_1d_array.v4f32(
+ && all(Tv(1) == t1DArray.Load(int3(0, 0, 0)))
+
+ // METAL: t2DArray{{.*}}.read(
+ // METALLIB: call {{.*}}.read_texture_2d_array.v4f32(
+ && all(Tv(1) == t2DArray.Load(int4(0, 0, 0, 0)))
+
+ // Offset variant
+
+#if defined(NEED_TO_TEST_FOR_METAL_CAPABILITY)
+ // Metal doesn't support offset variants for Load
+ && all(Tv(1) == t1D.Load(int2(0, 0), 0))
+ && all(Tv(1) == t2D.Load(int3(0, 0, 0), int2(0,0)))
+ && all(Tv(1) == t3D.Load(int4(0, 0, 0, 0), int3(0, 0, 0)))
+ && all(Tv(1) == t1DArray.Load(int3(0, 0, 0), 0))
+ && all(Tv(1) == t2DArray.Load(int4(0, 0, 0, 0), int2(0, 0)))
+#endif
+ ;
+
+ return result;
+}
+
+
+[numthreads(1, 1, 1)]
+void computeMain()
+{
+ // SPIR: OpEntryPoint
+ // HLSL: void computeMain(
+
+ bool result = true
+ && TEST_texture_float()
+ ;
+
+ // FUNCTIONAL: 1
+ outputBuffer[0] = int(result);
+}
generated by cgit v1.2.3 (git 2.39.1) at 2026-06-01 13:34:27 +0000