Initial commit

QuickSilver/CUDA/src/AtomicMacro.hh

@@ -0,0 +1,184 @@
+/*
+Modifications Copyright (C) 2023 Intel Corporation
+
+Redistribution and use in source and binary forms, with or without modification,
+are permitted provided that the following conditions are met:
+
+1. Redistributions of source code must retain the above copyright notice,
+   this list of conditions and the following disclaimer.
+2. Redistributions in binary form must reproduce the above copyright notice,
+   this list of conditions and the following disclaimer in the documentation
+   and/or other materials provided with the distribution.
+3. Neither the name of the copyright holder nor the names of its contributors
+   may be used to endorse or promote products derived from this software
+   without specific prior written permission.
+
+THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO,
+THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+ARE DISCLAIMED.  IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS
+BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY,
+OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT
+OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
+OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY,
+WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE
+OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE,
+EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+
+SPDX-License-Identifier: BSD-3-Clause
+*/
+
+/*
+Copyright 2019 Advanced Micro Devices
+
+Redistribution and use in source and binary forms, with or without modification, are permitted provided that the following conditions are met:
+
+1. Redistributions of source code must retain the above copyright notice, this list of conditions and the following disclaimer.
+
+2. Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following disclaimer in the documentation and/or other materials provided with the distribution.
+
+3. Neither the name of the copyright holder nor the names of its contributors may be used to endorse or promote products derived from this software without specific prior written permission.
+
+THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+*/
+
+// Determine which atomics to use based on platform being compiled for
+//
+
+#ifndef ATOMICS_HD
+#define ATOMICS_HD
+
+#include <thread>
+#include <mutex>
+#include <algorithm>
+
+inline __device__ double ull2d(const unsigned long long &val)
+{
+    return *((double *)&val);
+}
+
+#ifdef HAVE_CUDA
+#include <cuda.h>
+#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 600
+
+#else
+inline __device__ double atomicAdd(double *pointer, double val)
+{
+    // A workaround dealing with the fact that atomic doubles don't work with all versions of CUDA.
+    unsigned long long int *pointer_as_p2ull = (unsigned long long int *)pointer;
+    unsigned long long int old = *pointer_as_p2ull, check_value;
+    do
+    {
+        check_value = old;
+        old = atomicCAS(pointer_as_p2ull, check_value, __double_as_longlong(val + ull2d(check_value)));
+    } while (check_value != old);
+    return ull2d(old);
+};
+#endif
+
+#endif
+
+#ifdef HAVE_OPENMP
+#define USE_OPENMP_ATOMICS
+#elif HAVE_OPENMP_TARGET
+#define USE_OPENMP_ATOMICS
+#endif
+
+#if defined(HAVE_CUDA)
+
+// If in a CUDA GPU section use the CUDA atomics
+#ifdef __CUDA_ARCH__
+
+// Currently not atomic here. But its only used when it does not necissarially need to be atomic.
+#define ATOMIC_WRITE(x, v) \
+    x = v;
+
+#define ATOMIC_ADD(x, v) \
+    atomicAdd(&x, v);
+
+#define ATOMIC_UPDATE(x) \
+    atomicAdd(&x, 1);
+
+#define ATOMIC_CAPTURE(x, v, p) \
+    p = atomicAdd(&x, v);
+// If in a CPU OpenMP section use the OpenMP atomics
+#elif defined(USE_OPENMP_ATOMICS)
+#define ATOMIC_WRITE(x, v)      \
+    _Pragma("omp atomic write") \
+        x = v;
+
+#define ATOMIC_ADD(x, v)  \
+    _Pragma("omp atomic") \
+        x += v;
+
+#define ATOMIC_UPDATE(x)         \
+    _Pragma("omp atomic update") \
+        x++;
+
+#define ATOMIC_CAPTURE(x, v, p)   \
+    _Pragma("omp atomic capture") \
+    {                             \
+        p = x;                    \
+        x = x + v;                \
+    }
+
+// If in a serial section, no need to use atomics
+#else
+#define ATOMIC_WRITE(x, v) \
+    x = v;
+
+#define ATOMIC_UPDATE(x) \
+    x++;
+
+#define ATOMIC_ADD(x, v) \
+    x += v;
+
+#define ATOMIC_CAPTURE(x, v, p) \
+    {                           \
+        p = x;                  \
+        x = x + v;              \
+    }
+
+#endif
+
+// If in a OpenMP section use the OpenMP atomics
+#elif defined(USE_OPENMP_ATOMICS)
+#define ATOMIC_WRITE(x, v)      \
+    _Pragma("omp atomic write") \
+        x = v;
+
+#define ATOMIC_ADD(x, v)  \
+    _Pragma("omp atomic") \
+        x += v;
+
+#define ATOMIC_UPDATE(x)         \
+    _Pragma("omp atomic update") \
+        x++;
+
+#define ATOMIC_CAPTURE(x, v, p)   \
+    _Pragma("omp atomic capture") \
+    {                             \
+        p = x;                    \
+        x = x + v;                \
+    }
+
+// If in a serial section, no need to use atomics
+#else
+#define ATOMIC_WRITE(x, v) \
+    x = v;
+
+#define ATOMIC_UPDATE(x) \
+    x++;
+
+#define ATOMIC_ADD(x, v) \
+    x += v;
+
+#define ATOMIC_CAPTURE(x, v, p) \
+    {                           \
+        p = x;                  \
+        x = x + v;              \
+    }
+#endif
+
+#endif

QuickSilver/CUDA/src/BulkStorage.hh

@@ -0,0 +1,98 @@
+#ifndef BULK_STORAGE_HH
+#define BULK_STORAGE_HH
+
+#include "MemoryControl.hh"
+
+template <typename T>
+class BulkStorage
+{
+ public:
+   BulkStorage()
+   : _bulkStorage(0),
+     _refCount(0),
+     _size(0),
+     _capacity(0),
+     _memPolicy(MemoryControl::AllocationPolicy::UNDEFINED_POLICY)
+   {
+      _refCount = new int;
+      *_refCount = 1;
+   }
+
+   BulkStorage(const BulkStorage& aa)
+   : _bulkStorage(aa._bulkStorage), _refCount(aa._refCount), _size(aa._size), _capacity(aa._capacity), _memPolicy(aa._memPolicy)
+   {
+      ++(*_refCount);
+   }
+
+   ~BulkStorage()
+   {
+      // Check for instances that never allocated memory.
+      // I'm not exactly sure how this can happen, but it does.
+      if (_bulkStorage == 0) 
+         return;
+
+      --(*_refCount);
+      if (*_refCount > 0)
+         return;
+
+      MemoryControl::deallocate(_bulkStorage, _capacity, _memPolicy);
+      delete _refCount;
+   }
+
+   /// Needed for copy-swap idiom
+   void swap(BulkStorage<T>& other)
+   {
+      std::swap(_bulkStorage, other._bulkStorage);
+      std::swap(_refCount,    other._refCount);
+      std::swap(_size,        other._size);
+      std::swap(_capacity,    other._capacity);
+      std::swap(_memPolicy,   other._memPolicy);
+   }
+
+   /// Implement assignment using copy-swap idiom
+   BulkStorage& operator=(const BulkStorage& aa)
+   {
+      if (&aa != this)
+      {
+         BulkStorage<T> temp(aa);
+         this->swap(temp);
+      }
+      return *this;
+   }
+
+   void setCapacity(int capacity, MemoryControl::AllocationPolicy policy)
+   {
+      qs_assert(_bulkStorage == 0);
+      _bulkStorage = MemoryControl::allocate<T>(capacity, policy);
+      _capacity = capacity;
+      _memPolicy = policy;
+   }
+
+   T* getBlock(int nItems)
+   {
+      T* blockStart = _bulkStorage + _size;
+      _size += nItems;
+      qs_assert(_size <= _capacity);
+      return blockStart;
+   }
+
+
+ private:
+
+   // This class doesn't have well defined copy semantics.  However,
+   // just disabling copy operations breaks the build since we haven't
+   // been consistent about dealing with copy semantics in classes like
+   // MC_Mesh_Domain.
+
+
+
+   T* _bulkStorage;
+   int* _refCount;
+   int _size;
+   int _capacity;
+   MemoryControl::AllocationPolicy _memPolicy;
+
+};
+
+
+#endif