set a boolean from multiple threads

Question

I have a situation where if a condition happens in any of the threadgroup threads, then a boolean should be set and a final task performed by a single thread.

Does this look like the correct way to accomplish that?

threadgroup uint myBool; // threadgroup bool crashes the compiler

// Initialize myBool
if(tid == 0)
{
   myBool = false;
}

// Ensure all threads see value for myBool
threadgroup_barrier(mem_flags::mem_threadgroup);

// Do various caluclations
...

if(someFancyCondition)
{
   myBool = true;
}

// Ensure correct final value for myBool
threadgroup_barrier(mem_flags::mem_threadgroup);

if(tid == 0 and myBool)
{
   // do stuff
}

Metal

856

Posted by

Audulus

Reply

Add a Comment

Answer 1

Hello

To be honest, I don't think it is safe enough. Had some problems in the past with more than one thread accessing single variable at the same time. If I remember well, it sometimes caused garbage being written to the said variable, i.e. one thread wrote 1, another wrote 2, and I from time to time I saw some big numbers in there.

I guess that to be safe, you should either 1) Make all threads set their own flag and then perform parallel reduction on that values or 2) I guess this is what atomics are for. Unfortunately I never used them, so can't offer sound advice, but description is in Metal Shading Language Specification, 5.12

Regards

Michal

PS. I had some crashes when using bools too, submitted rdars and these got fixed. Perhaps you should do the same?

Posted by

MikeAlpha

Add a Comment

Answer 2

Thanks for your help Michal!

Sadly, when I try to add atomics, the compiler crashes as well (on the device, so it's harder to submit a bug report). Sigh.

I'll try the parallel reduction technique.

Posted by

Audulus

Add a Comment

set a boolean from multiple threads

Replies