CUDA barrier/atomic support (#1218)

* * Improved fastRemoveAt
* Fixed off by one bug
* Fixed const safeness with List<>

* Made List begin and end const safe.

* Revert to previous RefPtr usage.

* Fix bug with casting.

* Tabs -> spaces.
Small fixes/improvements to List.

* Improve comment on List.

* Group shared/atomic test works on CUDA.

* * Enabled CUDA tests for atomics tests
* Enabled DX12 test for atomics-buffer.slang

Not clear just yet how to implement that for CUDA - it will work with StructuredBuffer.

* hasContent -> isNonEmpty

* Remove unneeded comment.

1 files changed, 37 insertions, 1 deletions

diff --git a/source/slang/hlsl.meta.slang b/source/slang/hlsl.meta.slang
index a646c1b77..669fbb440 100644
--- a/source/slang/hlsl.meta.slang
+++ b/source/slang/hlsl.meta.slang
@@ -347,10 +347,13 @@ __generic<T : __BuiltinType, let N : int, let M : int> bool all(matrix<T,N,M> x)
 
 // Barrier for writes to all memory spaces (HLSL SM 5.0)
 __target_intrinsic(glsl, "memoryBarrier(), groupMemoryBarrier(), memoryBarrierImage(), memoryBarrierBuffer()")
+// TODO(JS): Doesn't seem to be weaker form of sync, so use this?
+__target_intrinsic(cuda, "__syncthreads()")
 void AllMemoryBarrier();
 
 // Thread-group sync and barrier for writes to all memory spaces (HLSL SM 5.0)
 __target_intrinsic(glsl, "memoryBarrier(), groupMemoryBarrier(), memoryBarrierImage(), memoryBarrierBuffer(), barrier()")
+__target_intrinsic(cuda, "__syncthreads()")
 void AllMemoryBarrierWithGroupSync();
 
 // Test if any components is non-zero (HLSL SM 1.0)
@@ -648,6 +651,7 @@ __target_intrinsic(glsl, "memoryBarrier(), memoryBarrierImage(), memoryBarrierBu
 void DeviceMemoryBarrier();
 
 __target_intrinsic(glsl, "memoryBarrier(), memoryBarrierImage(), memoryBarrierBuffer(), barrier()")
+__target_intrinsic(glsl, "__syncthreads()")
 void DeviceMemoryBarrierWithGroupSync();
 
 // Vector distance
@@ -812,99 +816,131 @@ float2 GetRenderTargetSamplePosition(int Index);
 __target_intrinsic(glsl, "groupMemoryBarrier")
 void GroupMemoryBarrier();
 
+
 __target_intrinsic(glsl, "groupMemoryBarrier(), barrier()")
+__target_intrinsic(cuda, "__syncthreads()")
 void GroupMemoryBarrierWithGroupSync();
 
 // Atomics
 
 __target_intrinsic(glsl, "$atomicAdd($A, $1)")
+__target_intrinsic(cuda, "atomicAdd(&$0, $1)")
 void InterlockedAdd(__ref  int dest,  int value);
 
 __target_intrinsic(glsl, "$atomicAdd($A, $1)")
+__target_intrinsic(cuda, "atomicAdd((uint*)&$0, $1)")
 void InterlockedAdd(__ref uint dest, uint value);
 
 __target_intrinsic(glsl, "($2 = $atomicAdd($A, $1))")
+__target_intrinsic(cuda, "($2 = atomicAdd(&$0, $1))")
 void InterlockedAdd(__ref  int dest,  int value, out  int original_value);
 
 __target_intrinsic(glsl, "($2 = $atomicAdd($A, $1))")
+__target_intrinsic(cuda, "($2 = (uint)atomicAdd((uint*)&$0, $1))")
 void InterlockedAdd(__ref uint dest, uint value, out uint original_value);
 
 __target_intrinsic(glsl, "$atomicAnd($A, $1)")
+__target_intrinsic(cuda, "atomicAnd(&$0, $1)")
 void InterlockedAnd(__ref  int dest,  int value);
 
 __target_intrinsic(glsl, "$atomicAnd($A, $1)")
+__target_intrinsic(cuda, "atomicAnd((int*)&$0, $1)")
 void InterlockedAnd(__ref uint dest, uint value);
 
 __target_intrinsic(glsl, "($2 = $atomicAnd($A, $1))")
+__target_intrinsic(cuda, "($2 = atomicAnd(&$0, $1))")
 void InterlockedAnd(__ref  int dest,  int value, out  int original_value);
 
 __target_intrinsic(glsl, "($2 = $atomicAnd($A, $1))")
+__target_intrinsic(cuda, "($2 = atomicAnd((int*)&$0, $1))")
 void InterlockedAnd(__ref uint dest, uint value, out uint original_value);
 
 __target_intrinsic(glsl, "($3 = $atomicCompSwap($A, $1, $2))")
+__target_intrinsic(cuda, "($3 = atomicCAS(&$0, $1, $2))")
 void InterlockedCompareExchange(__ref  int dest,  int compare_value,  int value, out  int original_value);
 
 __target_intrinsic(glsl, "($3 = $atomicCompSwap($A, $1, $2))")
+__target_intrinsic(cuda, "($3 = (uint)atomicCAS((int*)&$0, $1, $2))")
 void InterlockedCompareExchange(__ref uint dest, uint compare_value, uint value, out uint original_value);
 
 __target_intrinsic(glsl, "$atomicCompSwap($A, $1, $2)")
-void InterlockedCompareStore(__ref  int dest,  int compare_value,  int value);
+__target_intrinsic(cuda, "atomicCAS(&$0, $1, $2)")
+void InterlockedCompareStore(__ref int dest,  int compare_value,  int value);
 
 __target_intrinsic(glsl, "$atomicCompSwap($A, $1, $2)")
+__target_intrinsic(cuda, "atomicCAS((int*)&$0, $1, $2)")
 void InterlockedCompareStore(__ref uint dest, uint compare_value, uint value);
 
 __target_intrinsic(glsl, "($2 = $atomicExchange($A, $1))")
+__target_intrinsic(cuda, "($2 = atomicExch(&$0, $1))")
 void InterlockedExchange(__ref  int dest,  int value, out  int original_value);
 
 __target_intrinsic(glsl, "($2 = $atomicExchange($A, $1))")
+__target_intrinsic(cuda, "($2 = (uint)atomicExch((int*)&$0, $1))")
 void InterlockedExchange(__ref uint dest, uint value, out uint original_value);
 
 __target_intrinsic(glsl, "$atomicMax($A, $1)")
+__target_intrinsic(cuda, "atomicMax(&$0, $1)")
 void InterlockedMax(__ref  int dest,  int value);
 
 __target_intrinsic(glsl, "$atomicMax($A, $1)")
+__target_intrinsic(cuda, "atomicMax((int*)&$0, $1)")
 void InterlockedMax(__ref uint dest, uint value);
 
 __target_intrinsic(glsl, "($2 = $atomicMax($A, $1))")
+__target_intrinsic(cuda, "($2 = atomicMax(&$0, $1))")
 void InterlockedMax(__ref  int dest,  int value, out  int original_value);
 
 __target_intrinsic(glsl, "($2 = $atomicMax($A, $1))")
+__target_intrinsic(cuda, "($2 = (uint)atomicMax((int*)&$0, $1))")
 void InterlockedMax(__ref uint dest, uint value, out uint original_value);
 
 __target_intrinsic(glsl, "$atomicMin($A, $1)")
+__target_intrinsic(cuda, "atomicMin(&$0, $1)")
 void InterlockedMin(__ref  int dest,  int value);
 
 __target_intrinsic(glsl, "$atomicMin($A, $1)")
+__target_intrinsic(cuda, "atomicMin((int*)&$0, $1)")
 void InterlockedMin(__ref uint dest, uint value);
 
 __target_intrinsic(glsl, "($2 = $atomicMin($A, $1))")
+__target_intrinsic(cuda, "($2 = atomicMin(&$0, $1))")
 void InterlockedMin(__ref  int dest,  int value, out  int original_value);
 
 __target_intrinsic(glsl, "($2 = $atomicMin($A, $1))")
+__target_intrinsic(cuda, "($2 = (uint)atomicMin((int*)&$0, $1))")
 void InterlockedMin(__ref uint dest, uint value, out uint original_value);
 
 __target_intrinsic(glsl, "$atomicOr($A, $1)")
+__target_intrinsic(cuda, "atomicOr(&$0, $1)")
 void InterlockedOr(__ref  int dest,  int value);
 
 __target_intrinsic(glsl, "$atomicOr($A, $1)")
+__target_intrinsic(cuda, "atomicOr((int*)&$0, $1)")
 void InterlockedOr(__ref uint dest, uint value);
 
 __target_intrinsic(glsl, "($2 = $atomicOr($A, $1))")
+__target_intrinsic(cuda, "($2 = atomicOr(&$0, $1))")
 void InterlockedOr(__ref  int dest,  int value, out  int original_value);
 
 __target_intrinsic(glsl, "($2 = $atomicOr($A, $1))")
+__target_intrinsic(cuda, "($2 = (uint)atomicOr((int*)&$0, $1))")
 void InterlockedOr(__ref uint dest, uint value, out uint original_value);
 
 __target_intrinsic(glsl, "$atomicXor($A, $1)")
+__target_intrinsic(cuda, "atomicXor(&$0, $1)")
 void InterlockedXor(__ref  int dest,  int value);
 
 __target_intrinsic(glsl, "$atomicXor($A, $1)")
+__target_intrinsic(cuda, "atomicXor((int*)&$0, $1)")
 void InterlockedXor(__ref uint dest, uint value);
 
 __target_intrinsic(glsl, "($2 = $atomicXor($A, $1))")
+__target_intrinsic(cuda, "($2 = atomicXor(&$0, $1))")
 void InterlockedXor(__ref  int dest,  int value, out  int original_value);
 
 __target_intrinsic(glsl, "($2 = $atomicXor($A, $1))")
+__target_intrinsic(cuda, "($2 = (uint)atomicXor((int*)&$0, $1))")
 void InterlockedXor(__ref uint dest, uint value, out uint original_value);
 
 // Is floating-point value finite?