diff options
| author | jsmall-nvidia <jsmall@nvidia.com> | 2020-02-14 15:06:35 -0500 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2020-02-14 15:06:35 -0500 |
| commit | 2c097545eaa324a91a035327abad2e8b4fa60469 (patch) | |
| tree | 95fd3890f2bfb0184ddbc7f1008de30698651473 /source/slang/hlsl.meta.slang | |
| parent | dfd3d263704445b6dcebea54dc47193897548822 (diff) | |
Feature/cuda coverage (#1223)
* Add cubemap support.
* Add CUDA fence instrinsics.
* Added Gather for CUDA.
* Use the CUDA driver API as much as possible.
* * Support 1D texture on CPU
* WIP on 1D texture on CUDA
* Added simplified texture test
* Fix test.
* Improve texture-simple tests.
Co-authored-by: Tim Foley <tfoleyNV@users.noreply.github.com>
Diffstat (limited to 'source/slang/hlsl.meta.slang')
| -rw-r--r-- | source/slang/hlsl.meta.slang | 5 |
1 files changed, 3 insertions, 2 deletions
diff --git a/source/slang/hlsl.meta.slang b/source/slang/hlsl.meta.slang index f8ae340bc..f7707cc6d 100644 --- a/source/slang/hlsl.meta.slang +++ b/source/slang/hlsl.meta.slang @@ -347,8 +347,7 @@ __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()") +__target_intrinsic(cuda, "__threadfence()") void AllMemoryBarrier(); // Thread-group sync and barrier for writes to all memory spaces (HLSL SM 5.0) @@ -648,6 +647,7 @@ __generic<T : __BuiltinFloatingPointType, let N : int> T determinant(matrix<T,N, // Barrier for device memory __target_intrinsic(glsl, "memoryBarrier(), memoryBarrierImage(), memoryBarrierBuffer()") +__target_intrinsic(cuda, "__threadfence()") void DeviceMemoryBarrier(); __target_intrinsic(glsl, "memoryBarrier(), memoryBarrierImage(), memoryBarrierBuffer(), barrier()") @@ -814,6 +814,7 @@ float2 GetRenderTargetSamplePosition(int Index); // Group memory barrier __target_intrinsic(glsl, "groupMemoryBarrier") +__target_intrinsic(cuda, "__threadfence_block") void GroupMemoryBarrier(); |