diff --git a/CHANGELOG.md b/CHANGELOG.md index ab4e955..c780923 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -7,6 +7,15 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 ## [Unreleased] +### Fixed (MSL) + +- **64-bit atomic capability validation** ([#79](https://github.com/gogpu/naga/issues/79)) — + `msl.Compile` now returns a descriptive error for 64-bit atomic operations + that Metal cannot represent instead of emitting invalid intrinsics. Matching + Rust Naga, result-discarded `min`/`max` in storage address space remain + supported; loads, stores, full operations, result-producing min/max, and + workgroup min/max are rejected. + ## [0.17.15] - 2026-06-15 ### Fixed (MSL) diff --git a/msl/internal/codegen/expressions.go b/msl/internal/codegen/expressions.go index 9dfde63..83f79f7 100644 --- a/msl/internal/codegen/expressions.go +++ b/msl/internal/codegen/expressions.go @@ -1170,6 +1170,9 @@ func (w *Writer) writeLoad(load ir.ExprLoad) error { // Matches Rust naga MSL put_unchecked_load. func (w *Writer) writeUncheckedLoad(pointer ir.ExpressionHandle) error { if w.isAtomicPointer(pointer) { + if err := w.validateAtomicOperation(pointer, ir.AtomicLoad{}, nil); err != nil { + return err + } w.write("metal::atomic_load_explicit(&") if err := w.writeExpression(pointer); err != nil { return err diff --git a/msl/internal/codegen/int64_atomics_test.go b/msl/internal/codegen/int64_atomics_test.go new file mode 100644 index 0000000..0edd167 --- /dev/null +++ b/msl/internal/codegen/int64_atomics_test.go @@ -0,0 +1,281 @@ +package codegen + +import ( + "strings" + "testing" + + "github.com/gogpu/naga/wgsl" +) + +func compileWGSLForInt64AtomicPolicy(t *testing.T, source string) (string, error) { + t.Helper() + + lexer := wgsl.NewLexer(source) + tokens, err := lexer.Tokenize() + if err != nil { + t.Fatalf("tokenize WGSL: %v", err) + } + parser := wgsl.NewParser(tokens) + ast, err := parser.Parse() + if err != nil { + t.Fatalf("parse WGSL: %v", err) + } + module, err := wgsl.LowerWithSource(ast, source) + if err != nil { + t.Fatalf("lower WGSL: %v", err) + } + + options := DefaultOptions() + options.LangVersion = Version2_4 + options.FakeMissingBindings = true + code, _, err := Compile(module, options) + return code, err +} + +func TestMSLInt64AtomicPolicy(t *testing.T) { + t.Run("allows result-discarded storage min max", func(t *testing.T) { + const source = ` +@group(0) @binding(0) +var unsigned_value: atomic; +@group(0) @binding(1) +var signed_value: atomic; + +@compute @workgroup_size(1) +fn main() { + atomicMin(&unsigned_value, 1lu); + atomicMax(&unsigned_value, 2lu); + atomicMin(&signed_value, 1li); + atomicMax(&signed_value, 2li); +} +` + + code, err := compileWGSLForInt64AtomicPolicy(t, source) + if err != nil { + t.Fatalf("compile supported 64-bit storage min/max: %v", err) + } + for _, want := range []string{"atomic_min_explicit", "atomic_max_explicit"} { + if !strings.Contains(code, want) { + t.Errorf("MSL does not contain %q:\n%s", want, code) + } + } + for _, unwanted := range []string{"atomic_fetch_min_explicit", "atomic_fetch_max_explicit"} { + if strings.Contains(code, unwanted) { + t.Errorf("MSL contains unsupported result-producing intrinsic %q:\n%s", unwanted, code) + } + } + }) + + t.Run("allows nested storage min max", func(t *testing.T) { + const source = ` +struct AtomicPair { + values: array, 2>, +} + +@group(0) @binding(0) +var pairs: array; + +@compute @workgroup_size(1) +fn main() { + atomicMin(&pairs[1].values[0], 1lu); + atomicMax(&pairs[0].values[1], 2lu); +} +` + + code, err := compileWGSLForInt64AtomicPolicy(t, source) + if err != nil { + t.Fatalf("compile nested 64-bit storage min/max: %v", err) + } + if !strings.Contains(code, "atomic_min_explicit") || !strings.Contains(code, "atomic_max_explicit") { + t.Fatalf("nested storage min/max intrinsics missing:\n%s", code) + } + }) + + rejected := []struct { + name string + operation string + source string + }{ + { + name: "load", + operation: "load", + source: ` +@group(0) @binding(0) var value: atomic; +@compute @workgroup_size(1) fn main() { + let old = atomicLoad(&value); + _ = old; +} +`, + }, + { + name: "store", + operation: "store", + source: ` +@group(0) @binding(0) var value: atomic; +@compute @workgroup_size(1) fn main() { + atomicStore(&value, 1lu); +} +`, + }, + { + name: "add", + operation: "add", + source: ` +@group(0) @binding(0) var value: atomic; +@compute @workgroup_size(1) fn main() { + atomicAdd(&value, 1lu); +} +`, + }, + { + name: "subtract", + operation: "subtract", + source: ` +@group(0) @binding(0) var value: atomic; +@compute @workgroup_size(1) fn main() { + atomicSub(&value, 1lu); +} +`, + }, + { + name: "and", + operation: "and", + source: ` +@group(0) @binding(0) var value: atomic; +@compute @workgroup_size(1) fn main() { + atomicAnd(&value, 1lu); +} +`, + }, + { + name: "or", + operation: "or", + source: ` +@group(0) @binding(0) var value: atomic; +@compute @workgroup_size(1) fn main() { + atomicOr(&value, 1lu); +} +`, + }, + { + name: "xor", + operation: "xor", + source: ` +@group(0) @binding(0) var value: atomic; +@compute @workgroup_size(1) fn main() { + atomicXor(&value, 1lu); +} +`, + }, + { + name: "exchange", + operation: "exchange", + source: ` +@group(0) @binding(0) var value: atomic; +@compute @workgroup_size(1) fn main() { + atomicExchange(&value, 1lu); +} +`, + }, + { + name: "compare exchange", + operation: "compare exchange", + source: ` +@group(0) @binding(0) var value: atomic; +@compute @workgroup_size(1) fn main() { + let result = atomicCompareExchangeWeak(&value, 0lu, 1lu); + _ = result.old_value; +} +`, + }, + { + name: "result-producing min", + operation: "min", + source: ` +@group(0) @binding(0) var value: atomic; +@compute @workgroup_size(1) fn main() { + let old = atomicMin(&value, 1lu); + _ = old; +} +`, + }, + { + name: "result-producing max", + operation: "max", + source: ` +@group(0) @binding(0) var value: atomic; +@compute @workgroup_size(1) fn main() { + let old = atomicMax(&value, 1lu); + _ = old; +} +`, + }, + { + name: "workgroup min", + operation: "min", + source: ` +var value: atomic; +@compute @workgroup_size(1) fn main() { + atomicMin(&value, 1lu); +} +`, + }, + { + name: "workgroup max", + operation: "max", + source: ` +var value: atomic; +@compute @workgroup_size(1) fn main() { + atomicMax(&value, 1lu); +} +`, + }, + } + + for _, test := range rejected { + t.Run("rejects "+test.name, func(t *testing.T) { + code, err := compileWGSLForInt64AtomicPolicy(t, test.source) + if err == nil { + t.Fatalf("expected 64-bit atomic %s rejection, got MSL:\n%s", test.name, code) + } + for _, want := range []string{ + "64-bit atomic " + test.operation, + "only result-discarded min/max in the storage address space", + } { + if !strings.Contains(err.Error(), want) { + t.Errorf("error %q does not contain %q", err, want) + } + } + if code != "" { + t.Errorf("Compile returned source with error:\n%s", code) + } + }) + } + + t.Run("preserves 32-bit atomics", func(t *testing.T) { + const source = ` +@group(0) @binding(0) var value: atomic; +@compute @workgroup_size(1) fn main() { + atomicStore(&value, 1u); + let loaded = atomicLoad(&value); + let added = atomicAdd(&value, loaded); + let exchanged = atomicExchange(&value, added); + _ = exchanged; +} +` + + code, err := compileWGSLForInt64AtomicPolicy(t, source) + if err != nil { + t.Fatalf("compile 32-bit atomic control: %v", err) + } + for _, want := range []string{ + "atomic_store_explicit", + "atomic_load_explicit", + "atomic_fetch_add_explicit", + "atomic_exchange_explicit", + } { + if !strings.Contains(code, want) { + t.Errorf("32-bit MSL does not contain %q:\n%s", want, code) + } + } + }) +} diff --git a/msl/internal/codegen/msl_xcrun_test.go b/msl/internal/codegen/msl_xcrun_test.go index 1a3a427..be56696 100644 --- a/msl/internal/codegen/msl_xcrun_test.go +++ b/msl/internal/codegen/msl_xcrun_test.go @@ -6,6 +6,7 @@ import ( "testing" "github.com/gogpu/naga" + "github.com/gogpu/naga/ir" ) func TestMSLCompilesWithXcrun(t *testing.T) { @@ -55,3 +56,44 @@ fn fs_main(input: VertexOutput) -> @location(0) vec4 { } verifyMSLWithXcrun(t, mslSource) } + +func TestMSLInt64AtomicMinMaxCompilesWithXcrun(t *testing.T) { + const wgslSource = ` +@group(0) @binding(0) +var value: atomic; + +@compute @workgroup_size(1) +fn main() { + atomicMin(&value, 1lu); + atomicMax(&value, 2lu); +} +` + + ast, err := naga.Parse(wgslSource) + if err != nil { + t.Fatalf("naga.Parse failed: %v", err) + } + module, err := naga.LowerWithSource(ast, wgslSource) + if err != nil { + t.Fatalf("naga.LowerWithSource failed: %v", err) + } + + bufferSlot := uint8(0) + options := DefaultOptions() + options.LangVersion = Version2_4 + options.PerEntryPointMap = map[string]EntryPointResources{ + "main": { + Resources: map[ir.ResourceBinding]BindTarget{ + {Group: 0, Binding: 0}: { + Buffer: &bufferSlot, + Mutable: true, + }, + }, + }, + } + mslSource, _, err := Compile(module, options) + if err != nil { + t.Fatalf("msl.Compile failed: %v", err) + } + verifyMSLWithXcrun(t, mslSource) +} diff --git a/msl/internal/codegen/statements.go b/msl/internal/codegen/statements.go index 74b16b8..4d4d8e7 100644 --- a/msl/internal/codegen/statements.go +++ b/msl/internal/codegen/statements.go @@ -1170,6 +1170,9 @@ func (w *Writer) writeUncheckedStore(store ir.StmtStore) error { // Rust naga emits StmtStore (not StmtAtomic) for atomicStore(), // and the MSL backend detects atomic pointers at render time. if w.isAtomicPointer(store.Pointer) { + if err := w.validateAtomicOperation(store.Pointer, ir.AtomicStore{}, nil); err != nil { + return err + } w.WriteIndent() w.write("%satomic_store_explicit(&", Namespace) if err := w.writeExpression(store.Pointer); err != nil { @@ -1198,6 +1201,73 @@ func (w *Writer) writeUncheckedStore(store ir.StmtStore) error { return nil } +// validateAtomicOperation enforces the subset of 64-bit atomics supported by +// the MSL backend. This matches Rust naga's SHADER_INT64_ATOMIC_MIN_MAX +// capability: only result-discarded min/max operations in storage address +// space are supported. MSL has no SHADER_INT64_ATOMIC_ALL_OPS equivalent. +func (w *Writer) validateAtomicOperation( + pointer ir.ExpressionHandle, + fun ir.AtomicFunction, + result *ir.ExpressionHandle, +) error { + if w.currentFunction == nil { + return nil + } + + scalar := w.resolveAtomicScalarFromPointer(w.currentFunction, pointer) + if scalar == nil || scalar.Width != 8 { + return nil + } + + resultDiscardedMinMax := false + switch fun.(type) { + case ir.AtomicMin, ir.AtomicMax: + resultDiscardedMinMax = result == nil + } + if resultDiscardedMinMax { + if space, ok := w.getPointerAddressSpace(pointer); ok && space == ir.SpaceStorage { + return nil + } + } + + return fmt.Errorf( + "64-bit atomic %s is unsupported by MSL; only result-discarded min/max in the storage address space are supported", + atomicFunctionName(fun), + ) +} + +// atomicFunctionName returns the WGSL-facing name used in MSL capability +// errors. Keeping these stable makes backend errors actionable to callers. +func atomicFunctionName(fun ir.AtomicFunction) string { + switch f := fun.(type) { + case ir.AtomicAdd: + return "add" + case ir.AtomicSubtract: + return "subtract" + case ir.AtomicAnd: + return "and" + case ir.AtomicExclusiveOr: + return "xor" + case ir.AtomicInclusiveOr: + return "or" + case ir.AtomicMin: + return "min" + case ir.AtomicMax: + return "max" + case ir.AtomicExchange: + if f.Compare != nil { + return "compare exchange" + } + return "exchange" + case ir.AtomicLoad: + return "load" + case ir.AtomicStore: + return "store" + default: + return fmt.Sprintf("%T", fun) + } +} + // isAtomicPointer checks if an expression is a pointer to an atomic type. // Matches Rust naga's TypeInner::is_atomic_pointer. // Note: our type resolution strips the pointer wrapper for access chains, @@ -1576,6 +1646,10 @@ func (w *Writer) writeImageAtomic(imgAtomic ir.StmtImageAtomic) error { // writeAtomic writes an atomic operation statement. func (w *Writer) writeAtomic(atomic ir.StmtAtomic) error { + if err := w.validateAtomicOperation(atomic.Pointer, atomic.Fun, atomic.Result); err != nil { + return err + } + // Determine the function based on atomic operation type var funcName string switch f := atomic.Fun.(type) { diff --git a/snapshot/snapshot_test.go b/snapshot/snapshot_test.go index b81596f..4fbe8f3 100644 --- a/snapshot/snapshot_test.go +++ b/snapshot/snapshot_test.go @@ -84,7 +84,17 @@ func TestSnapshots(t *testing.T) { }) t.Run("msl", func(t *testing.T) { - code := compileMSL(t, module) + code, compileErr := compileMSL(t, module) + if compileErr != nil { + target, configErr := rustTargetForShader(shader.name, "METAL") + if configErr != nil { + t.Fatalf("read Rust target config: %v", configErr) + } + if target == rustTargetExcluded { + t.Skipf("Rust config explicitly excludes METAL: %v", compileErr) + } + t.Fatalf("MSL compile failed: %v", compileErr) + } compareGolden(t, filepath.Join("testdata", "golden", "msl", shader.name+".msl"), code) }) }) @@ -1311,12 +1321,12 @@ func glslStageName(stage ir.ShaderStage) string { // compileMSL compiles the IR module to MSL source. // Uses metal1.0 to match Rust naga default for snapshot tests. -func compileMSL(t *testing.T, module *ir.Module) string { +func compileMSL(t *testing.T, module *ir.Module) (string, error) { return compileMSLWithVersion(t, module, msl.Version1_0) } // compileMSLWithVersion compiles the IR module to MSL source with a specific version. -func compileMSLWithVersion(t *testing.T, module *ir.Module, version msl.Version) string { +func compileMSLWithVersion(t *testing.T, module *ir.Module, version msl.Version) (string, error) { t.Helper() opts := msl.DefaultOptions() @@ -1324,10 +1334,10 @@ func compileMSLWithVersion(t *testing.T, module *ir.Module, version msl.Version) opts.FakeMissingBindings = true code, _, err := msl.Compile(module, opts) if err != nil { - t.Skipf("MSL compile failed (skipping): %v", err) + return "", err } - return code + return code, nil } // compileMSLWithOpts compiles the IR module to MSL source with fully specified options. @@ -1349,6 +1359,44 @@ func compileMSLWithOpts(t *testing.T, module *ir.Module, opts msl.Options) (stri // can access them without the full Rust reference checkout. const rustTomlDir = "testdata/config" +type rustTargetClassification uint8 + +const ( + rustTargetUndeclared rustTargetClassification = iota + rustTargetEnabled + rustTargetExcluded +) + +var rustTargetsPattern = regexp.MustCompile(`(?m)^\s*targets\s*=\s*"([^"]*)"`) + +// classifyRustTarget reports whether a Rust naga TOML config explicitly +// enables or excludes a backend target. Missing target declarations stay +// undeclared so backend compile errors fail instead of being hidden. +func classifyRustTarget(content, target string) rustTargetClassification { + match := rustTargetsPattern.FindStringSubmatch(content) + if len(match) != 2 { + return rustTargetUndeclared + } + for item := range strings.SplitSeq(match[1], "|") { + if strings.TrimSpace(item) == target { + return rustTargetEnabled + } + } + return rustTargetExcluded +} + +func rustTargetForShader(shaderName, target string) (rustTargetClassification, error) { + tomlPath := filepath.Join(rustTomlDir, shaderName+".toml") + data, err := os.ReadFile(tomlPath) + if os.IsNotExist(err) { + return rustTargetUndeclared, nil + } + if err != nil { + return rustTargetUndeclared, fmt.Errorf("read %s: %w", tomlPath, err) + } + return classifyRustTarget(string(data), target), nil +} + // readRustMSLConfig reads MSL options from a Rust naga test TOML config file. // Parses lang_version, fake_missing_bindings, bounds_check_policies, // and per_entry_point_map with resource bindings. diff --git a/snapshot/target_config_test.go b/snapshot/target_config_test.go new file mode 100644 index 0000000..9a78a3f --- /dev/null +++ b/snapshot/target_config_test.go @@ -0,0 +1,61 @@ +package snapshot_test + +import "testing" + +func TestRustTargetClassification(t *testing.T) { + tests := []struct { + name string + content string + target string + want rustTargetClassification + }{ + { + name: "enabled among pipe separated targets", + content: `targets = "SPIRV | METAL | HLSL | WGSL"`, + target: "METAL", + want: rustTargetEnabled, + }, + { + name: "explicitly excluded", + content: `targets = "SPIRV | HLSL | WGSL"`, + target: "METAL", + want: rustTargetExcluded, + }, + { + name: "whitespace around declaration", + content: " targets = \"METAL | WGSL\"\n", + target: "METAL", + want: rustTargetEnabled, + }, + { + name: "no targets declaration", + content: "[msl]\nlang_version = [2, 4]\n", + target: "METAL", + want: rustTargetUndeclared, + }, + { + name: "comment is not a declaration", + content: `# targets = "METAL"`, + target: "METAL", + want: rustTargetUndeclared, + }, + } + + for _, test := range tests { + t.Run(test.name, func(t *testing.T) { + if got := classifyRustTarget(test.content, test.target); got != test.want { + t.Fatalf("classifyRustTarget() = %d, want %d", got, test.want) + } + }) + } + + t.Run("missing config stays undeclared", func(t *testing.T) { + got, err := rustTargetForShader("definitely-not-a-real-shader", "METAL") + if err != nil { + t.Fatalf("rustTargetForShader() error: %v", err) + } + if got != rustTargetUndeclared { + t.Fatalf("rustTargetForShader() = %d, want undeclared", got) + } + }) +} diff --git a/snapshot/testdata/golden/msl/atomicCompareExchange-int64.msl b/snapshot/testdata/golden/msl/atomicCompareExchange-int64.msl deleted file mode 100644 index c71e7ff..0000000 --- a/snapshot/testdata/golden/msl/atomicCompareExchange-int64.msl +++ /dev/null @@ -1,180 +0,0 @@ -// language: metal1.0 -#include -#include - -using metal::uint; -struct DefaultConstructible { - template - operator T() && { - return T {}; - } -}; - -struct type_2 { - metal::atomic_long inner[128]; -}; -struct type_4 { - metal::atomic_ulong inner[128]; -}; -struct _atomic_compare_exchange_result_Sint_8_ { - long old_value; - bool exchanged; - char _pad2[7]; -}; -struct _atomic_compare_exchange_result_Uint_8_ { - ulong old_value; - bool exchanged; - char _pad2[7]; -}; - -template -_atomic_compare_exchange_result_Sint_8_ naga_atomic_compare_exchange_weak_explicit( - device A *atomic_ptr, - long cmp, - long v -) { - bool swapped = metal::atomic_compare_exchange_weak_explicit( - atomic_ptr, &cmp, v, - metal::memory_order_relaxed, metal::memory_order_relaxed - ); - return _atomic_compare_exchange_result_Sint_8_{cmp, swapped}; -} -template -_atomic_compare_exchange_result_Sint_8_ naga_atomic_compare_exchange_weak_explicit( - threadgroup A *atomic_ptr, - long cmp, - long v -) { - bool swapped = metal::atomic_compare_exchange_weak_explicit( - atomic_ptr, &cmp, v, - metal::memory_order_relaxed, metal::memory_order_relaxed - ); - return _atomic_compare_exchange_result_Sint_8_{cmp, swapped}; -} - -template -_atomic_compare_exchange_result_Uint_8_ naga_atomic_compare_exchange_weak_explicit( - device A *atomic_ptr, - ulong cmp, - ulong v -) { - bool swapped = metal::atomic_compare_exchange_weak_explicit( - atomic_ptr, &cmp, v, - metal::memory_order_relaxed, metal::memory_order_relaxed - ); - return _atomic_compare_exchange_result_Uint_8_{cmp, swapped}; -} -template -_atomic_compare_exchange_result_Uint_8_ naga_atomic_compare_exchange_weak_explicit( - threadgroup A *atomic_ptr, - ulong cmp, - ulong v -) { - bool swapped = metal::atomic_compare_exchange_weak_explicit( - atomic_ptr, &cmp, v, - metal::memory_order_relaxed, metal::memory_order_relaxed - ); - return _atomic_compare_exchange_result_Uint_8_{cmp, swapped}; -} -constant uint SIZE = 128u; - -kernel void test_atomic_compare_exchange_i64_( - device type_2& arr_i64_ [[user(fake0)]] -) { - uint i = 0u; - long old = {}; - bool exchanged = {}; - uint2 loop_bound = uint2(4294967295u); - bool loop_init = true; - while(true) { - if (metal::all(loop_bound == uint2(0u))) { break; } - loop_bound -= uint2(loop_bound.y == 0u, 1u); - if (!loop_init) { - uint _e26 = i; - i = _e26 + 1u; - } - loop_init = false; - uint _e2 = i; - if (_e2 < SIZE) { - } else { - break; - } - { - uint _e6 = i; - long _e8 = uint(_e6) < 128 ? metal::atomic_load_explicit(&arr_i64_.inner[_e6], metal::memory_order_relaxed) : DefaultConstructible(); - old = _e8; - exchanged = false; - uint2 loop_bound_1 = uint2(4294967295u); - while(true) { - if (metal::all(loop_bound_1 == uint2(0u))) { break; } - loop_bound_1 -= uint2(loop_bound_1.y == 0u, 1u); - bool _e12 = exchanged; - if (!(_e12)) { - } else { - break; - } - { - long _e14 = old; - long new_ = as_type(as_type(as_type(_e14) + as_type(10L))); - uint _e19 = i; - long _e21 = old; - _atomic_compare_exchange_result_Sint_8_ _e22 = naga_atomic_compare_exchange_weak_explicit(&uint(_e19) < 128 ? arr_i64_.inner[_e19] : DefaultConstructible(), _e21, new_); - old = _e22.old_value; - exchanged = _e22.exchanged; - } - } - } - } - return; -} - - -kernel void test_atomic_compare_exchange_u64_( - device type_4& arr_u64_ [[user(fake0)]] -) { - uint i_1 = 0u; - ulong old_1 = {}; - bool exchanged_1 = {}; - uint2 loop_bound_2 = uint2(4294967295u); - bool loop_init_1 = true; - while(true) { - if (metal::all(loop_bound_2 == uint2(0u))) { break; } - loop_bound_2 -= uint2(loop_bound_2.y == 0u, 1u); - if (!loop_init_1) { - uint _e26 = i_1; - i_1 = _e26 + 1u; - } - loop_init_1 = false; - uint _e2 = i_1; - if (_e2 < SIZE) { - } else { - break; - } - { - uint _e6 = i_1; - ulong _e8 = uint(_e6) < 128 ? metal::atomic_load_explicit(&arr_u64_.inner[_e6], metal::memory_order_relaxed) : DefaultConstructible(); - old_1 = _e8; - exchanged_1 = false; - uint2 loop_bound_3 = uint2(4294967295u); - while(true) { - if (metal::all(loop_bound_3 == uint2(0u))) { break; } - loop_bound_3 -= uint2(loop_bound_3.y == 0u, 1u); - bool _e12 = exchanged_1; - if (!(_e12)) { - } else { - break; - } - { - ulong _e14 = old_1; - ulong new_1 = as_type(_e14 + 10uL); - uint _e19 = i_1; - ulong _e21 = old_1; - _atomic_compare_exchange_result_Uint_8_ _e22 = naga_atomic_compare_exchange_weak_explicit(&uint(_e19) < 128 ? arr_u64_.inner[_e19] : DefaultConstructible(), _e21, new_1); - old_1 = _e22.old_value; - exchanged_1 = _e22.exchanged; - } - } - } - } - return; -} diff --git a/snapshot/testdata/golden/msl/atomicOps-int64.msl b/snapshot/testdata/golden/msl/atomicOps-int64.msl deleted file mode 100644 index d2c5ded..0000000 --- a/snapshot/testdata/golden/msl/atomicOps-int64.msl +++ /dev/null @@ -1,200 +0,0 @@ -// language: metal1.0 -#include -#include - -using metal::uint; -struct DefaultConstructible { - template - operator T() && { - return T {}; - } -}; - -struct type_2 { - metal::atomic_long inner[2]; -}; -struct Struct { - metal::atomic_ulong atomic_scalar; - type_2 atomic_arr; -}; -struct _atomic_compare_exchange_result_Uint_8_ { - ulong old_value; - bool exchanged; - char _pad2[7]; -}; -struct _atomic_compare_exchange_result_Sint_8_ { - long old_value; - bool exchanged; - char _pad2[7]; -}; - -template -_atomic_compare_exchange_result_Uint_8_ naga_atomic_compare_exchange_weak_explicit( - device A *atomic_ptr, - ulong cmp, - ulong v -) { - bool swapped = metal::atomic_compare_exchange_weak_explicit( - atomic_ptr, &cmp, v, - metal::memory_order_relaxed, metal::memory_order_relaxed - ); - return _atomic_compare_exchange_result_Uint_8_{cmp, swapped}; -} -template -_atomic_compare_exchange_result_Uint_8_ naga_atomic_compare_exchange_weak_explicit( - threadgroup A *atomic_ptr, - ulong cmp, - ulong v -) { - bool swapped = metal::atomic_compare_exchange_weak_explicit( - atomic_ptr, &cmp, v, - metal::memory_order_relaxed, metal::memory_order_relaxed - ); - return _atomic_compare_exchange_result_Uint_8_{cmp, swapped}; -} - -template -_atomic_compare_exchange_result_Sint_8_ naga_atomic_compare_exchange_weak_explicit( - device A *atomic_ptr, - long cmp, - long v -) { - bool swapped = metal::atomic_compare_exchange_weak_explicit( - atomic_ptr, &cmp, v, - metal::memory_order_relaxed, metal::memory_order_relaxed - ); - return _atomic_compare_exchange_result_Sint_8_{cmp, swapped}; -} -template -_atomic_compare_exchange_result_Sint_8_ naga_atomic_compare_exchange_weak_explicit( - threadgroup A *atomic_ptr, - long cmp, - long v -) { - bool swapped = metal::atomic_compare_exchange_weak_explicit( - atomic_ptr, &cmp, v, - metal::memory_order_relaxed, metal::memory_order_relaxed - ); - return _atomic_compare_exchange_result_Sint_8_{cmp, swapped}; -} - -struct cs_mainInput { -}; -kernel void cs_main( - metal::uint3 id [[thread_position_in_threadgroup]] -, device metal::atomic_ulong& storage_atomic_scalar [[user(fake0)]] -, device type_2& storage_atomic_arr [[user(fake0)]] -, device Struct& storage_struct [[user(fake0)]] -) { - threadgroup metal::atomic_ulong workgroup_atomic_scalar; - threadgroup type_2 workgroup_atomic_arr; - threadgroup Struct workgroup_struct; - if (metal::all(id == metal::uint3(0u))) { - metal::atomic_store_explicit(&workgroup_atomic_scalar, 0, metal::memory_order_relaxed); - for (int __i0 = 0; __i0 < 2; __i0++) { - metal::atomic_store_explicit(&workgroup_atomic_arr.inner[__i0], 0, metal::memory_order_relaxed); - } - metal::atomic_store_explicit(&workgroup_struct.atomic_scalar, 0, metal::memory_order_relaxed); - for (int __i0 = 0; __i0 < 2; __i0++) { - metal::atomic_store_explicit(&workgroup_struct.atomic_arr.inner[__i0], 0, metal::memory_order_relaxed); - } - } - metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); - metal::atomic_store_explicit(&storage_atomic_scalar, 1uL, metal::memory_order_relaxed); - metal::atomic_store_explicit(&storage_atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - metal::atomic_store_explicit(&storage_struct.atomic_scalar, 1uL, metal::memory_order_relaxed); - metal::atomic_store_explicit(&storage_struct.atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - metal::atomic_store_explicit(&workgroup_atomic_scalar, 1uL, metal::memory_order_relaxed); - metal::atomic_store_explicit(&workgroup_atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - metal::atomic_store_explicit(&workgroup_struct.atomic_scalar, 1uL, metal::memory_order_relaxed); - metal::atomic_store_explicit(&workgroup_struct.atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); - ulong l0_ = metal::atomic_load_explicit(&storage_atomic_scalar, metal::memory_order_relaxed); - long l1_ = metal::atomic_load_explicit(&storage_atomic_arr.inner[1], metal::memory_order_relaxed); - ulong l2_ = metal::atomic_load_explicit(&storage_struct.atomic_scalar, metal::memory_order_relaxed); - long l3_ = metal::atomic_load_explicit(&storage_struct.atomic_arr.inner[1], metal::memory_order_relaxed); - ulong l4_ = metal::atomic_load_explicit(&workgroup_atomic_scalar, metal::memory_order_relaxed); - long l5_ = metal::atomic_load_explicit(&workgroup_atomic_arr.inner[1], metal::memory_order_relaxed); - ulong l6_ = metal::atomic_load_explicit(&workgroup_struct.atomic_scalar, metal::memory_order_relaxed); - long l7_ = metal::atomic_load_explicit(&workgroup_struct.atomic_arr.inner[1], metal::memory_order_relaxed); - metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); - ulong _e51 = metal::atomic_fetch_add_explicit(&storage_atomic_scalar, 1uL, metal::memory_order_relaxed); - long _e55 = metal::atomic_fetch_add_explicit(&storage_atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - ulong _e59 = metal::atomic_fetch_add_explicit(&storage_struct.atomic_scalar, 1uL, metal::memory_order_relaxed); - long _e64 = metal::atomic_fetch_add_explicit(&storage_struct.atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - ulong _e67 = metal::atomic_fetch_add_explicit(&workgroup_atomic_scalar, 1uL, metal::memory_order_relaxed); - long _e71 = metal::atomic_fetch_add_explicit(&workgroup_atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - ulong _e75 = metal::atomic_fetch_add_explicit(&workgroup_struct.atomic_scalar, 1uL, metal::memory_order_relaxed); - long _e80 = metal::atomic_fetch_add_explicit(&workgroup_struct.atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); - ulong _e83 = metal::atomic_fetch_sub_explicit(&storage_atomic_scalar, 1uL, metal::memory_order_relaxed); - long _e87 = metal::atomic_fetch_sub_explicit(&storage_atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - ulong _e91 = metal::atomic_fetch_sub_explicit(&storage_struct.atomic_scalar, 1uL, metal::memory_order_relaxed); - long _e96 = metal::atomic_fetch_sub_explicit(&storage_struct.atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - ulong _e99 = metal::atomic_fetch_sub_explicit(&workgroup_atomic_scalar, 1uL, metal::memory_order_relaxed); - long _e103 = metal::atomic_fetch_sub_explicit(&workgroup_atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - ulong _e107 = metal::atomic_fetch_sub_explicit(&workgroup_struct.atomic_scalar, 1uL, metal::memory_order_relaxed); - long _e112 = metal::atomic_fetch_sub_explicit(&workgroup_struct.atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); - metal::atomic_max_explicit(&storage_atomic_scalar, 1uL, metal::memory_order_relaxed); - metal::atomic_max_explicit(&storage_atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - metal::atomic_max_explicit(&storage_struct.atomic_scalar, 1uL, metal::memory_order_relaxed); - metal::atomic_max_explicit(&storage_struct.atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - metal::atomic_max_explicit(&workgroup_atomic_scalar, 1uL, metal::memory_order_relaxed); - metal::atomic_max_explicit(&workgroup_atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - metal::atomic_max_explicit(&workgroup_struct.atomic_scalar, 1uL, metal::memory_order_relaxed); - metal::atomic_max_explicit(&workgroup_struct.atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); - metal::atomic_min_explicit(&storage_atomic_scalar, 1uL, metal::memory_order_relaxed); - metal::atomic_min_explicit(&storage_atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - metal::atomic_min_explicit(&storage_struct.atomic_scalar, 1uL, metal::memory_order_relaxed); - metal::atomic_min_explicit(&storage_struct.atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - metal::atomic_min_explicit(&workgroup_atomic_scalar, 1uL, metal::memory_order_relaxed); - metal::atomic_min_explicit(&workgroup_atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - metal::atomic_min_explicit(&workgroup_struct.atomic_scalar, 1uL, metal::memory_order_relaxed); - metal::atomic_min_explicit(&workgroup_struct.atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); - ulong _e163 = metal::atomic_fetch_and_explicit(&storage_atomic_scalar, 1uL, metal::memory_order_relaxed); - long _e167 = metal::atomic_fetch_and_explicit(&storage_atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - ulong _e171 = metal::atomic_fetch_and_explicit(&storage_struct.atomic_scalar, 1uL, metal::memory_order_relaxed); - long _e176 = metal::atomic_fetch_and_explicit(&storage_struct.atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - ulong _e179 = metal::atomic_fetch_and_explicit(&workgroup_atomic_scalar, 1uL, metal::memory_order_relaxed); - long _e183 = metal::atomic_fetch_and_explicit(&workgroup_atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - ulong _e187 = metal::atomic_fetch_and_explicit(&workgroup_struct.atomic_scalar, 1uL, metal::memory_order_relaxed); - long _e192 = metal::atomic_fetch_and_explicit(&workgroup_struct.atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); - ulong _e195 = metal::atomic_fetch_or_explicit(&storage_atomic_scalar, 1uL, metal::memory_order_relaxed); - long _e199 = metal::atomic_fetch_or_explicit(&storage_atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - ulong _e203 = metal::atomic_fetch_or_explicit(&storage_struct.atomic_scalar, 1uL, metal::memory_order_relaxed); - long _e208 = metal::atomic_fetch_or_explicit(&storage_struct.atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - ulong _e211 = metal::atomic_fetch_or_explicit(&workgroup_atomic_scalar, 1uL, metal::memory_order_relaxed); - long _e215 = metal::atomic_fetch_or_explicit(&workgroup_atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - ulong _e219 = metal::atomic_fetch_or_explicit(&workgroup_struct.atomic_scalar, 1uL, metal::memory_order_relaxed); - long _e224 = metal::atomic_fetch_or_explicit(&workgroup_struct.atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); - ulong _e227 = metal::atomic_fetch_xor_explicit(&storage_atomic_scalar, 1uL, metal::memory_order_relaxed); - long _e231 = metal::atomic_fetch_xor_explicit(&storage_atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - ulong _e235 = metal::atomic_fetch_xor_explicit(&storage_struct.atomic_scalar, 1uL, metal::memory_order_relaxed); - long _e240 = metal::atomic_fetch_xor_explicit(&storage_struct.atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - ulong _e243 = metal::atomic_fetch_xor_explicit(&workgroup_atomic_scalar, 1uL, metal::memory_order_relaxed); - long _e247 = metal::atomic_fetch_xor_explicit(&workgroup_atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - ulong _e251 = metal::atomic_fetch_xor_explicit(&workgroup_struct.atomic_scalar, 1uL, metal::memory_order_relaxed); - long _e256 = metal::atomic_fetch_xor_explicit(&workgroup_struct.atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - ulong _e259 = metal::atomic_exchange_explicit(&storage_atomic_scalar, 1uL, metal::memory_order_relaxed); - long _e263 = metal::atomic_exchange_explicit(&storage_atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - ulong _e267 = metal::atomic_exchange_explicit(&storage_struct.atomic_scalar, 1uL, metal::memory_order_relaxed); - long _e272 = metal::atomic_exchange_explicit(&storage_struct.atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - ulong _e275 = metal::atomic_exchange_explicit(&workgroup_atomic_scalar, 1uL, metal::memory_order_relaxed); - long _e279 = metal::atomic_exchange_explicit(&workgroup_atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - ulong _e283 = metal::atomic_exchange_explicit(&workgroup_struct.atomic_scalar, 1uL, metal::memory_order_relaxed); - long _e288 = metal::atomic_exchange_explicit(&workgroup_struct.atomic_arr.inner[1], 1L, metal::memory_order_relaxed); - _atomic_compare_exchange_result_Uint_8_ _e292 = naga_atomic_compare_exchange_weak_explicit(&storage_atomic_scalar, 1uL, 2uL); - _atomic_compare_exchange_result_Sint_8_ _e297 = naga_atomic_compare_exchange_weak_explicit(&storage_atomic_arr.inner[1], 1L, 2L); - _atomic_compare_exchange_result_Uint_8_ _e302 = naga_atomic_compare_exchange_weak_explicit(&storage_struct.atomic_scalar, 1uL, 2uL); - _atomic_compare_exchange_result_Sint_8_ _e308 = naga_atomic_compare_exchange_weak_explicit(&storage_struct.atomic_arr.inner[1], 1L, 2L); - _atomic_compare_exchange_result_Uint_8_ _e312 = naga_atomic_compare_exchange_weak_explicit(&workgroup_atomic_scalar, 1uL, 2uL); - _atomic_compare_exchange_result_Sint_8_ _e317 = naga_atomic_compare_exchange_weak_explicit(&workgroup_atomic_arr.inner[1], 1L, 2L); - _atomic_compare_exchange_result_Uint_8_ _e322 = naga_atomic_compare_exchange_weak_explicit(&workgroup_struct.atomic_scalar, 1uL, 2uL); - _atomic_compare_exchange_result_Sint_8_ _e328 = naga_atomic_compare_exchange_weak_explicit(&workgroup_struct.atomic_arr.inner[1], 1L, 2L); - return; -}