initial commit, 4.5 stable

thirdparty/embree/kernels/common/buffer.h

@@ -0,0 +1,421 @@
+// Copyright 2009-2021 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+
+#pragma once
+
+#include "default.h"
+#include "device.h"
+
+namespace embree
+{
+  enum class BufferDataPointerType {
+    HOST = 0,
+    DEVICE = 1,
+    UNKNOWN = 2
+  };
+
+  /*! Implements an API data buffer object. This class may or may not own the data. */
+  class Buffer : public RefCount
+  {
+  private:
+    char* alloc(void* ptr_in, bool &shared, EmbreeMemoryType memoryType)
+    {
+      if (ptr_in)
+      {
+        shared = true;
+        return (char*)ptr_in;
+      }
+      else
+      {
+        shared = false;
+        device->memoryMonitor(this->bytes(), false);
+        size_t b = (this->bytes()+15) & ssize_t(-16);
+        return (char*)device->malloc(b,16,memoryType);
+      }
+    }
+
+  public:
+    Buffer(Device* device, size_t numBytes_in, void* ptr_in)
+      : device(device), numBytes(numBytes_in)
+    {
+      device->refInc();
+
+      ptr = alloc(ptr_in, shared, EmbreeMemoryType::USM_SHARED);
+#if defined(EMBREE_SYCL_SUPPORT)
+      dshared = true;
+      dptr = ptr;
+      modified = true;
+#endif
+    }
+
+    Buffer(Device* device, size_t numBytes_in, void* ptr_in, void* dptr_in)
+      : device(device), numBytes(numBytes_in)
+    {
+      device->refInc();
+
+#if defined(EMBREE_SYCL_SUPPORT)
+      modified = true;
+      if (device->is_gpu() && !device->has_unified_memory())
+      {
+        ptr  = alloc( ptr_in,  shared, EmbreeMemoryType::MALLOC);
+        dptr = alloc(dptr_in, dshared, EmbreeMemoryType::USM_DEVICE);
+      }
+      else if (device->is_gpu() && device->has_unified_memory())
+      {
+        ptr = alloc(ptr_in, shared, EmbreeMemoryType::USM_SHARED);
+
+        if (device->get_memory_type(ptr) != EmbreeMemoryType::USM_SHARED)
+        {
+          dptr = alloc(dptr_in, dshared, EmbreeMemoryType::USM_DEVICE);
+        }
+        else
+        {
+          dshared = true;
+          dptr = ptr;
+        }
+      }
+      else
+#endif
+      {
+        ptr = alloc(ptr_in, shared, EmbreeMemoryType::MALLOC);
+#if defined(EMBREE_SYCL_SUPPORT)
+        dshared = true;
+        dptr = ptr;
+#endif
+      }
+    }
+
+    /*! Buffer destruction */
+    virtual ~Buffer() {
+      free();
+      device->refDec();
+    }
+
+    /*! this class is not copyable */
+  private:
+    Buffer(const Buffer& other) DELETED; // do not implement
+    Buffer& operator =(const Buffer& other) DELETED; // do not implement
+
+  public:
+
+    /*! frees the buffer */
+    virtual void free()
+    {
+      if (!shared && ptr) {
+#if defined(EMBREE_SYCL_SUPPORT)
+        if (dptr == ptr) {
+          dptr = nullptr;
+        }
+#endif
+        device->free(ptr);
+        device->memoryMonitor(-ssize_t(this->bytes()), true);
+        ptr = nullptr;
+      }
+#if defined(EMBREE_SYCL_SUPPORT)
+      if (!dshared && dptr) {
+        device->free(dptr);
+        device->memoryMonitor(-ssize_t(this->bytes()), true);
+        dptr = nullptr;
+      }
+#endif
+    }
+
+    /*! gets buffer pointer */
+    void* data()
+    {
+      /* report error if buffer is not existing */
+      if (!device)
+        throw_RTCError(RTC_ERROR_INVALID_ARGUMENT, "invalid buffer specified");
+
+      /* return buffer */
+      return ptr;
+    }
+
+    /*! gets buffer pointer */
+    void* dataDevice()
+    {
+      /* report error if buffer is not existing */
+      if (!device)
+        throw_RTCError(RTC_ERROR_INVALID_ARGUMENT, "invalid buffer specified");
+
+      /* return buffer */
+#if defined(EMBREE_SYCL_SUPPORT)
+      return dptr;
+#else
+      return ptr;
+#endif
+    }
+
+    /*! returns pointer to first element */
+    __forceinline char* getPtr(BufferDataPointerType type) const
+    {
+      if (type == BufferDataPointerType::HOST) return getHostPtr();
+      else if (type == BufferDataPointerType::DEVICE) return getDevicePtr();
+
+      throw_RTCError(RTC_ERROR_INVALID_ARGUMENT, "invalid buffer data pointer type specified");
+      return nullptr;
+    }
+
+    /*! returns pointer to first element */
+    __forceinline virtual char* getHostPtr() const {
+      return ptr;
+    }
+
+    /*! returns pointer to first element */
+    __forceinline virtual char* getDevicePtr() const {
+#if defined(EMBREE_SYCL_SUPPORT)
+      return dptr;
+#else
+      return ptr;
+#endif
+    }
+
+    /*! returns the number of bytes of the buffer */
+    __forceinline size_t bytes() const {
+      return numBytes;
+    }
+
+    /*! returns true of the buffer is not empty */
+    __forceinline operator bool() const {
+      return ptr;
+    }
+
+    __forceinline void commit() {
+#if defined(EMBREE_SYCL_SUPPORT)
+      DeviceGPU* gpu_device = dynamic_cast<DeviceGPU*>(device);
+      if (gpu_device) {
+        sycl::queue queue(gpu_device->getGPUDevice());
+        commit(queue);
+        queue.wait_and_throw();
+      }
+      modified = false;
+#endif
+    }
+
+#if defined(EMBREE_SYCL_SUPPORT)
+    __forceinline sycl::event commit(sycl::queue queue) {
+      if (dptr == ptr)
+        return sycl::event();
+
+      modified = false;
+      return queue.memcpy(dptr, ptr, numBytes);
+    }
+#endif
+
+    __forceinline bool needsCommit() const {
+#if defined(EMBREE_SYCL_SUPPORT)
+     return (dptr == ptr) ? false : modified;
+#else
+      return false;
+#endif
+    }
+
+    __forceinline void setNeedsCommit(bool isModified = true) {
+#if defined(EMBREE_SYCL_SUPPORT)
+      modified = isModified;
+#endif
+    }
+
+    __forceinline void commitIfNeeded() {
+      if (needsCommit()) {
+        commit();
+      }
+    }
+
+  public:
+    Device* device;      //!< device to report memory usage to
+    size_t numBytes;     //!< number of bytes in the buffer
+    char* ptr;           //!< pointer to buffer data
+#if defined(EMBREE_SYCL_SUPPORT)
+    char* dptr;          //!< pointer to buffer data on device
+#endif
+    bool shared;         //!< set if memory is shared with application
+#if defined(EMBREE_SYCL_SUPPORT)
+    bool dshared;        //!< set if device memory is shared with application
+    bool modified;       //!< to be set when host memory has been modified and dev needs update
+#endif
+  };
+
+  /*! An untyped contiguous range of a buffer. This class does not own the buffer content. */
+  class RawBufferView
+  {
+  public:
+    /*! Buffer construction */
+    RawBufferView()
+      : ptr_ofs(nullptr), dptr_ofs(nullptr), stride(0), num(0), format(RTC_FORMAT_UNDEFINED), modCounter(1), modified(true), userData(0) {}
+
+  public:
+    /*! sets the buffer view */
+    void set(const Ref<Buffer>& buffer_in, size_t offset_in, size_t stride_in, size_t num_in, RTCFormat format_in)
+    {
+      if ((offset_in + stride_in * num_in) > (stride_in * buffer_in->numBytes))
+        throw_RTCError(RTC_ERROR_INVALID_ARGUMENT, "buffer range out of bounds");
+
+      ptr_ofs = buffer_in->getHostPtr() + offset_in;
+      dptr_ofs = buffer_in->getDevicePtr() + offset_in;
+      stride = stride_in;
+      num = num_in;
+      format = format_in;
+      modCounter++;
+      modified = true;
+      buffer = buffer_in;
+    }
+
+    /*! returns pointer to the i'th element */
+    __forceinline char* getPtr(BufferDataPointerType pointerType) const
+    {
+      if (pointerType == BufferDataPointerType::HOST)
+        return ptr_ofs;
+      else if (pointerType == BufferDataPointerType::DEVICE)
+        return dptr_ofs;
+
+      throw_RTCError(RTC_ERROR_INVALID_ARGUMENT, "invalid buffer data pointer type specified");
+      return nullptr;
+    }
+
+    /*! returns pointer to the first element */
+    __forceinline char* getPtr() const {
+      #if defined(__SYCL_DEVICE_ONLY__)
+        return dptr_ofs;
+      #else
+        return ptr_ofs;
+      #endif
+    }
+
+    /*! returns pointer to the i'th element */
+    __forceinline char* getPtr(size_t i) const
+    {
+      #if defined(__SYCL_DEVICE_ONLY__)
+        assert(i<num);
+        return dptr_ofs + i*stride;
+      #else
+        return ptr_ofs + i*stride;
+      #endif
+    }
+
+    /*! returns the number of elements of the buffer */
+    __forceinline size_t size() const {
+      return num;
+    }
+
+    /*! returns the number of bytes of the buffer */
+    __forceinline size_t bytes() const {
+      return num*stride;
+    }
+
+    /*! returns the buffer stride */
+    __forceinline unsigned getStride() const
+    {
+      assert(stride <= unsigned(inf));
+      return unsigned(stride);
+    }
+
+    /*! return the buffer format */
+    __forceinline RTCFormat getFormat() const {
+      return format;
+    }
+
+    /*! mark buffer as modified or unmodified */
+    __forceinline void setModified() {
+      modCounter++;
+      modified = true;
+      if (buffer) buffer->setNeedsCommit();
+    }
+
+    /*! mark buffer as modified or unmodified */
+    __forceinline bool isModified(unsigned int otherModCounter) const {
+      return modCounter > otherModCounter;
+    }
+
+     /*! mark buffer as modified or unmodified */
+    __forceinline bool isLocalModified() const {
+      return modified;
+    }
+
+    /*! clear local modified flag */
+    __forceinline void clearLocalModified() {
+      modified = false;
+    }
+
+    /*! returns true of the buffer is not empty */
+    __forceinline operator bool() const { 
+      return ptr_ofs;
+    }
+
+    /*! checks padding to 16 byte check, fails hard */
+    __forceinline void checkPadding16() const
+    {
+      if (ptr_ofs && num)
+        volatile int MAYBE_UNUSED w = *((int*)getPtr(size()-1)+3); // FIXME: is failing hard avoidable?
+    }
+
+  public:
+    char* ptr_ofs;      //!< base pointer plus offset
+    char* dptr_ofs;     //!< base pointer plus offset in device memory
+    size_t stride;      //!< stride of the buffer in bytes
+    size_t num;         //!< number of elements in the buffer
+    RTCFormat format;   //!< format of the buffer
+    unsigned int modCounter; //!< version ID of this buffer
+    bool modified;      //!< local modified data
+    int userData;       //!< special data
+    Ref<Buffer> buffer; //!< reference to the parent buffer
+  };
+
+  /*! A typed contiguous range of a buffer. This class does not own the buffer content. */
+  template<typename T>
+  class BufferView : public RawBufferView
+  {
+  public:
+    typedef T value_type;
+
+#if defined(__SYCL_DEVICE_ONLY__)
+    /*! access to the ith element of the buffer */
+    __forceinline       T& operator [](size_t i)       { assert(i<num); return *(T*)(dptr_ofs + i*stride); }
+    __forceinline const T& operator [](size_t i) const { assert(i<num); return *(T*)(dptr_ofs + i*stride); }
+#else
+    /*! access to the ith element of the buffer */
+    __forceinline       T& operator [](size_t i)       { assert(i<num); return *(T*)(ptr_ofs + i*stride); }
+    __forceinline const T& operator [](size_t i) const { assert(i<num); return *(T*)(ptr_ofs + i*stride); }
+#endif
+  };
+
+  template<>
+  class BufferView<Vec3fa> : public RawBufferView
+  {
+  public:
+    typedef Vec3fa value_type;
+
+#if defined(EMBREE_SYCL_SUPPORT) && defined(__SYCL_DEVICE_ONLY__)
+
+     /*! access to the ith element of the buffer */
+    __forceinline const Vec3fa operator [](size_t i) const
+    {
+      assert(i<num);
+      return Vec3fa::loadu(dptr_ofs + i*stride);
+    }
+    
+    /*! writes the i'th element */
+    __forceinline void store(size_t i, const Vec3fa& v)
+    {
+      assert(i<num);
+      Vec3fa::storeu(dptr_ofs + i*stride, v);
+    }
+    
+#else
+
+    /*! access to the ith element of the buffer */
+    __forceinline const Vec3fa operator [](size_t i) const
+    {
+      assert(i<num);
+      return Vec3fa(vfloat4::loadu((float*)(ptr_ofs + i*stride)));
+    }
+    
+    /*! writes the i'th element */
+    __forceinline void store(size_t i, const Vec3fa& v)
+    {
+      assert(i<num);
+      vfloat4::storeu((float*)(ptr_ofs + i*stride), (vfloat4)v);
+    }
+#endif
+  };
+}