Welcome to mirror list, hosted at ThFree Co, Russian Federation.

github.com/torch/cutorch.git - Unnamed repository; edit this file 'description' to name the repository.
summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat (limited to 'lib/THC/THCAtomics.cuh')
-rw-r--r--lib/THC/THCAtomics.cuh7
1 files changed, 7 insertions, 0 deletions
diff --git a/lib/THC/THCAtomics.cuh b/lib/THC/THCAtomics.cuh
index 7a0be48..400875c 100644
--- a/lib/THC/THCAtomics.cuh
+++ b/lib/THC/THCAtomics.cuh
@@ -102,9 +102,16 @@ static inline __device__ void atomicAdd(half *address, half val) {
do {
assumed = old;
+#if CUDA_VERSION < 9000
half hsum;
hsum.x = (size_t)address & 2 ? (old >> 16) : (old & 0xffff);
hsum = THCNumerics<half>::add(hsum, val);
+#else
+ __half_raw hsum;
+ hsum.x = (size_t)address & 2 ? (old >> 16) : (old & 0xffff);
+ half tmpres = THCNumerics<half>::add(hsum, val);
+ hsum = __half_raw(tmpres);
+#endif
old = (size_t)address & 2 ? (old & 0xffff) | (hsum.x << 16) : (old & 0xffff0000) | hsum.x;
old = atomicCAS(address_as_ui, assumed, old);
} while (assumed != old);