Is there a way where the gpu can change one variable simultaneously across multiple threads?

Question

Look at this code, main.swift:

import Metal

let device = MTLCreateSystemDefaultDevice()!
let lib = device.makeDefaultLibrary()!
let function = lib.makeFunction(name: "testout")!
let pso : MTLComputePipelineState
do {
    pso = try device.makeComputePipelineState(function: function)
}

let commandQueue = device.makeCommandQueue()!
// generate buffer
let buffer = device.makeBuffer(length: MemoryLayout<Float>.stride, options: .storageModeShared)!
// set the variable "a" in our testout.metal to the value -5
buffer.contents().assumingMemoryBound(to: Float.self).assign(repeating: -5, count: 1)

let commandBuffer = commandQueue.makeCommandBuffer()!
let computeEncoder = commandBuffer.makeComputeCommandEncoder()!
computeEncoder.setComputePipelineState(pso)
computeEncoder.setBuffer(buffer, offset: 0, index: 0)

// run it 3 times
let gridSize = MTLSizeMake(3, 1, 1)

computeEncoder.dispatchThreads(gridSize, threadsPerThreadgroup: gridSize)
computeEncoder.endEncoding()
commandBuffer.commit()
commandBuffer.waitUntilCompleted()

print(buffer.contents().assumingMemoryBound(to: Float.self).pointee)

testout.metal:

#include <metal_stdlib>
using namespace metal;

kernel void testout (device float* a) {
    a[0] += 2;
}

Is there a function / program that can output 1? Is it even possible? (It runs 3 times, original value is -5, metal increases a by 2, -5 + (3 * 2) = 1)

Any answer would be appreciated :)

You would use atomic operations. See [here](https://stackoverflow.com/a/47215871/1312143). — Ken Thomases, May 25 '20 at 04:45
Thanks, but can you set the initial value to values less than 0? — Eshaan Barkataki, May 25 '20 at 23:58
If you use signed types (`atomic_int` in the shader, and `Int32` for the buffer types in Swift), then yes. — Ken Thomases, May 26 '20 at 02:11
Thanks, but is there a piece of apple documentation that states `Int32` works rather than `Int` for atomic operations? If not, how did you find out? (Just want to get better at figuring out problems myself) — Eshaan Barkataki, May 26 '20 at 22:04
As mentioned in my linked answer, it's not documented what the in-memory representation of the atomic types is. The [Metal Shading Language spec](https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf) does document that `int` is a 32-bit type. Swift's `Int` is not always 32-bit. — Ken Thomases, May 27 '20 at 03:18

Eshaan Barkataki · Accepted Answer · 2020-05-26T21:34:21.127

0

Thanks to Ken Thomases, all I had to do is to change testout.metal into this:

kernel void testout (device atomic_int &a) {
    atomic_fetch_add_explicit(&a, 2, memory_order_relaxed);
}

And in the main script, I had to change the datatype of the buffer into Int32.

edited May 26 '20 at 21:34

answered May 25 '20 at 23:58

Eshaan Barkataki

35
1
7

Is there a way where the gpu can change one variable simultaneously across multiple threads?

1 Answers1