dawn-cmake/test/intrinsics/gen/modf/a128ab.wgsl.expected.msl
Ben Clayton 03c8393213 [writer/msl]: Implement modf and frexp
And remove the u32 overload of frexp (it's not in the spec).

Brings the number of failing tint end to end tests for MSL down to 19/1098.

The WG still haven't found consensus on reworking these two intrinsics.
It's very likely that their signature will change so that they return a structure instead of returning a value and outputing another as a pointer.

Until the WG makes a decision, let's implement these according to the current spec.
Some overloads are still failing due to MSL missing overloads of the pointer parameter being in the `threadgroup` address space.

I'm holding off fixing these until we know what's happening with these intrinsics.

See also:
https://github.com/gpuweb/gpuweb/issues/1480
https://github.com/gpuweb/gpuweb/issues/1846

Change-Id: Ib6764e6659d840db41bc65fed2b8b283d1056c3d
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57421
Commit-Queue: Ben Clayton <bclayton@chromium.org>
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
2021-07-08 21:21:27 +00:00

61 lines
3.9 KiB
Plaintext

SKIP: FAILED
#include <metal_stdlib>
using namespace metal;
void modf_a128ab(threadgroup float2* const tint_symbol_1) {
float2 res = modf(float2(), *(&(*(tint_symbol_1))));
}
kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
threadgroup float2 tint_symbol_2;
if ((local_invocation_index == 0u)) {
tint_symbol_2 = float2();
}
threadgroup_barrier(mem_flags::mem_threadgroup);
modf_a128ab(&(tint_symbol_2));
return;
}
Compilation failed:
program_source:5:16: error: no matching function for call to 'modf'
float2 res = modf(float2(), *(&(*(tint_symbol_1))));
^~~~
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:4854:19: note: candidate function not viable: address space mismatch in 2nd argument ('threadgroup float2' (vector of 2 'float' values)), parameter type must be 'metal::float2 &' (aka 'float2 &')
METAL_FUNC float2 modf(float2 x, thread float2 &intval)
^
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3142:17: note: candidate function not viable: no known conversion from 'float2' (vector of 2 'float' values) to 'half' for 1st argument
METAL_FUNC half modf(half x, thread half &intval)
^
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3386:18: note: candidate function not viable: no known conversion from 'float2' (vector of 2 'float' values) to 'metal::half2' (aka 'half2') for 1st argument
METAL_FUNC half2 modf(half2 x, thread half2 &intval)
^
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3630:18: note: candidate function not viable: no known conversion from 'float2' (vector of 2 'float' values) to 'metal::half3' (aka 'half3') for 1st argument
METAL_FUNC half3 modf(half3 x, thread half3 &intval)
^
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3874:18: note: candidate function not viable: no known conversion from 'float2' (vector of 2 'float' values) to 'metal::half4' (aka 'half4') for 1st argument
METAL_FUNC half4 modf(half4 x, thread half4 &intval)
^
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:4610:18: note: candidate function not viable: no known conversion from 'float2' (vector of 2 'float' values) to 'float' for 1st argument
METAL_FUNC float modf(float x, thread float &intval)
^
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:5098:19: note: candidate function not viable: no known conversion from 'float2' (vector of 2 'float' values) to 'metal::float3' (aka 'float3') for 1st argument
METAL_FUNC float3 modf(float3 x, thread float3 &intval)
^
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:5342:19: note: candidate function not viable: no known conversion from 'float2' (vector of 2 'float' values) to 'metal::float4' (aka 'float4') for 1st argument
METAL_FUNC float4 modf(float4 x, thread float4 &intval)
^
program_source:10:31: warning: equality comparison with extraneous parentheses
if ((local_invocation_index == 0u)) {
~~~~~~~~~~~~~~~~~~~~~~~^~~~~
program_source:10:31: note: remove extraneous parentheses around the comparison to silence this warning
if ((local_invocation_index == 0u)) {
~ ^ ~
program_source:10:31: note: use '=' to turn this equality comparison into an assignment
if ((local_invocation_index == 0u)) {
^~
=