Long live the slayer!

Initial commit of OpenGL Runtime to repository.
Based on SHA1 61823aaccc6510699a54b34a2fe3f7523dab3b4e
of qt3dstudio repository.

Task-number: QT3DS-3600
Change-Id: Iaeb80237399f0e5656a19ebec9d1ab3a681d8832
Reviewed-by: Pasi Keränen <pasi.keranen@qt.io>
 

9 files changed, 1732 insertions, 0 deletions

diff --git a/src/hdr/CUDABSDFMipmap.cpp b/src/hdr/CUDABSDFMipmap.cpp
new file mode 100644
index 0000000..94f6712
--- /dev/null
+++ b/src/hdr/CUDABSDFMipmap.cpp
@@ -0,0 +1,145 @@
+/****************************************************************************
+**
+** Copyright (C) 2017 The Qt Company Ltd.
+** Contact: https://www.qt.io/licensing/
+**
+** This file is part of Qt 3D Studio.
+**
+** $QT_BEGIN_LICENSE:GPL$
+** Commercial License Usage
+** Licensees holding valid commercial Qt licenses may use this file in
+** accordance with the commercial license agreement provided with the
+** Software or, alternatively, in accordance with the terms contained in
+** a written agreement between you and The Qt Company. For licensing terms
+** and conditions see https://www.qt.io/terms-conditions. For further
+** information use the contact form at https://www.qt.io/contact-us.
+**
+** GNU General Public License Usage
+** Alternatively, this file may be used under the terms of the GNU
+** General Public License version 3 or (at your option) any later version
+** approved by the KDE Free Qt Foundation. The licenses are as published by
+** the Free Software Foundation and appearing in the file LICENSE.GPL3
+** included in the packaging of this file. Please review the following
+** information to ensure the GNU General Public License requirements will
+** be met: https://www.gnu.org/licenses/gpl-3.0.html.
+**
+** $QT_END_LICENSE$
+**
+****************************************************************************/
+
+#ifdef PLATFORM_HAS_CUDA
+
+#include "CUDABSDFMipmap.h"
+#include "cuda.h"
+#include "cuda_runtime.h"
+#include "cuda_gl_interop.h"
+#include "render/backends/Qt3DSRenderBackend.h"
+#include "render/Qt3DSRenderTexture2D.h"
+#include "foundation/Qt3DSRefCounted.h"
+#include "nv_log.h"
+
+using namespace qt3ds;
+using namespace qt3ds::render;
+using namespace qt3ds::foundation;
+
+__host__ void jerror1(cudaError error);
+#ifdef _DEBUG
+#define CHECK_AND_HANDLE_CUDA_ERROR(func)                                                          \
+    func;                                                                                          \
+    {                                                                                              \
+        cudaError error = cudaGetLastError();                                                      \
+        if (error != cudaSuccess) {                                                                \
+            printf("%s\n", cudaGetErrorString(error));                                             \
+            jerror1(error);                                                                        \
+            QT3DS_ASSERT(false);                                                                      \
+        }                                                                                          \
+    }
+#else
+#define CHECK_AND_HANDLE_CUDA_ERROR(func) func;
+#endif
+
+CUDABSDFMipMap::CUDABSDFMipMap(NVRenderContext *inNVRenderContext, int inWidth, int inHeight,
+                               NVRenderTexture2D &inTexture2D,
+                               NVRenderTextureFormats::Enum inDestFormat, NVFoundationBase &inFnd)
+    : BSDFMipMap(inNVRenderContext, inWidth, inHeight, inTexture2D, inDestFormat, inFnd)
+    , m_TextureBinded(false)
+{
+
+    // CHECK_AND_HANDLE_CUDA_ERROR( cudaFree( 0 ); )
+    m_Pitches = (size_t *)QT3DS_ALLOC(m_Foundation.getAllocator(),
+                                   sizeof(size_t) * m_MaxMipMapLevel + 1, "BSDF MipMap pitches");
+    md_MipMapsData = (void **)QT3DS_ALLOC(m_Foundation.getAllocator(),
+                                       sizeof(void *) * m_MaxMipMapLevel + 1, "BSDF MipMap data");
+    // CHECK_AND_HANDLE_CUDA_ERROR();
+    size_t imagePitch;
+    int width = m_Width;
+    int height = m_Height;
+
+    for (int i = 0; i <= m_MaxMipMapLevel; ++i) {
+        imagePitch = m_SizeOfFormat * width;
+        // checkCudaErrors(cudaMalloc((void **)&cuda_dest_resource[mip], size_tex_data));
+        CHECK_AND_HANDLE_CUDA_ERROR(
+            cudaMallocPitch((void **)&md_MipMapsData[i], &m_Pitches[i], imagePitch, height);)
+        CHECK_AND_HANDLE_CUDA_ERROR(cudaMemset(md_MipMapsData[i], -1, m_Pitches[i] * height);)
+
+        width = width > 2 ? width >> 1 : 1;
+        height = height > 2 ? height >> 1 : 1;
+    }
+}
+
+CUDABSDFMipMap::~CUDABSDFMipMap()
+{
+    //	CHECK_AND_HANDLE_CUDA_ERROR( cudaDeviceReset(); )
+    CHECK_AND_HANDLE_CUDA_ERROR(cudaDeviceSynchronize();)
+    for (int i = 0; i <= m_MaxMipMapLevel; ++i) {
+        CHECK_AND_HANDLE_CUDA_ERROR(cudaFree(md_MipMapsData[i]);)
+    }
+    QT3DS_FREE(m_Foundation.getAllocator(), md_MipMapsData);
+    QT3DS_FREE(m_Foundation.getAllocator(), m_Pitches);
+}
+
+void CUDABSDFMipMap::BindTexture()
+{
+    if (!m_TextureBinded) {
+        m_TextureBinded = true;
+
+        int width = m_Width;
+        int height = m_Height;
+        for (int i = 0; i <= m_MaxMipMapLevel; ++i) {
+            // if you wwant to change some texture filter settings use m_Texture2D object
+            m_Texture2D.SetTextureData(NVDataRef<QT3DSU8>(), (QT3DSU8)i, width, height,
+                                       NVRenderTextureFormats::RGBA16F,
+                                       NVRenderTextureFormats::RGBA16F);
+
+            width = width > 2 ? width >> 1 : 1;
+            height = height > 2 ? height >> 1 : 1;
+        }
+        // CHECK_AND_HANDLE_CUDA_ERROR( cudaGraphicsGLRegisterImage( &m_CudaMipMapResource,
+        // (GLuint)m_TextureHandle, GL_TEXTURE_2D, cudaGraphicsRegisterFlagsWriteDiscard |
+        // cudaGraphicsRegisterFlagsTextureGather) )
+        CHECK_AND_HANDLE_CUDA_ERROR(cudaGraphicsGLRegisterImage(
+            &m_CudaMipMapResource, (GLuint)m_TextureHandle, GL_TEXTURE_2D,
+            cudaGraphicsRegisterFlagsWriteDiscard | cudaGraphicsRegisterFlagsTextureGather))
+    }
+}
+
+void CUDABSDFMipMap::TransferTexture() // after cuda generation
+{
+    cudaArray *texturePtr;
+    CHECK_AND_HANDLE_CUDA_ERROR(cudaGraphicsMapResources(1, &m_CudaMipMapResource, 0))
+    int width = m_Width;
+    int height = m_Height;
+    for (int i = 0; i <= m_MaxMipMapLevel; ++i) {
+        CHECK_AND_HANDLE_CUDA_ERROR(
+            cudaGraphicsSubResourceGetMappedArray(&texturePtr, m_CudaMipMapResource, 0, i))
+        CHECK_AND_HANDLE_CUDA_ERROR(cudaMemcpy2DToArray(texturePtr, 0, 0, md_MipMapsData[i],
+                                                        m_Pitches[i], width * m_SizeOfFormat,
+                                                        height, cudaMemcpyDeviceToDevice));
+
+        width = width > 2 ? width >> 1 : 1;
+        height = height > 2 ? height >> 1 : 1;
+    }
+    CHECK_AND_HANDLE_CUDA_ERROR(cudaGraphicsUnmapResources(1, &m_CudaMipMapResource, 0))
+}
+
+#endif
diff --git a/src/hdr/CUDABSDFMipmap.h b/src/hdr/CUDABSDFMipmap.h
new file mode 100644
index 0000000..33e0b91
--- /dev/null
+++ b/src/hdr/CUDABSDFMipmap.h
@@ -0,0 +1,77 @@
+/****************************************************************************
+**
+** Copyright (C) 2017 The Qt Company Ltd.
+** Contact: https://www.qt.io/licensing/
+**
+** This file is part of Qt 3D Studio.
+**
+** $QT_BEGIN_LICENSE:GPL$
+** Commercial License Usage
+** Licensees holding valid commercial Qt licenses may use this file in
+** accordance with the commercial license agreement provided with the
+** Software or, alternatively, in accordance with the terms contained in
+** a written agreement between you and The Qt Company. For licensing terms
+** and conditions see https://www.qt.io/terms-conditions. For further
+** information use the contact form at https://www.qt.io/contact-us.
+**
+** GNU General Public License Usage
+** Alternatively, this file may be used under the terms of the GNU
+** General Public License version 3 or (at your option) any later version
+** approved by the KDE Free Qt Foundation. The licenses are as published by
+** the Free Software Foundation and appearing in the file LICENSE.GPL3
+** included in the packaging of this file. Please review the following
+** information to ensure the GNU General Public License requirements will
+** be met: https://www.gnu.org/licenses/gpl-3.0.html.
+**
+** $QT_END_LICENSE$
+**
+****************************************************************************/
+
+#ifndef CUDABSDfMIPMAP_H
+#define CUDABSDFMIPMAP_H
+#include "foundation/Qt3DSVec3.h"
+#include "foundation/Qt3DSSimpleTypes.h"
+#include "foundation/Qt3DSPerfTimer.h"
+#include "foundation/Qt3DSAtomic.h"
+#include "render/Qt3DSRenderBaseTypes.h"
+#include "render/backends/Qt3DSRenderBackend.h"
+#include "MipmapBSDF.h"
+
+#include "Qt3DSRenderLoadedTexture.h"
+
+#include "Qt3DSTypes.h"
+struct cudaGraphicsResource;
+#ifdef _LINUXPLATFORM
+#define __declspec(dllexport)
+#define __cdecl
+
+#endif
+using namespace qt3ds::render;
+
+class CUDABSDFMipMap : public BSDFMipMap
+{
+public:
+    CUDABSDFMipMap(NVRenderContext *inNVRenderContext, int inWidth, int inHeight,
+                   NVRenderTexture2D &inTexture, NVRenderTextureFormats::Enum inDestFormat,
+                   qt3ds::NVFoundationBase &inFnd);
+    ~CUDABSDFMipMap();
+    void Build(void *inTextureData, int inTextureDataSize,
+               NVRenderBackend::NVRenderBackendTextureObject inTextureHandle,
+               NVRenderTextureFormats::Enum inFormat);
+    QT3DS_IMPLEMENT_REF_COUNT_ADDREF_RELEASE(m_Foundation);
+
+private:
+    void CreateBsdfMipMaps(qt3ds::render::SLoadedTexture &inLoadedImage, void **result, int width,
+                           int height); //, qt3ds::foundation::IPerfTimer& inPerfTimer);
+
+    void BindTexture();
+    void TransferTexture();
+
+    cudaGraphicsResource *m_CudaMipMapResource;
+    void **md_MipMapsData;
+    size_t *m_Pitches;
+    NVRenderBackend::NVRenderBackendTextureObject m_TextureHandle;
+    bool m_TextureBinded;
+};
+
+#endif
diff --git a/src/hdr/GLComputeMipMap.h b/src/hdr/GLComputeMipMap.h
new file mode 100644
index 0000000..a1a7ce5
--- /dev/null
+++ b/src/hdr/GLComputeMipMap.h
@@ -0,0 +1,74 @@
+/****************************************************************************
+**
+** Copyright (C) 2008-2012 NVIDIA Corporation.
+** Copyright (C) 2017 The Qt Company Ltd.
+** Contact: https://www.qt.io/licensing/
+**
+** This file is part of Qt 3D Studio.
+**
+** $QT_BEGIN_LICENSE:GPL$
+** Commercial License Usage
+** Licensees holding valid commercial Qt licenses may use this file in
+** accordance with the commercial license agreement provided with the
+** Software or, alternatively, in accordance with the terms contained in
+** a written agreement between you and The Qt Company. For licensing terms
+** and conditions see https://www.qt.io/terms-conditions. For further
+** information use the contact form at https://www.qt.io/contact-us.
+**
+** GNU General Public License Usage
+** Alternatively, this file may be used under the terms of the GNU
+** General Public License version 3 or (at your option) any later version
+** approved by the KDE Free Qt Foundation. The licenses are as published by
+** the Free Software Foundation and appearing in the file LICENSE.GPL3
+** included in the packaging of this file. Please review the following
+** information to ensure the GNU General Public License requirements will
+** be met: https://www.gnu.org/licenses/gpl-3.0.html.
+**
+** $QT_END_LICENSE$
+**
+****************************************************************************/
+
+#ifndef GLCOMPUTE_BSDF_MIMAP_H
+#define GLCOMPUTE_BSDF_MIMAP_H
+
+#include "render/Qt3DSRenderBaseTypes.h"
+#include "render/backends/Qt3DSRenderBackend.h"
+#include "MipmapBSDF.h"
+
+#include "Qt3DSRenderLoadedTexture.h"
+#include "Qt3DSTypes.h"
+
+using namespace qt3ds::render;
+
+class qt3ds::render::NVRenderContext;
+class qt3ds::render::NVRenderShaderProgram;
+class qt3ds::render::NVRenderTexture2D;
+
+class GLComputeMipMap : public BSDFMipMap
+{
+public:
+    GLComputeMipMap(NVRenderContext *inNVRenderContext, int inWidth, int inHeight,
+                    NVRenderTexture2D &inTexture, NVRenderTextureFormats::Enum inDestFormat,
+                    qt3ds::NVFoundationBase &inFnd);
+    ~GLComputeMipMap();
+    void Build(void *inTextureData, int inTextureDataSize,
+               NVRenderBackend::NVRenderBackendTextureObject inTextureHandle,
+               NVRenderTextureFormats::Enum inFormat);
+    QT3DS_IMPLEMENT_REF_COUNT_ADDREF_RELEASE(m_Foundation);
+
+private:
+    void CreateLevel0Tex(void *inTextureData, int inTextureDataSize,
+                         NVRenderTextureFormats::Enum inFormat);
+
+    NVScopedRefCounted<NVRenderShaderProgram> m_BSDFProgram;
+    NVScopedRefCounted<NVRenderShaderProgram> m_UploadProgram_RGBA8;
+    NVScopedRefCounted<NVRenderShaderProgram> m_UploadProgram_RGB8;
+    NVScopedRefCounted<NVRenderTexture2D> m_Level0Tex;
+    bool m_TextureCreated;
+
+    void createComputeProgram(NVRenderContext *context);
+    NVRenderShaderProgram *getOrCreateUploadComputeProgram(NVRenderContext *context,
+                                                           NVRenderTextureFormats::Enum inFormat);
+};
+
+#endif
diff --git a/src/hdr/GLComputeMipmap.cpp b/src/hdr/GLComputeMipmap.cpp
new file mode 100644
index 0000000..36b550d
--- /dev/null
+++ b/src/hdr/GLComputeMipmap.cpp
@@ -0,0 +1,394 @@
+/****************************************************************************
+**
+** Copyright (C) 2008-2012 NVIDIA Corporation.
+** Copyright (C) 2017 The Qt Company Ltd.
+** Contact: https://www.qt.io/licensing/
+**
+** This file is part of Qt 3D Studio.
+**
+** $QT_BEGIN_LICENSE:GPL$
+** Commercial License Usage
+** Licensees holding valid commercial Qt licenses may use this file in
+** accordance with the commercial license agreement provided with the
+** Software or, alternatively, in accordance with the terms contained in
+** a written agreement between you and The Qt Company. For licensing terms
+** and conditions see https://www.qt.io/terms-conditions. For further
+** information use the contact form at https://www.qt.io/contact-us.
+**
+** GNU General Public License Usage
+** Alternatively, this file may be used under the terms of the GNU
+** General Public License version 3 or (at your option) any later version
+** approved by the KDE Free Qt Foundation. The licenses are as published by
+** the Free Software Foundation and appearing in the file LICENSE.GPL3
+** included in the packaging of this file. Please review the following
+** information to ensure the GNU General Public License requirements will
+** be met: https://www.gnu.org/licenses/gpl-3.0.html.
+**
+** $QT_END_LICENSE$
+**
+****************************************************************************/
+
+#include "GLComputeMipMap.h"
+#include "render/Qt3DSRenderTexture2D.h"
+#include "render/Qt3DSRenderShaderProgram.h"
+#include "render/Qt3DSRenderContext.h"
+#include "nv_log.h"
+#include <string>
+
+using namespace qt3ds;
+using namespace qt3ds::render;
+using namespace qt3ds::foundation;
+
+static const char *computeUploadShader(std::string &prog, NVRenderTextureFormats::Enum inFormat,
+                                       bool binESContext)
+{
+    if (binESContext) {
+        prog += "#version 310 es\n"
+                "#extension GL_ARB_compute_shader : enable\n"
+                "precision highp float;\n"
+                "precision highp int;\n"
+                "precision mediump image2D;\n";
+    } else {
+        prog += "#version 430\n"
+                "#extension GL_ARB_compute_shader : enable\n";
+    }
+
+    if (inFormat == NVRenderTextureFormats::RGBA8) {
+        prog += "// Set workgroup layout;\n"
+                "layout (local_size_x = 16, local_size_y = 16) in;\n\n"
+                "layout (rgba8, binding = 1) uniform image2D inputImage;\n\n"
+                "layout (rgba16f, binding = 2) uniform image2D outputImage;\n\n"
+                "void main()\n"
+                "{\n"
+                "  if ( gl_GlobalInvocationID.x >= gl_NumWorkGroups.x || gl_GlobalInvocationID.y "
+                ">= gl_NumWorkGroups.y )\n"
+                "    return;\n"
+                "  vec4 value = imageLoad(inputImage, ivec2(gl_GlobalInvocationID.xy));\n"
+                "  imageStore( outputImage, ivec2(gl_GlobalInvocationID.xy), value );\n"
+                "}\n";
+    } else {
+        prog += "float convertToFloat( in uint inValue )\n"
+                "{\n"
+                "  uint v = inValue & uint(0xFF);\n"
+                "  float f = float(v)/256.0;\n"
+                "  return f;\n"
+                "}\n";
+
+        prog += "int getMod( in int inValue, in int mod )\n"
+                "{\n"
+                "  int v = mod * (inValue/mod);\n"
+                "  return inValue - v;\n"
+                "}\n";
+
+        prog += "vec4 getRGBValue( in int byteNo, vec4 inVal, vec4 inVal1 )\n"
+                "{\n"
+                "  vec4 result= vec4(0.0);\n"
+                "  if( byteNo == 0) {\n"
+                "    result.r = inVal.r;\n"
+                "    result.g = inVal.g;\n"
+                "    result.b = inVal.b;\n"
+                "  }\n"
+                "  else if( byteNo == 1) {\n"
+                "    result.r = inVal.g;\n"
+                "    result.g = inVal.b;\n"
+                "    result.b = inVal.a;\n"
+                "  }\n"
+                "  else if( byteNo == 2) {\n"
+                "    result.r = inVal.b;\n"
+                "    result.g = inVal.a;\n"
+                "    result.b = inVal1.r;\n"
+                "  }\n"
+                "  else if( byteNo == 3) {\n"
+                "    result.r = inVal.a;\n"
+                "    result.g = inVal1.r;\n"
+                "    result.b = inVal1.g;\n"
+                "  }\n"
+                "  return result;\n"
+                "}\n";
+
+        prog += "// Set workgroup layout;\n"
+                "layout (local_size_x = 16, local_size_y = 16) in;\n\n"
+                "layout (rgba8, binding = 1) uniform image2D inputImage;\n\n"
+                "layout (rgba16f, binding = 2) uniform image2D outputImage;\n\n"
+                "void main()\n"
+                "{\n"
+                "  vec4 result = vec4(0.0);\n"
+                "  if ( gl_GlobalInvocationID.x >= gl_NumWorkGroups.x || gl_GlobalInvocationID.y "
+                ">= gl_NumWorkGroups.y )\n"
+                "    return;\n"
+                "  int xpos = (int(gl_GlobalInvocationID.x)*3)/4;\n"
+                "  int xmod = getMod(int(gl_GlobalInvocationID.x)*3, 4);\n"
+                "  ivec2 readPos = ivec2(xpos, gl_GlobalInvocationID.y);\n"
+                "  vec4 value = imageLoad(inputImage, readPos);\n"
+                "  vec4 value1 = imageLoad(inputImage, ivec2(readPos.x + 1, readPos.y));\n"
+                "  result = getRGBValue( xmod, value, value1);\n"
+                "  imageStore( outputImage, ivec2(gl_GlobalInvocationID.xy), result );\n"
+                "}\n";
+    }
+    return prog.c_str();
+}
+
+static const char *computeWorkShader(std::string &prog, bool binESContext)
+{
+    if (binESContext) {
+        prog += "#version 310 es\n"
+                "#extension GL_ARB_compute_shader : enable\n"
+                "precision highp float;\n"
+                "precision highp int;\n"
+                "precision mediump image2D;\n";
+    } else {
+        prog += "#version 430\n"
+                "#extension GL_ARB_compute_shader : enable\n";
+    }
+
+    prog += "int wrapMod( in int a, in int base )\n"
+            "{\n"
+            "  return ( a >= 0 ) ? a % base : -(a % base) + base;\n"
+            "}\n";
+
+    prog += "void getWrappedCoords( inout int sX, inout int sY, in int width, in int height )\n"
+            "{\n"
+            "  if (sY < 0) { sX -= width >> 1; sY = -sY; }\n"
+            "  if (sY >= height) { sX += width >> 1; sY = height - sY; }\n"
+            "  sX = wrapMod( sX, width );\n"
+            "}\n";
+
+    prog += "// Set workgroup layout;\n"
+            "layout (local_size_x = 16, local_size_y = 16) in;\n\n"
+            "layout (rgba16f, binding = 1) uniform image2D inputImage;\n\n"
+            "layout (rgba16f, binding = 2) uniform image2D outputImage;\n\n"
+            "void main()\n"
+            "{\n"
+            "  int prevWidth = int(gl_NumWorkGroups.x) << 1;\n"
+            "  int prevHeight = int(gl_NumWorkGroups.y) << 1;\n"
+            "  if ( gl_GlobalInvocationID.x >= gl_NumWorkGroups.x || gl_GlobalInvocationID.y >= "
+            "gl_NumWorkGroups.y )\n"
+            "    return;\n"
+            "  vec4 accumVal = vec4(0.0);\n"
+            "  for ( int sy = -2; sy <= 2; ++sy )\n"
+            "  {\n"
+            "    for ( int sx = -2; sx <= 2; ++sx )\n"
+            "    {\n"
+            "      int sampleX = sx + (int(gl_GlobalInvocationID.x) << 1);\n"
+            "      int sampleY = sy + (int(gl_GlobalInvocationID.y) << 1);\n"
+            "      getWrappedCoords(sampleX, sampleY, prevWidth, prevHeight);\n"
+            "	   if ((sampleY * prevWidth + sampleX) < 0 )\n"
+            "        sampleY = prevHeight + sampleY;\n"
+            "      ivec2 pos = ivec2(sampleX, sampleY);\n"
+            "      vec4 value = imageLoad(inputImage, pos);\n"
+            "      float filterPdf = 1.0 / ( 1.0 + float(sx*sx + sy*sy)*2.0 );\n"
+            "      filterPdf /= 4.71238898;\n"
+            "      accumVal[0] += filterPdf * value.r;\n"
+            "	   accumVal[1] += filterPdf * value.g;\n"
+            "	   accumVal[2] += filterPdf * value.b;\n"
+            "	   accumVal[3] += filterPdf * value.a;\n"
+            "    }\n"
+            "  }\n"
+            "  imageStore( outputImage, ivec2(gl_GlobalInvocationID.xy), accumVal );\n"
+            "}\n";
+
+    return prog.c_str();
+}
+
+GLComputeMipMap::GLComputeMipMap(NVRenderContext *inNVRenderContext, int inWidth, int inHeight,
+                                 NVRenderTexture2D &inTexture2D,
+                                 NVRenderTextureFormats::Enum inDestFormat, NVFoundationBase &inFnd)
+    : BSDFMipMap(inNVRenderContext, inWidth, inHeight, inTexture2D, inDestFormat, inFnd)
+    , m_BSDFProgram(NULL)
+    , m_UploadProgram_RGBA8(NULL)
+    , m_UploadProgram_RGB8(NULL)
+    , m_Level0Tex(NULL)
+    , m_TextureCreated(false)
+{
+}
+
+GLComputeMipMap::~GLComputeMipMap()
+{
+    m_UploadProgram_RGB8 = NULL;
+    m_UploadProgram_RGBA8 = NULL;
+    m_BSDFProgram = NULL;
+    m_Level0Tex = NULL;
+}
+
+inline NVConstDataRef<QT3DSI8> toRef(const char *data)
+{
+    size_t len = strlen(data) + 1;
+    return NVConstDataRef<QT3DSI8>((const QT3DSI8 *)data, (QT3DSU32)len);
+}
+
+static bool isGLESContext(NVRenderContext *context)
+{
+    NVRenderContextType ctxType = context->GetRenderContextType();
+
+    // Need minimum of GL3 or GLES3
+    if (ctxType == NVRenderContextValues::GLES2 || ctxType == NVRenderContextValues::GLES3
+        || ctxType == NVRenderContextValues::GLES31) {
+        return true;
+    }
+
+    return false;
+}
+
+#define WORKGROUP_SIZE 16
+
+void GLComputeMipMap::createComputeProgram(NVRenderContext *context)
+{
+    std::string computeProg;
+
+    if (!m_BSDFProgram) {
+        m_BSDFProgram = context
+                            ->CompileComputeSource(
+                                "Compute BSDF mipmap shader",
+                                toRef(computeWorkShader(computeProg, isGLESContext(context))))
+                            .mShader;
+    }
+}
+
+NVRenderShaderProgram *
+GLComputeMipMap::getOrCreateUploadComputeProgram(NVRenderContext *context,
+                                                 NVRenderTextureFormats::Enum inFormat)
+{
+    std::string computeProg;
+
+    if (inFormat == NVRenderTextureFormats::RGB8) {
+        if (!m_UploadProgram_RGB8) {
+            m_UploadProgram_RGB8 =
+                context
+                    ->CompileComputeSource(
+                        "Compute BSDF mipmap level 0 RGB8 shader",
+                        toRef(computeUploadShader(computeProg, inFormat, isGLESContext(context))))
+                    .mShader;
+        }
+
+        return m_UploadProgram_RGB8;
+    } else {
+        if (!m_UploadProgram_RGBA8) {
+            m_UploadProgram_RGBA8 =
+                context
+                    ->CompileComputeSource(
+                        "Compute BSDF mipmap level 0 RGBA8 shader",
+                        toRef(computeUploadShader(computeProg, inFormat, isGLESContext(context))))
+                    .mShader;
+        }
+
+        return m_UploadProgram_RGBA8;
+    }
+}
+
+void GLComputeMipMap::CreateLevel0Tex(void *inTextureData, int inTextureDataSize,
+                                      NVRenderTextureFormats::Enum inFormat)
+{
+    NVRenderTextureFormats::Enum theFormat = inFormat;
+    int theWidth = m_Width;
+
+    // Since we cannot use RGB format in GL compute
+    // we treat it as a RGBA component format
+    if (inFormat == NVRenderTextureFormats::RGB8) {
+        // This works only with 4 byte aligned data
+        QT3DS_ASSERT(m_Width % 4 == 0);
+        theFormat = NVRenderTextureFormats::RGBA8;
+        theWidth = (m_Width * 3) / 4;
+    }
+
+    if (m_Level0Tex == NULL) {
+        m_Level0Tex = m_NVRenderContext->CreateTexture2D();
+        m_Level0Tex->SetTextureStorage(1, theWidth, m_Height, theFormat, theFormat,
+                                       NVDataRef<QT3DSU8>((QT3DSU8 *)inTextureData, inTextureDataSize));
+    } else {
+        m_Level0Tex->SetTextureSubData(NVDataRef<QT3DSU8>((QT3DSU8 *)inTextureData, inTextureDataSize), 0,
+                                       0, 0, theWidth, m_Height, theFormat);
+    }
+}
+
+void GLComputeMipMap::Build(void *inTextureData, int inTextureDataSize,
+                            NVRenderBackend::NVRenderBackendTextureObject,
+                            NVRenderTextureFormats::Enum inFormat)
+{
+    bool needMipUpload = (inFormat != m_DestinationFormat);
+    // re-upload data
+    if (!m_TextureCreated) {
+        m_Texture2D.SetTextureStorage(
+            m_MaxMipMapLevel + 1, m_Width, m_Height, m_DestinationFormat, inFormat, (needMipUpload)
+                ? NVDataRef<QT3DSU8>()
+                : NVDataRef<QT3DSU8>((QT3DSU8 *)inTextureData, inTextureDataSize));
+        m_Texture2D.addRef();
+        // create a compute shader (if not aloread done) which computes the BSDF mipmaps for this
+        // texture
+        createComputeProgram(m_NVRenderContext);
+
+        if (!m_BSDFProgram) {
+            QT3DS_ASSERT(false);
+            return;
+        }
+
+        m_TextureCreated = true;
+    } else if (!needMipUpload) {
+        m_Texture2D.SetTextureSubData(NVDataRef<QT3DSU8>((QT3DSU8 *)inTextureData, inTextureDataSize), 0,
+                                      0, 0, m_Width, m_Height, inFormat);
+    }
+
+    if (needMipUpload) {
+        CreateLevel0Tex(inTextureData, inTextureDataSize, inFormat);
+    }
+
+    NVScopedRefCounted<NVRenderImage2D> theInputImage;
+    NVScopedRefCounted<NVRenderImage2D> theOutputImage;
+    theInputImage =
+        m_NVRenderContext->CreateImage2D(&m_Texture2D, NVRenderImageAccessType::ReadWrite);
+    theOutputImage =
+        m_NVRenderContext->CreateImage2D(&m_Texture2D, NVRenderImageAccessType::ReadWrite);
+
+    if (needMipUpload && m_Level0Tex) {
+        NVRenderShaderProgram *uploadProg =
+            getOrCreateUploadComputeProgram(m_NVRenderContext, inFormat);
+        if (!uploadProg)
+            return;
+
+        m_NVRenderContext->SetActiveShader(uploadProg);
+
+        NVScopedRefCounted<NVRenderImage2D> theInputImage0;
+        theInputImage0 =
+            m_NVRenderContext->CreateImage2D(m_Level0Tex, NVRenderImageAccessType::ReadWrite);
+
+        theInputImage0->SetTextureLevel(0);
+        NVRenderCachedShaderProperty<NVRenderImage2D *> theCachedinputImage0("inputImage",
+                                                                             *uploadProg);
+        theCachedinputImage0.Set(theInputImage0);
+
+        theOutputImage->SetTextureLevel(0);
+        NVRenderCachedShaderProperty<NVRenderImage2D *> theCachedOutputImage("outputImage",
+                                                                             *uploadProg);
+        theCachedOutputImage.Set(theOutputImage);
+
+        m_NVRenderContext->DispatchCompute(uploadProg, m_Width, m_Height, 1);
+
+        // sync
+        NVRenderBufferBarrierFlags flags(NVRenderBufferBarrierValues::ShaderImageAccess);
+        m_NVRenderContext->SetMemoryBarrier(flags);
+    }
+
+    int width = m_Width >> 1;
+    int height = m_Height >> 1;
+
+    m_NVRenderContext->SetActiveShader(m_BSDFProgram);
+
+    for (int i = 1; i <= m_MaxMipMapLevel; ++i) {
+        theOutputImage->SetTextureLevel(i);
+        NVRenderCachedShaderProperty<NVRenderImage2D *> theCachedOutputImage("outputImage",
+                                                                             *m_BSDFProgram);
+        theCachedOutputImage.Set(theOutputImage);
+        theInputImage->SetTextureLevel(i - 1);
+        NVRenderCachedShaderProperty<NVRenderImage2D *> theCachedinputImage("inputImage",
+                                                                            *m_BSDFProgram);
+        theCachedinputImage.Set(theInputImage);
+
+        m_NVRenderContext->DispatchCompute(m_BSDFProgram, width, height, 1);
+
+        width = width > 2 ? width >> 1 : 1;
+        height = height > 2 ? height >> 1 : 1;
+
+        // sync
+        NVRenderBufferBarrierFlags flags(NVRenderBufferBarrierValues::ShaderImageAccess);
+        m_NVRenderContext->SetMemoryBarrier(flags);
+    }
+}
diff --git a/src/hdr/HDR.cpp b/src/hdr/HDR.cpp
new file mode 100644
index 0000000..9abe2b3
--- /dev/null
+++ b/src/hdr/HDR.cpp
@@ -0,0 +1,30 @@
+/****************************************************************************
+**
+** Copyright (C) 2017 The Qt Company Ltd.
+** Contact: https://www.qt.io/licensing/
+**
+** This file is part of Qt 3D Studio.
+**
+** $QT_BEGIN_LICENSE:GPL$
+** Commercial License Usage
+** Licensees holding valid commercial Qt licenses may use this file in
+** accordance with the commercial license agreement provided with the
+** Software or, alternatively, in accordance with the terms contained in
+** a written agreement between you and The Qt Company. For licensing terms
+** and conditions see https://www.qt.io/terms-conditions. For further
+** information use the contact form at https://www.qt.io/contact-us.
+**
+** GNU General Public License Usage
+** Alternatively, this file may be used under the terms of the GNU
+** General Public License version 3 or (at your option) any later version
+** approved by the KDE Free Qt Foundation. The licenses are as published by
+** the Free Software Foundation and appearing in the file LICENSE.GPL3
+** included in the packaging of this file. Please review the following
+** information to ensure the GNU General Public License requirements will
+** be met: https://www.gnu.org/licenses/gpl-3.0.html.
+**
+** $QT_END_LICENSE$
+**
+****************************************************************************/
+
+#include "HDR.h"
diff --git a/src/hdr/HDR.h b/src/hdr/HDR.h
new file mode 100644
index 0000000..fbc4ea4
--- /dev/null
+++ b/src/hdr/HDR.h
@@ -0,0 +1,239 @@
+/****************************************************************************
+**
+** Copyright (C) 2008-2012 NVIDIA Corporation.
+** Copyright (C) 2017 The Qt Company Ltd.
+** Contact: https://www.qt.io/licensing/
+**
+** This file is part of Qt 3D Studio.
+**
+** $QT_BEGIN_LICENSE:GPL$
+** Commercial License Usage
+** Licensees holding valid commercial Qt licenses may use this file in
+** accordance with the commercial license agreement provided with the
+** Software or, alternatively, in accordance with the terms contained in
+** a written agreement between you and The Qt Company. For licensing terms
+** and conditions see https://www.qt.io/terms-conditions. For further
+** information use the contact form at https://www.qt.io/contact-us.
+**
+** GNU General Public License Usage
+** Alternatively, this file may be used under the terms of the GNU
+** General Public License version 3 or (at your option) any later version
+** approved by the KDE Free Qt Foundation. The licenses are as published by
+** the Free Software Foundation and appearing in the file LICENSE.GPL3
+** included in the packaging of this file. Please review the following
+** information to ensure the GNU General Public License requirements will
+** be met: https://www.gnu.org/licenses/gpl-3.0.html.
+**
+** $QT_END_LICENSE$
+**
+****************************************************************************/
+
+#pragma once
+#ifndef HDR_H
+#define HDR_H
+
+#include "foundation/Qt3DSVec3.h"
+//#include "HDR.h"
+
+namespace qt3ds {
+class QT3DSVec3;
+
+namespace HDR {
+
+    template <int N>
+    class HDRConfiguration;
+
+    template <int N>
+    class Histogram
+    {
+    public:
+        /**
+         * @brief build the histogram based on 2^10 binning.
+         *
+         * @param[in]	inImage			Pointer to image
+         * @param[in] 	inWidth			Width of image
+         * @param[in] 	inHeight		Height of image
+         * @param[out]	outHistogram
+         *
+         * @return No return
+         */
+        static void Build(QT3DSVec3 *inImage, int inWidth, int inHeight, QT3DSVec3 *outHistogram)
+        {
+            long noPixels = inWidth * inHeight;
+
+            for (int i = 0; i < noPixels; ++i) {
+                outHistogram[(int)inImage[i].x].x++;
+                outHistogram[(int)inImage[i].y].y++;
+                outHistogram[(int)inImage[i].z].z++;
+            }
+        }
+    };
+
+    template <int N>
+    class HDR
+    {
+
+    public:
+        HDR(HDRConfiguration<N> *inConfiguration) { mHDRConfiguration = inConfiguration; }
+
+        void Build(QT3DSVec3 **inImages, int inNoImages, int inWidth, int inHeight, float *inExposures,
+                   QT3DSVec3 *outRadiance)
+        {
+            for (int x = 0; x < inWidth; ++x) {
+                for (int y = 0; y < inHeight; ++y) {
+                    QT3DSVec3 divisor(0.0f);
+                    QT3DSVec3 dividend(0.0f);
+                    for (int i = 0; i < inNoImages; ++i) {
+                        QT3DSVec3 pixel = inImages[i][x + y * inWidth];
+
+                        dividend += (LUT(mHDRConfiguration->weights, pixel) * inExposures[i])
+                                        .multiply(LUT(mHDRConfiguration->CRF, pixel));
+                        divisor += LUT(mHDRConfiguration->weights, pixel) * inExposures[i]
+                            * inExposures[i];
+                    }
+                    divisor.x = 1.0f / divisor.x;
+                    divisor.y = 1.0f / divisor.y;
+                    divisor.z = 1.0f / divisor.z;
+
+                    outRadiance[x + y * inWidth] = dividend.multiply(divisor);
+                }
+            }
+        }
+
+    private:
+        /**
+         * @brief Return the value based off from array as a LUT
+         *
+         * @param[in]	inLUT			The LUT table of interest
+         * @param[in]	inValue			The value that you have
+         * @param[out]	QT3DSVec3			The corresponding value of array based on input
+         * tuple
+         *
+         * @return No return
+         */
+        QT3DSVec3 LUT(float *inLUT, QT3DSVec3 inValue)
+        {
+            return QT3DSVec3(inLUT[(int)inValue.x], inLUT[(int)inValue.y], inLUT[(int)inValue.z]);
+        }
+
+        QT3DSVec3 LUT(QT3DSVec3 *inLUT, QT3DSVec3 inValue)
+        {
+            return QT3DSVec3(inLUT[(int)inValue.x].x, inLUT[(int)inValue.y].y,
+                          inLUT[(int)inValue.z].z);
+        }
+
+        HDRConfiguration<N> *mHDRConfiguration;
+    };
+
+    template <int N>
+    class HDRConfiguration
+    {
+    public:
+        HDRConfiguration()
+        {
+            threshold = 0.1f;
+            maxIterations = 30;
+            GenerateRobertsonWeighting();
+        }
+
+        void SetCRF(QT3DSVec3 *inCRF) { memcpy(CRF, inCRF, sizeof(QT3DSVec3) * N); }
+
+        /**
+         * @brief build the camera response function
+         *
+         * @param[in]	inImages		Pointer to images
+         * @param[in] 	inNoImages		Number of images
+         * @param[in] 	inWidth			Width of image
+         * @param[in] 	inHeight		Height of image
+         * @param[out]	outHistogram
+         *
+         * @return No return
+         */
+        void BuildCRF(QT3DSVec3 **inImages, int inNoImages, int inWidth, int inHeight,
+                      float *inExposures)
+        {
+            QT3DSVec3 histogram[N];
+            QT3DSVec3 newCRF[N];
+
+            HDR<N> hdr(this);
+
+            memset(histogram, 0, sizeof(QT3DSVec3) * N);
+
+            for (int i = 0; i < inNoImages; ++i) {
+                Histogram<N>::Build(inImages[i], inWidth, inHeight, histogram);
+            }
+
+            for (int i = 0; i < N; ++i) {
+                histogram[i].x = histogram[i].x > 0 ? 1 / histogram[i].x : 0;
+                histogram[i].y = histogram[i].y > 0 ? 1 / histogram[i].y : 0;
+                histogram[i].z = histogram[i].z > 0 ? 1 / histogram[i].z : 0;
+            }
+
+            QT3DSVec3 *radiance = new QT3DSVec3[inWidth * inHeight * sizeof(QT3DSVec3)];
+
+            // iteration 0, linearize CRF
+            for (int i = 0; i < N; ++i) {
+                CRF[i] = QT3DSVec3((float)i) * 2.0f / N;
+            }
+
+            for (int iteration = 0; iteration < maxIterations; ++iteration) {
+                hdr.Build(inImages, inNoImages, inWidth, inHeight, inExposures, radiance);
+
+                memset(newCRF, 0, sizeof(QT3DSVec3) * N);
+
+                for (int i = 0; i < inNoImages; ++i) {
+                    for (int x = 0; x < inWidth; ++x) {
+                        for (int y = 0; y < inHeight; ++y) {
+                            long offset = x + y * inWidth;
+                            QT3DSVec3 pixel = inImages[i][offset];
+                            newCRF[(int)pixel.x].x += (inExposures[i] * radiance[offset].x);
+                            newCRF[(int)pixel.y].y += (inExposures[i] * radiance[offset].y);
+                            newCRF[(int)pixel.z].z += (inExposures[i] * radiance[offset].z);
+                        }
+                    }
+                }
+
+                float difference = 0;
+                for (int i = 0; i < N; ++i) {
+                    newCRF[i] = newCRF[i].multiply(histogram[i]);
+                }
+
+                QT3DSVec3 middle = newCRF[N / 2];
+                for (int i = 0; i < N; ++i) {
+                    newCRF[i].x = newCRF[i].x / middle.x;
+                    newCRF[i].y = newCRF[i].y / middle.y;
+                    newCRF[i].z = newCRF[i].z / middle.z;
+                    difference += (CRF[i] - newCRF[i]).magnitude();
+                }
+                for (int i = 0; i < N; ++i) {
+                    CRF[i] = newCRF[i];
+                }
+                if (difference < threshold) {
+                    break;
+                }
+            }
+            delete[] radiance;
+        }
+
+        float weights[N];
+        QT3DSVec3 CRF[N];
+
+    private:
+        void GenerateRobertsonWeighting()
+        {
+            // Dynamic Range Improvement Through Multiple Exposures (5)
+            // gaussian random weighting
+            float divisor = (N - 1) * (N - 1) / 4.0;
+            for (int i = 0; i < N; ++i) {
+                float dividend = (i - (N - 1) / 2.0f);
+                dividend *= dividend;
+                weights[i] = exp(-4.0f * dividend / divisor);
+            }
+        }
+
+        int maxIterations;
+        float threshold;
+    };
+}
+}
+#endif
diff --git a/src/hdr/MipmapBSDF.cpp b/src/hdr/MipmapBSDF.cpp
new file mode 100644
index 0000000..fd57015
--- /dev/null
+++ b/src/hdr/MipmapBSDF.cpp
@@ -0,0 +1,265 @@
+/****************************************************************************
+**
+** Copyright (C) 2017 The Qt Company Ltd.
+** Contact: https://www.qt.io/licensing/
+**
+** This file is part of Qt 3D Studio.
+**
+** $QT_BEGIN_LICENSE:GPL$
+** Commercial License Usage
+** Licensees holding valid commercial Qt licenses may use this file in
+** accordance with the commercial license agreement provided with the
+** Software or, alternatively, in accordance with the terms contained in
+** a written agreement between you and The Qt Company. For licensing terms
+** and conditions see https://www.qt.io/terms-conditions. For further
+** information use the contact form at https://www.qt.io/contact-us.
+**
+** GNU General Public License Usage
+** Alternatively, this file may be used under the terms of the GNU
+** General Public License version 3 or (at your option) any later version
+** approved by the KDE Free Qt Foundation. The licenses are as published by
+** the Free Software Foundation and appearing in the file LICENSE.GPL3
+** included in the packaging of this file. Please review the following
+** information to ensure the GNU General Public License requirements will
+** be met: https://www.gnu.org/licenses/gpl-3.0.html.
+**
+** $QT_END_LICENSE$
+**
+****************************************************************************/
+
+#include "MipmapBSDF.h"
+#include "GLComputeMipMap.h"
+
+#ifdef PLATFORM_HAS_CUDA
+#include "cuda.h"
+#include "cuda_runtime.h"
+#include "CUDABSDFMipmap.h"
+#endif
+
+#include "render/Qt3DSRenderContext.h"
+#include "render/Qt3DSRenderTexture2D.h"
+#include "foundation/Qt3DSRefCounted.h"
+#include "nv_log.h"
+
+using namespace qt3ds;
+using namespace qt3ds::render;
+using namespace qt3ds::foundation;
+
+BSDFMipMap::BSDFMipMap(NVRenderContext *inNVRenderContext, int inWidth, int inHeight,
+                       NVRenderTexture2D &inTexture2D, NVRenderTextureFormats::Enum inDestFormat,
+                       NVFoundationBase &inFnd)
+    : m_Foundation(inFnd)
+    , m_Texture2D(inTexture2D)
+    , m_Width(inWidth)
+    , m_Height(inHeight)
+    , m_DestinationFormat(inDestFormat)
+    , m_NVRenderContext(inNVRenderContext)
+{
+    // Calculate mip level
+    int maxDim = inWidth >= inHeight ? inWidth : inHeight;
+
+    m_MaxMipMapLevel = static_cast<int>(logf((float)maxDim) / logf(2.0f));
+    // no concept of sizeOfFormat just does'nt make sense
+    m_SizeOfFormat = NVRenderTextureFormats::getSizeofFormat(m_DestinationFormat);
+    m_NoOfComponent = NVRenderTextureFormats::getNumberOfComponent(m_DestinationFormat);
+}
+
+BSDFMipMap *BSDFMipMap::Create(NVRenderContext *inNVRenderContext, int inWidth, int inHeight,
+                               NVRenderTexture2D &inTexture2D,
+                               NVRenderTextureFormats::Enum inDestFormat,
+                               qt3ds::NVFoundationBase &inFnd)
+{
+    BSDFMipMap *theBSDFMipMap = NULL;
+#ifdef PLATFORM_HAS_CUDA
+    int deviceCount;
+    cudaError_t e = cudaGetDeviceCount(&deviceCount);
+#endif
+
+    if (inNVRenderContext->IsComputeSupported()) {
+        theBSDFMipMap = QT3DS_NEW(inFnd.getAllocator(), GLComputeMipMap)(
+            inNVRenderContext, inWidth, inHeight, inTexture2D, inDestFormat, inFnd);
+    } else
+#ifdef PLATFORM_HAS_CUDA
+        if (e == cudaSuccess && deviceCount > 0) {
+        theBSDFMipMap = QT3DS_NEW(inFnd.getAllocator(), CUDABSDFMipMap)(
+            inNVRenderContext, inWidth, inHeight, inTexture2D, inDestFormat, inFnd);
+    } else
+#endif
+        if (!theBSDFMipMap) {
+        theBSDFMipMap = QT3DS_NEW(inFnd.getAllocator(), BasicBSDFMipMap)(
+            inNVRenderContext, inWidth, inHeight, inTexture2D, inDestFormat, inFnd);
+    }
+
+    return theBSDFMipMap;
+}
+
+BasicBSDFMipMap::BasicBSDFMipMap(NVRenderContext *inNVRenderContext, int inWidth, int inHeight,
+                                 NVRenderTexture2D &inTexture2D,
+                                 NVRenderTextureFormats::Enum inDestFormat, NVFoundationBase &inFnd)
+    : BSDFMipMap(inNVRenderContext, inWidth, inHeight, inTexture2D, inDestFormat, inFnd)
+{
+}
+
+BSDFMipMap::~BSDFMipMap()
+{
+}
+
+void BasicBSDFMipMap::Build(void *inTextureData, int inTextureDataSize,
+                            NVRenderBackend::NVRenderBackendTextureObject,
+                            NVRenderTextureFormats::Enum inFormat)
+{
+
+    m_InternalFormat = inFormat;
+    m_SizeOfInternalFormat = NVRenderTextureFormats::getSizeofFormat(m_InternalFormat);
+    m_InternalNoOfComponent = NVRenderTextureFormats::getNumberOfComponent(m_InternalFormat);
+
+    m_Texture2D.SetTextureData(NVDataRef<QT3DSU8>((QT3DSU8 *)inTextureData, inTextureDataSize), 0,
+                               m_Width, m_Height, inFormat, m_DestinationFormat);
+
+    STextureData theMipImage;
+    STextureData prevImage;
+    prevImage.data = inTextureData;
+    prevImage.dataSizeInBytes = inTextureDataSize;
+    prevImage.format = inFormat;
+    int curWidth = m_Width;
+    int curHeight = m_Height;
+    int size = NVRenderTextureFormats::getSizeofFormat(m_InternalFormat);
+    for (int idx = 1; idx <= m_MaxMipMapLevel; ++idx) {
+        theMipImage =
+            CreateBsdfMipLevel(theMipImage, prevImage, curWidth, curHeight); //, m_PerfTimer );
+        curWidth = curWidth >> 1;
+        curHeight = curHeight >> 1;
+        curWidth = curWidth >= 1 ? curWidth : 1;
+        curHeight = curHeight >= 1 ? curHeight : 1;
+        inTextureDataSize = curWidth * curHeight * size;
+
+        m_Texture2D.SetTextureData(toU8DataRef((char *)theMipImage.data, (QT3DSU32)inTextureDataSize),
+                                   (QT3DSU8)idx, (QT3DSU32)curWidth, (QT3DSU32)curHeight, theMipImage.format,
+                                   m_DestinationFormat);
+
+        if (prevImage.data == inTextureData)
+            prevImage = STextureData();
+
+        STextureData temp = prevImage;
+        prevImage = theMipImage;
+        theMipImage = temp;
+    }
+    QT3DS_FREE(m_Foundation.getAllocator(), theMipImage.data);
+    QT3DS_FREE(m_Foundation.getAllocator(), prevImage.data);
+}
+
+inline int BasicBSDFMipMap::wrapMod(int a, int base)
+{
+    return (a >= 0) ? a % base : (a % base) + base;
+}
+
+inline void BasicBSDFMipMap::getWrappedCoords(int &sX, int &sY, int width, int height)
+{
+    if (sY < 0) {
+        sX -= width >> 1;
+        sY = -sY;
+    }
+    if (sY >= height) {
+        sX += width >> 1;
+        sY = height - sY;
+    }
+    sX = wrapMod(sX, width);
+}
+
+STextureData BasicBSDFMipMap::CreateBsdfMipLevel(STextureData &inCurMipLevel,
+                                                 STextureData &inPrevMipLevel, int width,
+                                                 int height) //, IPerfTimer& inPerfTimer )
+{
+    //	SStackPerfTimer __timer( inPerfTimer, "Image BSDF Mip Level" );
+    STextureData retval;
+    int newWidth = width >> 1;
+    int newHeight = height >> 1;
+    newWidth = newWidth >= 1 ? newWidth : 1;
+    newHeight = newHeight >= 1 ? newHeight : 1;
+
+    if (inCurMipLevel.data) {
+        retval = inCurMipLevel;
+        retval.dataSizeInBytes =
+            newWidth * newHeight * NVRenderTextureFormats::getSizeofFormat(inPrevMipLevel.format);
+    } else {
+        retval.dataSizeInBytes =
+            newWidth * newHeight * NVRenderTextureFormats::getSizeofFormat(inPrevMipLevel.format);
+        retval.format = inPrevMipLevel.format; // inLoadedImage.format;
+        retval.data = m_Foundation.getAllocator().allocate(
+            retval.dataSizeInBytes, "Bsdf Scaled Image Data", __FILE__, __LINE__);
+    }
+
+    for (int y = 0; y < newHeight; ++y) {
+        for (int x = 0; x < newWidth; ++x) {
+            float accumVal[4];
+            accumVal[0] = 0;
+            accumVal[1] = 0;
+            accumVal[2] = 0;
+            accumVal[3] = 0;
+            for (int sy = -2; sy <= 2; ++sy) {
+                for (int sx = -2; sx <= 2; ++sx) {
+                    int sampleX = sx + (x << 1);
+                    int sampleY = sy + (y << 1);
+                    getWrappedCoords(sampleX, sampleY, width, height);
+
+                    // Cauchy filter (this is simply because it's the easiest to evaluate, and
+                    // requires no complex
+                    // functions).
+                    float filterPdf = 1.f / (1.f + float(sx * sx + sy * sy) * 2.f);
+                    // With FP HDR formats, we're not worried about intensity loss so much as
+                    // unnecessary energy gain,
+                    // whereas with LDR formats, the fear with a continuous normalization factor is
+                    // that we'd lose
+                    // intensity and saturation as well.
+                    filterPdf /= (NVRenderTextureFormats::getSizeofFormat(retval.format) >= 8)
+                        ? 4.71238898f
+                        : 4.5403446f;
+                    // filterPdf /= 4.5403446f;		// Discrete normalization factor
+                    // filterPdf /= 4.71238898f;		// Continuous normalization factor
+                    float curPix[4];
+                    QT3DSI32 byteOffset = (sampleY * width + sampleX)
+                        * NVRenderTextureFormats::getSizeofFormat(retval.format);
+                    if (byteOffset < 0) {
+                        sampleY = height + sampleY;
+                        byteOffset = (sampleY * width + sampleX)
+                            * NVRenderTextureFormats::getSizeofFormat(retval.format);
+                    }
+
+                    NVRenderTextureFormats::decodeToFloat(inPrevMipLevel.data, byteOffset, curPix,
+                                                          retval.format);
+
+                    accumVal[0] += filterPdf * curPix[0];
+                    accumVal[1] += filterPdf * curPix[1];
+                    accumVal[2] += filterPdf * curPix[2];
+                    accumVal[3] += filterPdf * curPix[3];
+                }
+            }
+
+            /*
+            // Re-adjustment after the fact for the RGBD hack.
+            if (retval.format == NVRenderTextureFormats::RGBA8 || retval.format ==
+            NVRenderTextureFormats::SRGB8A8)
+            {
+                    float divVal = (accumVal[0] > accumVal[1]) ? accumVal[0] : accumVal[1];
+                    divVal = (divVal > accumVal[2]) ? divVal : accumVal[2];
+                    if (divVal > 1.0)
+                    {
+                            divVal = 1.0f / divVal;
+                            accumVal[0] *= divVal;
+                            accumVal[1] *= divVal;
+                            accumVal[2] *= divVal;
+                            accumVal[3] = divVal;
+                    }
+                    else
+                            accumVal[3] = 1.0f;
+            }
+            */
+            QT3DSU32 newIdx =
+                (y * newWidth + x) * NVRenderTextureFormats::getSizeofFormat(retval.format);
+
+            NVRenderTextureFormats::encodeToPixel(accumVal, retval.data, newIdx, retval.format);
+        }
+    }
+
+    return retval;
+}
diff --git a/src/hdr/MipmapBSDF.cu b/src/hdr/MipmapBSDF.cu
new file mode 100644
index 0000000..6ddac4b
--- /dev/null
+++ b/src/hdr/MipmapBSDF.cu
@@ -0,0 +1,404 @@
+/****************************************************************************
+**
+** Copyright (C) 2017 The Qt Company Ltd.
+** Contact: https://www.qt.io/licensing/
+**
+** This file is part of Qt 3D Studio.
+**
+** $QT_BEGIN_LICENSE:GPL$
+** Commercial License Usage
+** Licensees holding valid commercial Qt licenses may use this file in
+** accordance with the commercial license agreement provided with the
+** Software or, alternatively, in accordance with the terms contained in
+** a written agreement between you and The Qt Company. For licensing terms
+** and conditions see https://www.qt.io/terms-conditions. For further
+** information use the contact form at https://www.qt.io/contact-us.
+**
+** GNU General Public License Usage
+** Alternatively, this file may be used under the terms of the GNU
+** General Public License version 3 or (at your option) any later version
+** approved by the KDE Free Qt Foundation. The licenses are as published by
+** the Free Software Foundation and appearing in the file LICENSE.GPL3
+** included in the packaging of this file. Please review the following
+** information to ensure the GNU General Public License requirements will
+** be met: https://www.gnu.org/licenses/gpl-3.0.html.
+**
+** $QT_END_LICENSE$
+**
+****************************************************************************/
+
+#if defined (_PLATFORM_USE_EGL)
+#include <GLES31/gl31.h>
+#include <GLES31/gl2ext.h>
+#endif
+
+#include "CUDABSDFMipmap.h"
+#include "cuda.h"
+#include "cuda_runtime.h"
+#include "cuda_gl_interop.h"
+#include <iostream>
+
+using namespace nv;
+using namespace nv::render;
+__host__ void jerror1(cudaError error)
+{
+   static  int i = 0;
+   ++i;
+}
+#ifdef _DEBUG
+#define CHECK_AND_HANDLE_CUDA_ERROR(func)													\
+		func;																			\
+		{																				\
+			cudaError error = cudaGetLastError();												\
+			if ( error != cudaSuccess )													\
+			{																			\
+				printf("%s\n", cudaGetErrorString(error));								\
+                jerror1(error);\
+				NV_ASSERT( false );														\
+			}																			\
+		}
+#else
+#define CHECK_AND_HANDLE_CUDA_ERROR(func)													\
+func;
+#endif
+
+__device__ inline int wrapMod( int a, int base )
+{
+	int ret = a % base;
+	if (ret < 0 ) ret += base;
+	return ret;
+}
+
+__device__ inline void getWrappedCoords( int &sX, int &sY, int width, int height )
+{
+	if (sY < 0) { sX -= width >> 1; sY = -sY; }
+	if (sY >= height) { sX += width >> 1; sY = height - sY; }
+	sX = wrapMod( sX, width );
+	sY = wrapMod( sY, height );
+}
+
+__device__ void decodeToFloat( void *inPtr, NVU32 byteOfs, float *outPtr, NVRenderTextureFormats::Enum inFmt, unsigned int numberOfComponent )
+{
+	outPtr[0] = 0.0f;	outPtr[1] = 0.0f;	outPtr[2] = 0.0f;	outPtr[3] = 0.0f;
+	NVU8 *src = reinterpret_cast<NVU8 *>(inPtr);
+	//float divisor;		// If we want to support RGBD?
+	switch(inFmt)
+	{
+	case NVRenderTextureFormats::Alpha8:
+		outPtr[0] = ((float)src[byteOfs]) / 255.0f;
+		break;
+
+	case NVRenderTextureFormats::Luminance8:
+	case NVRenderTextureFormats::LuminanceAlpha8:
+	case NVRenderTextureFormats::R8:
+	case NVRenderTextureFormats::RG8:
+	case NVRenderTextureFormats::RGB8:
+	case NVRenderTextureFormats::RGBA8:
+	case NVRenderTextureFormats::SRGB8:
+	case NVRenderTextureFormats::SRGB8A8:
+		// NOTE : RGBD Hack here for reference.  Not meant for installation.
+		//divisor = (NVRenderTextureFormats::getSizeofFormat(inFmt) == 4) ? ((float)src[byteOfs+3]) / 255.0f : 1.0f;
+		for ( NVU32 i = 0; i < numberOfComponent; ++i )
+		{
+			float val = ((float)src[byteOfs + i]) / 255.0f;
+			outPtr[i] = (i < 3) ? powf(val, 0.4545454545f) : val;
+			// Assuming RGBA8 actually means RGBD (which is stupid, I know)
+			//if ( NVRenderTextureFormats::getSizeofFormat(inFmt) == 4 ) { outPtr[i] /= divisor; }
+		}
+		//outPtr[3] = divisor;
+		break;
+
+	case NVRenderTextureFormats::RGBA32F:
+		outPtr[0] = reinterpret_cast<float *>(src+byteOfs)[0];
+		outPtr[1] = reinterpret_cast<float *>(src+byteOfs)[1];
+		outPtr[2] = reinterpret_cast<float *>(src+byteOfs)[2];
+		outPtr[3] = reinterpret_cast<float *>(src+byteOfs)[3];
+		break;
+	case NVRenderTextureFormats::RGB32F:
+		outPtr[0] = reinterpret_cast<float *>(src+byteOfs)[0];
+		outPtr[1] = reinterpret_cast<float *>(src+byteOfs)[1];
+		outPtr[2] = reinterpret_cast<float *>(src+byteOfs)[2];
+		break;
+
+	case NVRenderTextureFormats::RGBA16F:
+        /*
+		for ( NVU32 i = 0; i < 4; ++i )
+		{
+			// NOTE : This only works on the assumption that we don't have any denormals, Infs or NaNs.
+			// Every pixel in our source image should be "regular"
+			NVU16 h = reinterpret_cast<NVU16 *>(src + byteOfs)[i];
+			NVU32 sign = (h & 0x8000) << 16;
+			NVU32 exponent = (((((h & 0x7c00) >> 10) - 15) + 127) << 23);
+			NVU32 mantissa =  ((h & 0x3ff) << 13);
+			NVU32 result = sign | exponent | mantissa;
+					
+			if (h == 0 || h == 0x8000) { result = 0; }	// Special case for zero and negative zero
+			memcpy( reinterpret_cast<NVU32 *>(outPtr) + i, &result, 4 );
+		}*/
+
+		for ( NVU32 i = 0; i < 2; i++ )
+		{
+			// NOTE : This only works on the assumption that we don't have any denormals, Infs or NaNs.
+			// Every pixel in our source image should be "regular"
+
+			NVU32 h1 = reinterpret_cast<NVU32 *>(src + byteOfs)[i];
+
+			for ( NVU8 j = 0; j < 2; j++ ) 
+			{
+				NVU16 h = (h1 & (0x0000FFFF << j*16 )) >> j*16;
+				NVU32 sign = (h & 0x8000) << 16;
+				NVU32 exponent = (((((h & 0x7c00) >> 10) - 15) + 127) << 23);
+				NVU32 mantissa =  ((h & 0x3ff) << 13);
+				NVU32 result = sign | exponent | mantissa;
+					
+				if (h == 0 || h == 0x8000) { result = 0; }	// Special case for zero and negative zero
+				memcpy( reinterpret_cast<NVU32 *>(outPtr) + i*2 + j, &result, 4 );
+			}
+		}
+		break;
+
+	case NVRenderTextureFormats::R11G11B10:
+		// place holder
+		NV_ASSERT( false );
+		break;
+
+	default:
+		outPtr[0] = 0.0f;
+		outPtr[1] = 0.0f;
+		outPtr[2] = 0.0f;
+		outPtr[3] = 0.0f;
+		break;
+	}
+}
+
+void __device__ encodeToPixel( float *inPtr, void *outPtr, NVU32 byteOfs, NVRenderTextureFormats::Enum inFmt, unsigned int noOfComponent )
+{
+	NVU8 *dest = reinterpret_cast<NVU8 *>(outPtr);
+	switch(inFmt)
+	{
+	case NVRenderTextureFormats::Alpha8:
+		dest[byteOfs] = NVU8( inPtr[0] * 255.0f );
+		break;
+
+	case NVRenderTextureFormats::Luminance8:
+	case NVRenderTextureFormats::LuminanceAlpha8:
+	case NVRenderTextureFormats::R8:
+	case NVRenderTextureFormats::RG8:
+	case NVRenderTextureFormats::RGB8:
+	case NVRenderTextureFormats::RGBA8:
+	case NVRenderTextureFormats::SRGB8:
+	case NVRenderTextureFormats::SRGB8A8:
+		for ( NVU32 i = 0; i < noOfComponent; ++i )
+		{
+			inPtr[i] = (inPtr[i] > 1.0f) ? 1.0f : inPtr[i];
+			if (i < 3)
+				dest[byteOfs+i] = NVU8( powf( inPtr[i], 2.2f ) * 255.0f);
+			else
+				dest[byteOfs+i] = NVU8( inPtr[i] * 255.0f );
+		}
+		break;
+
+	case NVRenderTextureFormats::RGBA32F:
+		reinterpret_cast<float *>(dest+byteOfs)[0] = inPtr[0];
+		reinterpret_cast<float *>(dest+byteOfs)[1] = inPtr[1];
+		reinterpret_cast<float *>(dest+byteOfs)[2] = inPtr[2];
+		reinterpret_cast<float *>(dest+byteOfs)[3] = inPtr[3];
+		break;
+	case NVRenderTextureFormats::RGB32F:
+		reinterpret_cast<float *>(dest+byteOfs)[0] = inPtr[0];
+		reinterpret_cast<float *>(dest+byteOfs)[1] = inPtr[1];
+		reinterpret_cast<float *>(dest+byteOfs)[2] = inPtr[2];
+		break;
+
+	case NVRenderTextureFormats::RGBA16F:
+		for ( NVU32 i = 0; i < 4; ++i )
+		{
+			// NOTE : This also has the limitation of not handling  infs, NaNs and denormals, but it should be
+			// sufficient for our purposes.
+			if (inPtr[i] > 65519.0f) { inPtr[i] = 65519.0f; }
+			if (fabs(inPtr[i]) < 6.10352E-5f) { inPtr[i] = 0.0f; }
+			NVU32 f = reinterpret_cast<NVU32 *>(inPtr)[i];
+			NVU32 sign = (f & 0x80000000) >> 16;
+			NVI32 exponent = (f & 0x7f800000) >> 23;
+			NVU32 mantissa = (f >> 13) & 0x3ff;
+			exponent = exponent - 112;
+			if (exponent > 31) { exponent = 31; }
+			if (exponent < 0) { exponent = 0; }
+			exponent = exponent << 10;
+			reinterpret_cast<NVU16 *>(dest + byteOfs)[i] = NVU16(sign | exponent | mantissa);
+		}
+		break;
+
+	case NVRenderTextureFormats::R11G11B10:
+		// place holder
+		NV_ASSERT( false );
+		break;
+
+	default:
+		dest[byteOfs] = 0;
+		dest[byteOfs+1] = 0;
+		dest[byteOfs+2] = 0;
+		dest[byteOfs+3] = 0;
+		break;
+	}
+}
+
+void __global__ Convert3To4Component( cudaTextureObject_t tex, float *d_outBuffer, Q3DStudio::INT32 dpitch, Q3DStudio::INT32 width, Q3DStudio::INT32 height )
+{
+	float *dest = d_outBuffer;
+
+	int x = threadIdx.x + blockIdx.x * blockDim.x;
+	int y = threadIdx.y + blockIdx.y * blockDim.y;
+	if ( x >= width || y >= height ) 
+		return;
+	int inX = x * 3;
+	int outX = x * 4;
+	dest[outX + y * width * 4] = tex2D<float>(tex, inX, y);
+	dest[outX + y * width * 4 + 1] = tex2D<float>(tex, inX + 1, y);
+	dest[outX + y * width * 4 + 2] = tex2D<float>(tex, inX + 2, y);
+	dest[outX + y * width * 4 + 3] = 255 * 255;
+}
+
+void __global__ ConvertData( void* d_InBuffer, NVRenderTextureFormats::Enum inFmt, int inSizeOfFormat, int inNoOfComponent, int inPitch, 
+                             void* d_OutBuffer, NVRenderTextureFormats::Enum outFmt, int outSizeOfFormat, int outNoOfComponent, int outPitch, int width, int height )
+{
+
+	int x = threadIdx.x + blockIdx.x * blockDim.x;
+	int y = threadIdx.y + blockIdx.y * blockDim.y;
+	if ( x >= width || y >= height ) 
+		return;
+	float values[4];
+
+	decodeToFloat( d_InBuffer, (inPitch * y) + (x * inSizeOfFormat), values, inFmt, inNoOfComponent );
+	encodeToPixel( values, d_OutBuffer, (outPitch * y) + (x * outSizeOfFormat), outFmt, outSizeOfFormat );
+}
+
+void __global__ CreateBsdfMipLevel( cudaTextureObject_t tex, void *d_curBuffer, void *d_prevBuffer, 	Q3DStudio::INT32 pitch, Q3DStudio::INT32 width, Q3DStudio::INT32 height, 
+										nv::render::NVRenderTextureFormats::Enum inFormat, unsigned int sizeOfFormat )
+{
+	float accumVal[4];
+	//unsigned int sizeofFormat = getSizeofFormat(inFormat);
+	//__shared__ float dataBlock[ ]; //(32+4) * (32+4) * 12 
+	int x = threadIdx.x + blockIdx.x * blockDim.x;
+	int y = threadIdx.y + blockIdx.y * blockDim.y;
+
+	if ( x >= (width > 2 ? width >> 1 : 1) || y >= (height > 2 ? height >> 1 : 1)) return;
+
+	accumVal[0] = 0;	accumVal[1] = 0;	accumVal[2] = 0;	accumVal[3] = 0;
+
+	for ( int sy = -2; sy <= 2; ++sy )
+	{
+		for ( int sx = -2; sx <= 2; ++sx )
+		{
+			int sampleX = sx + (x << 1);
+			int sampleY = sy + (y << 1);
+			//getWrappedCoords(sampleX, sampleY, width, height);
+			// Cauchy filter (this is simply because it's the easiest to evaluate, and requires no complex
+			// functions).
+			float filterPdf = 1.f / ( 1.f + float(sx*sx + sy*sy)*2.f );
+			// With FP HDR formats, we're not worried about intensity loss so much as unnecessary energy gain, 
+			// whereas with LDR formats, the fear with a continuous normalization factor is that we'd lose
+			// intensity and saturation as well.
+			filterPdf /= sizeOfFormat >= 8 ? 4.71238898f : 4.5403446f;
+			//filterPdf /= 4.5403446f;		// Discrete normalization factor
+			//filterPdf /= 4.71238898f;		// Continuous normalization factor
+			//float curPix[4];
+			sampleX = sampleX*4;
+			getWrappedCoords(sampleX, sampleY, width*4, height);
+			accumVal[0] += filterPdf * tex2D<float>(tex, sampleX, sampleY);
+			accumVal[1] += filterPdf * tex2D<float>(tex, sampleX + 1, sampleY);
+			accumVal[2] += filterPdf * tex2D<float>(tex, sampleX + 2, sampleY);
+			accumVal[3] += filterPdf * tex2D<float>(tex, sampleX + 3, sampleY);
+		}
+		}
+
+	encodeToPixel(accumVal, d_curBuffer, y * pitch + x * sizeOfFormat, inFormat, sizeOfFormat);
+}
+
+struct SMipTextureData
+{
+	void* data;
+	unsigned int dataSize;
+	unsigned int mipLevel;
+	unsigned int width;
+	unsigned int height;
+	NVRenderTextureFormats::Enum format;	
+};
+
+__host__  void CUDABSDFMipMap::Build( void* inTextureData, int inTextureDataSize, NVRenderBackend::NVRenderBackendTextureObject inTextureHandle, NVRenderTextureFormats::Enum inFormat  )
+{
+	m_TextureHandle = inTextureHandle;
+	m_InternalFormat = inFormat;
+	m_SizeOfInternalFormat = NVRenderTextureFormats::getSizeofFormat( m_InternalFormat );
+	m_InternalNoOfComponent = NVRenderTextureFormats::getNumberOfComponent( m_InternalFormat );
+
+	m_Texture2D.SetTextureData(  NVDataRef<NVU8>( (NVU8*)inTextureData, inTextureDataSize )
+								, 0
+								, m_Width
+								, m_Height
+								, inFormat
+								, m_DestinationFormat );
+
+	size_t pitch;
+	float* d_inTextureData;
+
+	cudaMallocPitch(&d_inTextureData, &pitch, m_Width * m_SizeOfInternalFormat, m_Height); CHECK_AND_HANDLE_CUDA_ERROR();
+	CHECK_AND_HANDLE_CUDA_ERROR( cudaMemcpy2D( d_inTextureData, pitch, inTextureData, m_Width * m_SizeOfInternalFormat, m_Width * m_SizeOfInternalFormat, m_Height, cudaMemcpyHostToDevice ) );
+	{
+		dim3 blockDim(16, 16, 1);	
+		dim3 gridDim(ceil(m_Width / 16.0f), ceil(m_Height / 16.0f) ,1 );
+
+		//std::cerr << "if= " << m_InternalFormat << " sizeOut= " << m_SizeOfInternalFormat << " numOfIntComp" << m_InternalNoOfComponent << " pitch= " << pitch << " destFormat= " << m_DestinationFormat << " sizeFormat= " << m_SizeOfFormat << " numOfComp= " << m_NoOfComponent << " Pitch0=" << m_Pitches[0] << std::endl;
+		//NVLogWarn("cuda", "%i %i %i %i %i %i %i %i\n",(int)m_InternalFormat ,m_SizeOfInternalFormat ,m_InternalNoOfComponent , pitch, (int)m_DestinationFormat, m_SizeOfFormat, m_NoOfComponent ,m_Pitches[0]);
+        	ConvertData<<<gridDim, blockDim>>>( d_inTextureData, m_InternalFormat, m_SizeOfInternalFormat, m_InternalNoOfComponent, pitch,
+                        md_MipMapsData[0], m_DestinationFormat, m_SizeOfFormat, m_NoOfComponent, m_Pitches[0], m_Width, m_Height );
+	}
+	cudaFree(d_inTextureData);
+
+	int curWidth  = m_Width;
+	int curHeight = m_Height;
+
+	cudaTextureObject_t* tex;
+	tex = new cudaTextureObject_t[m_MaxMipMapLevel];
+	for ( int idx = 1; idx <= m_MaxMipMapLevel; ++idx )
+	{
+		tex[idx-1] = -1;
+		dim3 blockDim(16, 16, 1);	
+		dim3 gridDim(ceil(curWidth / 32.0f), ceil(curHeight / 32.0f) ,1 );
+
+		cudaResourceDesc resDesc;
+		memset(&resDesc, 0, sizeof(resDesc));
+		resDesc.res.pitch2D.desc.f = cudaChannelFormatKindFloat;
+		resDesc.res.pitch2D.desc.x = m_SizeOfFormat / m_NoOfComponent * 8; // bits per channel
+		resDesc.resType = cudaResourceTypePitch2D;
+		resDesc.res.pitch2D.devPtr = (char*)(md_MipMapsData[idx-1]);
+		resDesc.res.pitch2D.height = curHeight;
+		resDesc.res.pitch2D.width  = curWidth * m_NoOfComponent;
+		resDesc.res.pitch2D.pitchInBytes  = m_Pitches[idx-1];// aligned to texturePitchAlignment
+
+		cudaTextureDesc texDesc;
+		memset(&texDesc, 0, sizeof(texDesc));
+		texDesc.addressMode[0] = cudaAddressModeWrap;
+		texDesc.addressMode[1] = cudaAddressModeWrap;
+		texDesc.readMode = cudaReadModeElementType;
+		//texDesc.normalizedCoords = 1;
+
+		
+		CHECK_AND_HANDLE_CUDA_ERROR( cudaCreateTextureObject( &tex[idx-1], &resDesc, &texDesc, NULL ) );
+		CreateBsdfMipLevel<<<gridDim, blockDim>>>( tex[idx-1], (reinterpret_cast<NVU8 *>(md_MipMapsData[idx])), (reinterpret_cast<NVU8 *>(md_MipMapsData[idx-1])), m_Pitches[idx], curWidth, curHeight, m_DestinationFormat, m_SizeOfFormat );
+
+		curWidth  = curWidth > 2 ? curWidth >> 1 : 1;
+		curHeight = curHeight > 2 ? curHeight >> 1 : 1;
+	}
+
+	CHECK_AND_HANDLE_CUDA_ERROR( cudaDeviceSynchronize(); )
+	BindTexture();
+	TransferTexture();
+	for (int idx = 0; idx < m_MaxMipMapLevel;++idx )
+		cudaDestroyTextureObject(tex[idx]);
+//	CHECK_AND_HANDLE_CUDA_ERROR( cudaDeviceReset(); )
+	CHECK_AND_HANDLE_CUDA_ERROR( cudaDeviceSynchronize(); )
+
+    //NV_FREE( m_Foundation.getAllocator(), inTextureData );
+
+}
diff --git a/src/hdr/MipmapBSDF.h b/src/hdr/MipmapBSDF.h
new file mode 100644
index 0000000..be8f136
--- /dev/null
+++ b/src/hdr/MipmapBSDF.h
@@ -0,0 +1,104 @@
+/****************************************************************************
+**
+** Copyright (C) 2017 The Qt Company Ltd.
+** Contact: https://www.qt.io/licensing/
+**
+** This file is part of Qt 3D Studio.
+**
+** $QT_BEGIN_LICENSE:GPL$
+** Commercial License Usage
+** Licensees holding valid commercial Qt licenses may use this file in
+** accordance with the commercial license agreement provided with the
+** Software or, alternatively, in accordance with the terms contained in
+** a written agreement between you and The Qt Company. For licensing terms
+** and conditions see https://www.qt.io/terms-conditions. For further
+** information use the contact form at https://www.qt.io/contact-us.
+**
+** GNU General Public License Usage
+** Alternatively, this file may be used under the terms of the GNU
+** General Public License version 3 or (at your option) any later version
+** approved by the KDE Free Qt Foundation. The licenses are as published by
+** the Free Software Foundation and appearing in the file LICENSE.GPL3
+** included in the packaging of this file. Please review the following
+** information to ensure the GNU General Public License requirements will
+** be met: https://www.gnu.org/licenses/gpl-3.0.html.
+**
+** $QT_END_LICENSE$
+**
+****************************************************************************/
+
+#ifndef MIPMAPBSDF_H
+#define MIPMAPBSDF_H
+#include "foundation/Qt3DSVec3.h"
+#include "foundation/Qt3DSSimpleTypes.h"
+#include "foundation/Qt3DSPerfTimer.h"
+#include "foundation/Qt3DSAtomic.h"
+#include "render/Qt3DSRenderBaseTypes.h"
+#include "render/backends/Qt3DSRenderBackend.h"
+#include "render/Qt3DSRenderTexture2D.h"
+#include "render/backends/gl/Qt3DSOpenGLUtil.h"
+
+#include "Qt3DSRenderLoadedTexture.h"
+
+#include "Qt3DSTypes.h"
+#ifdef _LINUXPLATFORM
+#define __declspec(dllexport)
+#define __cdecl
+
+#endif
+using namespace qt3ds::render;
+
+class BSDFMipMap : public NVRefCounted
+{
+public:
+    BSDFMipMap(NVRenderContext *inNVRenderContext, int inWidth, int inHeight,
+               NVRenderTexture2D &inTexture, NVRenderTextureFormats::Enum inDestFormat,
+               qt3ds::NVFoundationBase &inFnd);
+    virtual ~BSDFMipMap();
+
+    virtual void Build(void *inTextureData, int inTextureDataSize,
+                       NVRenderBackend::NVRenderBackendTextureObject inTextureHandle,
+                       NVRenderTextureFormats::Enum inFormat) = 0;
+    static BSDFMipMap *Create(NVRenderContext *inNVRenderContext, int inWidth, int inHeight,
+                              NVRenderTexture2D &inTexture,
+                              NVRenderTextureFormats::Enum inDestFormat,
+                              qt3ds::NVFoundationBase &inFnd);
+
+protected:
+    volatile QT3DSI32 mRefCount; ///< reference count
+    NVFoundationBase &m_Foundation; ///< Foundation class for allocations and other base things
+
+    NVRenderTexture2D &m_Texture2D;
+    NVRenderTextureFormats::Enum m_InternalFormat;
+    NVRenderTextureFormats::Enum m_DestinationFormat;
+    int m_Width;
+    int m_Height;
+    int m_MaxMipMapLevel;
+    int m_SizeOfFormat;
+    int m_SizeOfInternalFormat;
+    int m_InternalNoOfComponent;
+    int m_NoOfComponent;
+    NVRenderContext *m_NVRenderContext;
+};
+
+class BasicBSDFMipMap : public BSDFMipMap
+{
+public:
+    BasicBSDFMipMap(NVRenderContext *inNVRenderContext, int inWidth, int inHeight,
+                    NVRenderTexture2D &inTexture, NVRenderTextureFormats::Enum inDestFormat,
+                    qt3ds::NVFoundationBase &inFnd);
+
+    void Build(void *inTextureData, int inTextureDataSize,
+               NVRenderBackend::NVRenderBackendTextureObject inTextureHandle,
+               NVRenderTextureFormats::Enum inFormat);
+
+    STextureData CreateBsdfMipLevel(STextureData &inCurMipLevel, STextureData &inPrevMipLevel,
+                                    int width, int height); //, IPerfTimer& inPerfTimer );
+    QT3DS_IMPLEMENT_REF_COUNT_ADDREF_RELEASE(m_Foundation);
+
+    int wrapMod(int a, int base);
+    void getWrappedCoords(int &sX, int &sY, int width, int height);
+    NVRenderBackend::NVRenderBackendTextureObject m_TextureHandle;
+};
+
+#endif