tint: spir-v reader: fix atomicCompareExchangeWeak with var comparator
Also fix HLSL generator to unwrap the ref type when emitting the comparator value. Bug: tint:1185 Change-Id: I01d04ca6357e72fd5ead0f25012ab39794e65da5 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/94522 Reviewed-by: Ben Clayton <bclayton@chromium.org> Commit-Queue: Antonio Maiorano <amaiorano@google.com> Kokoro: Kokoro <noreply+kokoro@google.com>
This commit is contained in:
parent
606abfbeae
commit
f99671b830
|
@ -1779,8 +1779,8 @@ bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out,
|
|||
|
||||
{ // T compare_value = <compare_value>;
|
||||
auto pre = line();
|
||||
if (!EmitTypeAndName(pre, TypeOf(compare_value), ast::StorageClass::kNone,
|
||||
ast::Access::kUndefined, compare)) {
|
||||
if (!EmitTypeAndName(pre, TypeOf(compare_value)->UnwrapRef(),
|
||||
ast::StorageClass::kNone, ast::Access::kUndefined, compare)) {
|
||||
return false;
|
||||
}
|
||||
pre << " = ";
|
||||
|
|
|
@ -3272,7 +3272,8 @@ bool Builder::GenerateAtomicBuiltin(const sem::Call* call,
|
|||
value,
|
||||
});
|
||||
case sem::BuiltinType::kAtomicCompareExchangeWeak: {
|
||||
auto comparator = GenerateExpression(call->Arguments()[1]->Declaration());
|
||||
auto comparator =
|
||||
GenerateExpressionWithLoadIfNeeded(call->Arguments()[1]->Declaration());
|
||||
if (comparator == 0) {
|
||||
return false;
|
||||
}
|
||||
|
|
|
@ -1,9 +1,7 @@
|
|||
SKIP: FAILED
|
||||
|
||||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 35
|
||||
; Bound: 36
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
|
@ -45,32 +43,30 @@ SKIP: FAILED
|
|||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
|
||||
%_ptr_Function___atomic_compare_exchange_resulti32 = OpTypePointer Function %__atomic_compare_exchange_resulti32
|
||||
%28 = OpConstantNull %__atomic_compare_exchange_resulti32
|
||||
%29 = OpConstantNull %__atomic_compare_exchange_resulti32
|
||||
%atomicCompareExchangeWeak_1bd40a = OpFunction %void None %5
|
||||
%8 = OpLabel
|
||||
%arg_1 = OpVariable %_ptr_Function_int Function %12
|
||||
%arg_2 = OpVariable %_ptr_Function_int Function %12
|
||||
%res = OpVariable %_ptr_Function___atomic_compare_exchange_resulti32 Function %28
|
||||
%res = OpVariable %_ptr_Function___atomic_compare_exchange_resulti32 Function %29
|
||||
OpStore %arg_1 %int_1
|
||||
OpStore %arg_2 %int_1
|
||||
%22 = OpAccessChain %_ptr_StorageBuffer_int %sb_rw %uint_0
|
||||
%23 = OpLoad %int %arg_2
|
||||
%24 = OpAtomicCompareExchange %int %22 %uint_1 %uint_0 %uint_0 %23 %arg_1
|
||||
%25 = OpIEqual %bool %24 %23
|
||||
%14 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %24 %25
|
||||
%24 = OpLoad %int %arg_1
|
||||
%25 = OpAtomicCompareExchange %int %22 %uint_1 %uint_0 %uint_0 %23 %24
|
||||
%26 = OpIEqual %bool %25 %23
|
||||
%14 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %25 %26
|
||||
OpStore %res %14
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%fragment_main = OpFunction %void None %5
|
||||
%30 = OpLabel
|
||||
%31 = OpFunctionCall %void %atomicCompareExchangeWeak_1bd40a
|
||||
%31 = OpLabel
|
||||
%32 = OpFunctionCall %void %atomicCompareExchangeWeak_1bd40a
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%compute_main = OpFunction %void None %5
|
||||
%33 = OpLabel
|
||||
%34 = OpFunctionCall %void %atomicCompareExchangeWeak_1bd40a
|
||||
%34 = OpLabel
|
||||
%35 = OpFunctionCall %void %atomicCompareExchangeWeak_1bd40a
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
1:1: AtomicCompareExchange: expected Comparator to be of type Result Type
|
||||
%24 = OpAtomicCompareExchange %int %22 %uint_1 %uint_0 %uint_0 %23 %arg_1
|
||||
|
||||
|
|
|
@ -1,9 +1,7 @@
|
|||
SKIP: FAILED
|
||||
|
||||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 33
|
||||
; Bound: 34
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
|
@ -43,32 +41,30 @@ SKIP: FAILED
|
|||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
|
||||
%_ptr_Function___atomic_compare_exchange_resultu32 = OpTypePointer Function %__atomic_compare_exchange_resultu32
|
||||
%26 = OpConstantNull %__atomic_compare_exchange_resultu32
|
||||
%27 = OpConstantNull %__atomic_compare_exchange_resultu32
|
||||
%atomicCompareExchangeWeak_63d8e6 = OpFunction %void None %5
|
||||
%8 = OpLabel
|
||||
%arg_1 = OpVariable %_ptr_Function_uint Function %12
|
||||
%arg_2 = OpVariable %_ptr_Function_uint Function %12
|
||||
%res = OpVariable %_ptr_Function___atomic_compare_exchange_resultu32 Function %26
|
||||
%res = OpVariable %_ptr_Function___atomic_compare_exchange_resultu32 Function %27
|
||||
OpStore %arg_1 %uint_1
|
||||
OpStore %arg_2 %uint_1
|
||||
%20 = OpAccessChain %_ptr_StorageBuffer_uint %sb_rw %uint_0
|
||||
%21 = OpLoad %uint %arg_2
|
||||
%22 = OpAtomicCompareExchange %uint %20 %uint_1 %uint_0 %uint_0 %21 %arg_1
|
||||
%23 = OpIEqual %bool %22 %21
|
||||
%14 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %22 %23
|
||||
%22 = OpLoad %uint %arg_1
|
||||
%23 = OpAtomicCompareExchange %uint %20 %uint_1 %uint_0 %uint_0 %21 %22
|
||||
%24 = OpIEqual %bool %23 %21
|
||||
%14 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %23 %24
|
||||
OpStore %res %14
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%fragment_main = OpFunction %void None %5
|
||||
%28 = OpLabel
|
||||
%29 = OpFunctionCall %void %atomicCompareExchangeWeak_63d8e6
|
||||
%29 = OpLabel
|
||||
%30 = OpFunctionCall %void %atomicCompareExchangeWeak_63d8e6
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%compute_main = OpFunction %void None %5
|
||||
%31 = OpLabel
|
||||
%32 = OpFunctionCall %void %atomicCompareExchangeWeak_63d8e6
|
||||
%32 = OpLabel
|
||||
%33 = OpFunctionCall %void %atomicCompareExchangeWeak_63d8e6
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
1:1: AtomicCompareExchange: expected Comparator to be of type Result Type
|
||||
%22 = OpAtomicCompareExchange %uint %20 %uint_1 %uint_0 %uint_0 %21 %arg_1
|
||||
|
||||
|
|
|
@ -1,17 +1,34 @@
|
|||
SKIP: FAILED
|
||||
struct atomic_compare_exchange_resultu32 {
|
||||
uint old_value;
|
||||
bool exchanged;
|
||||
};
|
||||
groupshared uint arg_0;
|
||||
|
||||
|
||||
var<workgroup> arg_0 : atomic<u32>;
|
||||
|
||||
fn atomicCompareExchangeWeak_83580d() {
|
||||
var arg_1 = 1u;
|
||||
var arg_2 = 1u;
|
||||
var res = atomicCompareExchangeWeak(&(arg_0), arg_1, arg_2);
|
||||
void atomicCompareExchangeWeak_83580d() {
|
||||
uint arg_1 = 1u;
|
||||
uint arg_2 = 1u;
|
||||
atomic_compare_exchange_resultu32 atomic_result = (atomic_compare_exchange_resultu32)0;
|
||||
uint atomic_compare_value = arg_1;
|
||||
InterlockedCompareExchange(arg_0, atomic_compare_value, arg_2, atomic_result.old_value);
|
||||
atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
|
||||
atomic_compare_exchange_resultu32 res = atomic_result;
|
||||
}
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn compute_main() {
|
||||
struct tint_symbol_1 {
|
||||
uint local_invocation_index : SV_GroupIndex;
|
||||
};
|
||||
|
||||
void compute_main_inner(uint local_invocation_index) {
|
||||
{
|
||||
uint atomic_result_1 = 0u;
|
||||
InterlockedExchange(arg_0, 0u, atomic_result_1);
|
||||
}
|
||||
GroupMemoryBarrierWithGroupSync();
|
||||
atomicCompareExchangeWeak_83580d();
|
||||
}
|
||||
|
||||
Failed to generate: error: unknown type in EmitType
|
||||
[numthreads(1, 1, 1)]
|
||||
void compute_main(tint_symbol_1 tint_symbol) {
|
||||
compute_main_inner(tint_symbol.local_invocation_index);
|
||||
return;
|
||||
}
|
||||
|
|
|
@ -1,9 +1,7 @@
|
|||
SKIP: FAILED
|
||||
|
||||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 40
|
||||
; Bound: 41
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
|
@ -39,37 +37,35 @@ SKIP: FAILED
|
|||
%uint_2 = OpConstant %uint 2
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Function___atomic_compare_exchange_resultu32 = OpTypePointer Function %__atomic_compare_exchange_resultu32
|
||||
%26 = OpConstantNull %__atomic_compare_exchange_resultu32
|
||||
%27 = OpTypeFunction %void %uint
|
||||
%27 = OpConstantNull %__atomic_compare_exchange_resultu32
|
||||
%28 = OpTypeFunction %void %uint
|
||||
%uint_264 = OpConstant %uint 264
|
||||
%atomicCompareExchangeWeak_83580d = OpFunction %void None %6
|
||||
%9 = OpLabel
|
||||
%arg_1 = OpVariable %_ptr_Function_uint Function %13
|
||||
%arg_2 = OpVariable %_ptr_Function_uint Function %13
|
||||
%res = OpVariable %_ptr_Function___atomic_compare_exchange_resultu32 Function %26
|
||||
%res = OpVariable %_ptr_Function___atomic_compare_exchange_resultu32 Function %27
|
||||
OpStore %arg_1 %uint_1
|
||||
OpStore %arg_2 %uint_1
|
||||
%21 = OpLoad %uint %arg_2
|
||||
%22 = OpAtomicCompareExchange %uint %arg_0 %uint_2 %uint_0 %uint_0 %21 %arg_1
|
||||
%23 = OpIEqual %bool %22 %21
|
||||
%15 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %22 %23
|
||||
%22 = OpLoad %uint %arg_1
|
||||
%23 = OpAtomicCompareExchange %uint %arg_0 %uint_2 %uint_0 %uint_0 %21 %22
|
||||
%24 = OpIEqual %bool %23 %21
|
||||
%15 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %23 %24
|
||||
OpStore %res %15
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%compute_main_inner = OpFunction %void None %27
|
||||
%compute_main_inner = OpFunction %void None %28
|
||||
%local_invocation_index = OpFunctionParameter %uint
|
||||
%30 = OpLabel
|
||||
%31 = OpLabel
|
||||
OpAtomicStore %arg_0 %uint_2 %uint_0 %13
|
||||
OpControlBarrier %uint_2 %uint_2 %uint_264
|
||||
%35 = OpFunctionCall %void %atomicCompareExchangeWeak_83580d
|
||||
%36 = OpFunctionCall %void %atomicCompareExchangeWeak_83580d
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%compute_main = OpFunction %void None %6
|
||||
%37 = OpLabel
|
||||
%39 = OpLoad %uint %local_invocation_index_1
|
||||
%38 = OpFunctionCall %void %compute_main_inner %39
|
||||
%38 = OpLabel
|
||||
%40 = OpLoad %uint %local_invocation_index_1
|
||||
%39 = OpFunctionCall %void %compute_main_inner %40
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
1:1: AtomicCompareExchange: expected Comparator to be of type Result Type
|
||||
%22 = OpAtomicCompareExchange %uint %arg_0 %uint_2 %uint_0 %uint_0 %21 %arg_1
|
||||
|
||||
|
|
|
@ -1,17 +1,34 @@
|
|||
SKIP: FAILED
|
||||
struct atomic_compare_exchange_resulti32 {
|
||||
int old_value;
|
||||
bool exchanged;
|
||||
};
|
||||
groupshared int arg_0;
|
||||
|
||||
|
||||
var<workgroup> arg_0 : atomic<i32>;
|
||||
|
||||
fn atomicCompareExchangeWeak_e88938() {
|
||||
var arg_1 = 1;
|
||||
var arg_2 = 1;
|
||||
var res = atomicCompareExchangeWeak(&(arg_0), arg_1, arg_2);
|
||||
void atomicCompareExchangeWeak_e88938() {
|
||||
int arg_1 = 1;
|
||||
int arg_2 = 1;
|
||||
atomic_compare_exchange_resulti32 atomic_result = (atomic_compare_exchange_resulti32)0;
|
||||
int atomic_compare_value = arg_1;
|
||||
InterlockedCompareExchange(arg_0, atomic_compare_value, arg_2, atomic_result.old_value);
|
||||
atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
|
||||
atomic_compare_exchange_resulti32 res = atomic_result;
|
||||
}
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn compute_main() {
|
||||
struct tint_symbol_1 {
|
||||
uint local_invocation_index : SV_GroupIndex;
|
||||
};
|
||||
|
||||
void compute_main_inner(uint local_invocation_index) {
|
||||
{
|
||||
int atomic_result_1 = 0;
|
||||
InterlockedExchange(arg_0, 0, atomic_result_1);
|
||||
}
|
||||
GroupMemoryBarrierWithGroupSync();
|
||||
atomicCompareExchangeWeak_e88938();
|
||||
}
|
||||
|
||||
Failed to generate: error: unknown type in EmitType
|
||||
[numthreads(1, 1, 1)]
|
||||
void compute_main(tint_symbol_1 tint_symbol) {
|
||||
compute_main_inner(tint_symbol.local_invocation_index);
|
||||
return;
|
||||
}
|
||||
|
|
|
@ -1,9 +1,7 @@
|
|||
SKIP: FAILED
|
||||
|
||||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 41
|
||||
; Bound: 42
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
|
@ -40,37 +38,35 @@ SKIP: FAILED
|
|||
%uint_2 = OpConstant %uint 2
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Function___atomic_compare_exchange_resulti32 = OpTypePointer Function %__atomic_compare_exchange_resulti32
|
||||
%27 = OpConstantNull %__atomic_compare_exchange_resulti32
|
||||
%28 = OpTypeFunction %void %uint
|
||||
%28 = OpConstantNull %__atomic_compare_exchange_resulti32
|
||||
%29 = OpTypeFunction %void %uint
|
||||
%uint_264 = OpConstant %uint 264
|
||||
%atomicCompareExchangeWeak_e88938 = OpFunction %void None %7
|
||||
%10 = OpLabel
|
||||
%arg_1 = OpVariable %_ptr_Function_int Function %14
|
||||
%arg_2 = OpVariable %_ptr_Function_int Function %14
|
||||
%res = OpVariable %_ptr_Function___atomic_compare_exchange_resulti32 Function %27
|
||||
%res = OpVariable %_ptr_Function___atomic_compare_exchange_resulti32 Function %28
|
||||
OpStore %arg_1 %int_1
|
||||
OpStore %arg_2 %int_1
|
||||
%22 = OpLoad %int %arg_2
|
||||
%23 = OpAtomicCompareExchange %int %arg_0 %uint_2 %uint_0 %uint_0 %22 %arg_1
|
||||
%24 = OpIEqual %bool %23 %22
|
||||
%16 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %23 %24
|
||||
%23 = OpLoad %int %arg_1
|
||||
%24 = OpAtomicCompareExchange %int %arg_0 %uint_2 %uint_0 %uint_0 %22 %23
|
||||
%25 = OpIEqual %bool %24 %22
|
||||
%16 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %24 %25
|
||||
OpStore %res %16
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%compute_main_inner = OpFunction %void None %28
|
||||
%compute_main_inner = OpFunction %void None %29
|
||||
%local_invocation_index = OpFunctionParameter %uint
|
||||
%31 = OpLabel
|
||||
%32 = OpLabel
|
||||
OpAtomicStore %arg_0 %uint_2 %uint_0 %14
|
||||
OpControlBarrier %uint_2 %uint_2 %uint_264
|
||||
%36 = OpFunctionCall %void %atomicCompareExchangeWeak_e88938
|
||||
%37 = OpFunctionCall %void %atomicCompareExchangeWeak_e88938
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%compute_main = OpFunction %void None %7
|
||||
%38 = OpLabel
|
||||
%40 = OpLoad %uint %local_invocation_index_1
|
||||
%39 = OpFunctionCall %void %compute_main_inner %40
|
||||
%39 = OpLabel
|
||||
%41 = OpLoad %uint %local_invocation_index_1
|
||||
%40 = OpFunctionCall %void %compute_main_inner %41
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
1:1: AtomicCompareExchange: expected Comparator to be of type Result Type
|
||||
%23 = OpAtomicCompareExchange %int %arg_0 %uint_2 %uint_0 %uint_0 %22 %arg_1
|
||||
|
||||
|
|
Loading…
Reference in New Issue