diff --git a/.gitignore b/.gitignore new file mode 100644 index 0000000..e87e1e7 --- /dev/null +++ b/.gitignore @@ -0,0 +1,70 @@ +CMakeCache.txt +CMakeFiles +CMakeScripts +Makefile +cmake_install.cmake +install_manifest.txt +# Compiled Object files +*.slo +*.lo +*.o +*.obj + +# Precompiled Headers +*.gch +*.pch + +# Compiled Dynamic libraries +*.so +*.dylib +*.dll + +# Fortran module files +*.mod + +# Compiled Static libraries +*.lai +*.la +*.a +*.lib + +# Executables +*.exe +*.out +*.app + +*.i +*.ii +*.gpu +*.ptx +*.cubin +*.fatbin + +# Qt-es + +/.qmake.cache +/.qmake.stash +*.pro.user +*.pro.user.* +*.qbs.user +*.qbs.user.* +*.moc +moc_*.cpp +qrc_*.cpp +ui_*.h +Makefile* +*build-* + +# QtCreator + +*.autosave + +# QtCtreator Qml +*.qmlproject.user +*.qmlproject.user.* + +# QtCtreator CMake +CMakeLists.txt.user + +build/* +.idea/* \ No newline at end of file diff --git a/CMakeLists.txt b/CMakeLists.txt new file mode 100644 index 0000000..14d0dfe --- /dev/null +++ b/CMakeLists.txt @@ -0,0 +1,49 @@ +cmake_minimum_required(VERSION 3.3) +project(SunVolumeRender) + +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11") +set(CMAKE_INCLUDE_CURRENT_DIR ON) + +# CUDA settings +find_package(CUDA QUIET REQUIRED) +list(APPEND CUDA_NVCC_FLAGS --compiler-options -fno-strict-aliasing -use_fast_math -Xptxas -v -maxrregcount=32) +list(APPEND CUDA_NVCC_FLAGS -gencode arch=compute_30,code=sm_30) +set(CUDA_INCLUDE_DIRS /usr/local/cuda/include/) +include_directories(${CUDA_INCLUDE_DIRS}) + +#Qt +find_package(Qt4 COMPONENTS QtCore QtGui QtOpenGL REQUIRED) +include(${QT_USE_FILE}) +set(QT_USE_QTOPENGL TRUE) +set(CMAKE_AUTOMOC ON) +set(CMAKE_AUTOUIC ON) +set(CMAKE_AUTORCC ON) +set(RESOURCES_FILES qdarkstyle/style.qrc) + +#VTK +set(VTK_DIR /Volumes/Free/SDK/CTK/build/VTK-build) +find_package(VTK REQUIRED) +include(${VTK_USE_FILE}) + +#CTK +set(CTK_DIR /Volumes/Free/SDK/CTK/build/) +#set(CTK_DIR /Users/sunwj/Desktop/SDK/CTK/build) +find_package(CTK REQUIRED) +include(${CTK_USE_FILE}) +set(CTK_LIBRARIES CTKCore CTKWidgets CTKVisualizationVTKCore CTKVisualizationVTKWidgets) + +#GLM +find_package(GLM REQUIRED) +include_directories(${GLM_INCLUDE_DIRS}) + +set(HOST_SOURCES main.cpp + gui/mainwindow.cpp + gui/canvas.cpp + gui/transferfunction.cpp) + +set(DEVICE_SOURCES pathtracer.cu) + +cuda_compile(DEVICE_OBJS ${DEVICE_SOURCES}) + +add_executable(SunVolumeRender ${HOST_SOURCES} ${RESOURCES_FILES} ${DEVICE_OBJS}) +target_link_libraries(SunVolumeRender ${QT_LIBRARIES} ${VTK_LIBRARIES} ${CTK_LIBRARIES} ${CUDA_LIBRARIES}) diff --git a/common.h b/common.h new file mode 100644 index 0000000..3e2a1f1 --- /dev/null +++ b/common.h @@ -0,0 +1,11 @@ +// +// Created by 孙万捷 on 16/5/19. +// + +#ifndef SUNVOLUMERENDER_COMMON_H +#define SUNVOLUMERENDER_COMMON_H + +#define WIDTH 640 +#define HEIGHT 640 + +#endif //SUNVOLUMERENDER_COMMON_H diff --git a/core/cuda_bbox.h b/core/cuda_bbox.h new file mode 100644 index 0000000..3697143 --- /dev/null +++ b/core/cuda_bbox.h @@ -0,0 +1,56 @@ +// +// Created by 孙万捷 on 16/5/19. +// + +#ifndef SUNVOLUMERENDER_CUDA_BBOX_H +#define SUNVOLUMERENDER_CUDA_BBOX_H + +#define GLM_FORCE_INLINE +#include + +#include + +#include "cuda_ray.h" + +class cudaBBox +{ +public: + __host__ cudaBBox() {} + __host__ cudaBBox(const glm::vec3& vmin, const glm::vec3& vmax) + { + Set(vmin, vmax); + } + + __host__ void Set(const glm::vec3& vmin, const glm::vec3& vmax) + { + this->vmin = vmin; + this->vmax = vmax; + + invSize = 1.f / (vmax - vmin); + } + + __device__ bool Intersect(const cudaRay& ray, float* tNear, float* tFar) const + { + auto invDir = 1.f / ray.dir; + auto tbot = invDir * (vmin - ray.orig); + auto ttop = invDir * (vmax - ray.orig); + + auto tmin = glm::min(tbot, ttop); + auto tmax = glm::max(tbot, ttop); + + float largest_tmin = fmaxf(tmin.x, fmaxf(tmin.y, tmin.z)); + float smallest_tmax = fminf(tmax.x, fminf(tmax.y, tmax.z)); + + *tNear = largest_tmin; + *tFar = smallest_tmax; + + return smallest_tmax > largest_tmin; + } + +public: + glm::vec3 vmin = glm::vec3(glm::uninitialize); + glm::vec3 vmax = glm::vec3(glm::uninitialize); + glm::vec3 invSize = glm::vec3(glm::uninitialize); +}; + +#endif //SUNVOLUMERENDER_CUDA_BBOX_H diff --git a/core/cuda_camera.h b/core/cuda_camera.h new file mode 100644 index 0000000..52f330b --- /dev/null +++ b/core/cuda_camera.h @@ -0,0 +1,76 @@ +// +// Created by 孙万捷 on 16/2/6. +// + +#ifndef SUNPATHTRACER_CAMERA_H +#define SUNPATHTRACER_CAMERA_H + +#include +#include + +#define GLM_FORCE_INLINE +#include + +#include "cuda_ray.h" + +class cudaCamera +{ +public: + __host__ __device__ cudaCamera() {} + + __host__ __device__ cudaCamera(const glm::vec3& _pos, const glm::vec3& _u, const glm::vec3& _v, const glm::vec3& _w, float fovx = 45.f, unsigned int _imageW = 640, unsigned int _imageH = 480) + { + Setup(_pos, _u, _v, _w, fovx, _imageW, _imageH); + } + + __host__ __device__ cudaCamera(const glm::vec3& _pos, const glm::vec3& target, const glm::vec3& up, float fovx = 45.f, unsigned int _imageW = 640, unsigned int _imageH = 480) + { + Setup(_pos, target, up, fovx, _imageW, _imageH); + } + + __host__ __device__ void Setup(const glm::vec3& _pos, const glm::vec3& _u, const glm::vec3& _v, const glm::vec3& _w, float fovx, unsigned int _imageW, unsigned int _imageH) + { + pos = _pos; + u = _u; + v = _v; + w = _w; + imageW = _imageW; + imageH = _imageH; + aspectRatio = (float)imageW / (float)imageH; + tanFovxOverTwo = tanf(fovx * 0.5f * M_PI / 180.f); + } + + __host__ __device__ void Setup(const glm::vec3& _pos, const glm::vec3& target, const glm::vec3& up, float fovx, unsigned int _imageW, unsigned int _imageH) + { + pos = _pos; + w = normalize(pos - target); + u = cross(up, w); + v = cross(w, u); + imageW = _imageW; + imageH = _imageH; + aspectRatio = (float)imageW / (float)imageH; + tanFovxOverTwo = tanf(fovx * 0.5f * M_PI / 180.f); + } + + // TODO: depth of field + __device__ void GenerateRay(unsigned int x, unsigned int y, curandState& rng, cudaRay* ray) const + { + float nx = 2.f * ((x + curand_uniform(&rng)) / (imageW - 1.f)) - 1.f; + float ny = 2.f * ((y + curand_uniform(&rng)) / (imageH - 1.f)) - 1.f; + + nx = nx * aspectRatio * tanFovxOverTwo; + ny = ny * tanFovxOverTwo; + + ray->orig = pos; + ray->dir = normalize(nx * u + ny * v - w); + } + +public: + unsigned int imageW, imageH; + float aspectRatio; + float tanFovxOverTwo; + glm::vec3 pos; + glm::vec3 u, v, w; +}; + +#endif //SUNPATHTRACER_CAMERA_H diff --git a/core/cuda_onb.h b/core/cuda_onb.h new file mode 100644 index 0000000..391d6f1 --- /dev/null +++ b/core/cuda_onb.h @@ -0,0 +1,52 @@ +// +// Created by 孙万捷 on 16/2/27. +// + +#ifndef SUNPATHTRACER_ONB_H +#define SUNPATHTRACER_ONB_H + +#define GLM_FORCE_INLINE +#include + +#include + +class cudaONB +{ +public: + __device__ cudaONB(const glm::vec3& _w) + { + InitFromW(_w); + } + + __device__ cudaONB(const glm::vec3& _v, const glm::vec3& _w) + { + InitFromVW(_v, _w); + } + + __device__ void InitFromW(const glm::vec3& _w) + { + w = _w; + if(fabsf(w.x) > fabsf(w.y)) + { + float invLength = rsqrtf(w.x * w.x + w.z * w.z); + v = glm::vec3(-w.z * invLength, 0.f, w.x * invLength); + } + else + { + float invLength = rsqrtf(w.y * w.y + w.z * w.z); + v = glm::vec3(0.f, w.z * invLength, -w.y * invLength); + } + u = cross(v, w); + } + + __device__ void InitFromVW(const glm::vec3& _v, const glm::vec3& _w) + { + w = _w; + u = cross(_v, w); + v = cross(w, u); + } +public: + glm::vec3 u, v, w; +}; + +#endif //SUNPATHTRACER_ONB_H diff --git a/core/cuda_ray.h b/core/cuda_ray.h new file mode 100644 index 0000000..8c1252d --- /dev/null +++ b/core/cuda_ray.h @@ -0,0 +1,44 @@ +// +// Created by 孙万捷 on 16/2/6. +// + +#ifndef SUNPATHTRACER_RAY_H +#define SUNPATHTRACER_RAY_H + +#include + +#define GLM_FORCE_INLINE +#include + +#include + +class cudaRay +{ +public: + __device__ cudaRay() + { + tMin = 1e-8; + tMax = FLT_MAX; + } + + __device__ cudaRay(const glm::vec3& orig, const glm::vec3& dir, float tMin = 1e-8, float tMax = FLT_MAX) + { + this->orig = orig; + this->dir = dir; + this->tMin = tMin; + this->tMax = tMax; + } + + __device__ glm::vec3 PointOnRay(float t) const + { + return orig + t * dir; + } + +public: + glm::vec3 orig; + glm::vec3 dir; + mutable float tMin; + mutable float tMax; +}; + +#endif //SUNPATHTRACER_RAY_H diff --git a/core/cuda_transfer_function.h b/core/cuda_transfer_function.h new file mode 100644 index 0000000..d3ae31c --- /dev/null +++ b/core/cuda_transfer_function.h @@ -0,0 +1,55 @@ +// +// Created by 孙万捷 on 16/5/20. +// + +#ifndef SUNVOLUMERENDER_CUDA_TRANSFER_FUNCTION_H +#define SUNVOLUMERENDER_CUDA_TRANSFER_FUNCTION_H + +#define GLM_FORCE_INLINE +#include + +#include + +class cudaTransferFunction +{ +public: + __host__ __device__ void Set(const cudaTextureObject_t& tex) + { + this->tex = tex; + } + + __device__ glm::vec4 operator ()(float intensity) + { +#ifdef __CUDACC__ + auto val = tex1D(tex, intensity); + return glm::vec4(val.x, val.y, val.z, val.w); +#else + return glm::vec4(0.f); +#endif + } + + __device__ glm::vec3 GetColor(float intensity) + { +#ifdef __CUDACC__ + auto val = tex1D(tex, intensity); + return glm::vec3(val.x, val.y, val.z); +#else + return glm::vec3(0.f); +#endif + } + + __device__ float GetOpacity(float intensity) + { +#ifdef __CUDACC__ + auto val = tex1D(tex, intensity); + return val.w; +#else + return 0.f; +#endif + } + +private: + cudaTextureObject_t tex; +}; + +#endif //SUNVOLUMERENDER_CUDA_TRANSFER_FUNCTION_H diff --git a/core/cuda_volume.h b/core/cuda_volume.h new file mode 100644 index 0000000..c1642a1 --- /dev/null +++ b/core/cuda_volume.h @@ -0,0 +1,79 @@ +// +// Created by 孙万捷 on 16/5/19. +// + +#ifndef SUNVOLUMERENDER_VOLUME_H +#define SUNVOLUMERENDER_VOLUME_H + +#include + +#define GLM_FORCE_INLINE +#include + +#include "core/cuda_bbox.h" + +class cudaVolume +{ +public: + __device__ float operator ()(const glm::vec3& pointInWorld) + { + return GetIntensity(pointInWorld); + } + + __device__ float operator ()(const glm::vec3& normalizedTexCoord, bool dummy) + { + return GetIntensityNTC(normalizedTexCoord); + } + + __device__ bool Intersect(const cudaRay& ray, float* tNear, float* tFar) + { + return bbox.Intersect(ray, tNear, tFar); + } + + __device__ glm::vec3 Gradient_CentralDiff(const glm::vec3& pointInWorld) const + { + auto xdiff = GetIntensity(pointInWorld + glm::vec3(spacing.x, 0.f, 0.f)) - GetIntensity(pointInWorld - glm::vec3(spacing.x, 0.f, 0.f)); + auto ydiff = GetIntensity(pointInWorld + glm::vec3(0.f, spacing.y, 0.f)) - GetIntensity(pointInWorld - glm::vec3(0.f, spacing.y, 0.f)); + auto zdiff = GetIntensity(pointInWorld + glm::vec3(0.f, 0.f, spacing.z)) - GetIntensity(pointInWorld - glm::vec3(0.f, 0.f, spacing.z)); + + return glm::vec3(xdiff, ydiff, zdiff) * 0.5f; + } + + __device__ glm::vec3 NormalizedGradient(const glm::vec3& pointInWorld) const + { + return glm::normalize(Gradient_CentralDiff(pointInWorld)); + } + +private: + __device__ glm::vec3 GetNormalizedTexCoord(const glm::vec3& pointInWorld) const + { + return (pointInWorld - bbox.vmin) * bbox.invSize; + } + + __device__ float GetIntensity(const glm::vec3& pointInWorld) const + { +#ifdef __CUDACC__ + auto texCoord = GetNormalizedTexCoord(pointInWorld); + return tex3D(tex, texCoord.x, texCoord.y, texCoord.z); +#else + return 0.f; +#endif + } + + __device__ float GetIntensityNTC(const glm::vec3& normalizedTexCoord) + { +#ifdef __CUDACC__ + return tex3D(tex, normalizedTexCoord.x, normalizedTexCoord.y, normalizedTexCoord.z); +#else + return 0.f; +#endif + } + +private: + cudaBBox bbox; + cudaTextureObject_t tex = 0; + glm::vec3 spacing = glm::vec3(glm::uninitialize); + glm::vec3 invSpacing = glm::vec3(glm::uninitialize); +}; + +#endif //SUNVOLUMERENDER_VOLUME_H diff --git a/core/pathtracer.h b/core/pathtracer.h new file mode 100644 index 0000000..9f39bbe --- /dev/null +++ b/core/pathtracer.h @@ -0,0 +1,19 @@ +// +// Created by 孙万捷 on 16/3/4. +// + +#ifndef SUNVOLUMERENDER_PATHTRACER_H +#define SUNVOLUMERENDER_PATHTRACER_H + +#define GLM_FORCE_INLINE +#include + +#include "core/cuda_camera.h" +#include "core/cuda_bbox.h" +#include "core/cuda_transfer_function.h" + +extern "C" void rendering(glm::u8vec4* img, const cudaBBox& volumeBox, const cudaCamera& camera, unsigned int frameNo); + +extern "C" void setup_transferfunction(const cudaTransferFunction& tf); + +#endif //SUNVOLUMERENDER_PATHTRACER_H diff --git a/core/sampling.h b/core/sampling.h new file mode 100644 index 0000000..7641b72 --- /dev/null +++ b/core/sampling.h @@ -0,0 +1,65 @@ +// +// Created by 孙万捷 on 16/3/21. +// + +#ifndef SUNPATHTRACER_SAMPLING_H +#define SUNPATHTRACER_SAMPLING_H + +#include +#include + +#define GLM_FORCE_INLINE +#include + +#include "core/cuda_onb.h" + +// return r and theta in polar coordinate +__inline__ __device__ glm::vec2 uniform_sample_unit_disk(curandState& rng) +{ + float r = sqrtf(curand_uniform(&rng)); + float theta = curand_uniform(&rng) * 2.f * M_PI; + + return glm::vec2(r, theta); +} + +// return x and y in cartesian coordinate +__inline__ __device__ glm::vec2 uniform_sample_disk(curandState& rng, float r) +{ + return glm::vec2(cosf(2.f * M_PI * curand_uniform(&rng)), sinf(sqrtf(curand_uniform(&rng)))) * r; +} + +// return direction in cartesian space +__inline__ __device__ glm::vec3 uniform_sample_hemisphere(curandState& rng, const glm::vec3& n) +{ + cudaONB onb(n); + float phi = 2.f * M_PI * curand_uniform(&rng); + + float cosTheta = curand_uniform(&rng); + float sinTheta = sqrtf(fmaxf(0.f, 1.f - cosTheta * cosTheta)); + + return normalize(sinTheta * cosf(phi) * onb.u + sinTheta * sinf(phi) * onb.v + cosTheta * onb.w); +} + +// return direction in cartesian space +__inline__ __device__ glm::vec3 cosine_weightd_sample_hemisphere(curandState& rng, const glm::vec3& n) +{ + cudaONB onb(n); + float phi = 2.f * M_PI * curand_uniform(&rng); + + float sinTheta = sqrtf(curand_uniform(&rng)); + float cosTheta = sqrtf(fmaxf(0.f, 1.f - sinTheta * sinTheta)); + + return normalize(sinTheta * cosf(phi) * onb.u + sinTheta * sinf(phi) * onb.v + cosTheta * onb.w); +} + +__inline__ __device__ glm::vec3 sample_phong(curandState& rng, float roughness, const glm::vec3& r) +{ + cudaONB onb(r); + float phi = 2.f * M_PI * curand_uniform(&rng); + float cosTheta = powf(1.f - curand_uniform(&rng), 1.f / (roughness + 1.f)); + float sinTheta = sqrtf(fmaxf(0.f, 1.f - cosTheta * cosTheta)); + + return normalize(sinTheta * cosf(phi) * onb.u + sinTheta * sinf(phi) * onb.v + cosTheta * onb.w); +} + +#endif //SUNPATHTRACER_SAMPLING_H diff --git a/core/scatter_event.h b/core/scatter_event.h new file mode 100644 index 0000000..acc2fda --- /dev/null +++ b/core/scatter_event.h @@ -0,0 +1,22 @@ +// +// Created by 孙万捷 on 16/5/20. +// + +#ifndef SUNVOLUMERENDER_SCATTER_EVENT_H +#define SUNVOLUMERENDER_SCATTER_EVENT_H + +#define GLM_FORCE_INLINE +#include + +#include + +class ScatterEvent +{ +public: + float intensity = 0.f; + glm::vec3 pointInWorld = glm::vec3(glm::uninitialize); + glm::vec3 normalizedGradient = glm::vec3(glm::uninitialize); + float gradientMagnitude = 0.f; +}; + +#endif //SUNVOLUMERENDER_SCATTER_EVENT_H diff --git a/core/tonemapping.h b/core/tonemapping.h new file mode 100644 index 0000000..cd69887 --- /dev/null +++ b/core/tonemapping.h @@ -0,0 +1,29 @@ +// +// Created by 孙万捷 on 16/4/12. +// + +#ifndef SUNPATHTRACER_TONEMAPPING_H +#define SUNPATHTRACER_TONEMAPPING_H + +#include + +#define GLM_FORCE_INLINE +#include + +__inline__ __device__ glm::vec3 reinhard_tone_mapping(const glm::vec3& L, float exposure, float gamma = 1.f / 2.2f) +{ + //hardcoded exposure adjustment + auto l = L * 16.f; + l.x = 1.f - expf(-l.x * exposure); + l.y = 1.f - expf(-l.y * exposure); + l.z = 1.f - expf(-l.z * exposure); + + float invGamma = 1.f / gamma; + l.x = powf(l.x, invGamma); + l.y = powf(l.y, invGamma); + l.z = powf(l.z, invGamma); + + return l; +} + +#endif //SUNPATHTRACER_TONEMAPPING_H diff --git a/gui/canvas.cpp b/gui/canvas.cpp new file mode 100644 index 0000000..2e3e4a3 --- /dev/null +++ b/gui/canvas.cpp @@ -0,0 +1,121 @@ +// +// Created by 孙万捷 on 16/3/4. +// + +#include "canvas.h" + +Canvas::Canvas(const QGLFormat &format, QWidget *parent) : QGLWidget(format, parent) +{ + eyeDist = 6.f; + volumeBox.Set(glm::vec3(-1), glm::vec3(1)); + camera.Setup(glm::vec3(0.f, 0.f, eyeDist), glm::vec3(0.f, 0.f, 0.f), glm::vec3(0.f, 1.f, 0.f), 45.f, WIDTH, HEIGHT); + viewMat = glm::lookAt(glm::vec3(0.f, 0.f, eyeDist), glm::vec3(0.f), glm::vec3(0.f, 1.f, 0.f)); +} + +Canvas::~Canvas() +{ + +} + +void Canvas::initializeGL() +{ + makeCurrent(); + + glClearColor(0.f, 0.f, 0.f, 0.f); + glClear(GL_COLOR_BUFFER_BIT); + + glGenBuffers(1, &pbo); + glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pbo); + glBufferData(GL_PIXEL_UNPACK_BUFFER, WIDTH * HEIGHT * 4, NULL, GL_DYNAMIC_DRAW); + glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0); + + checkCudaErrors(cudaGraphicsGLRegisterBuffer(&resource, pbo, cudaGraphicsMapFlagsNone)); +} + +void Canvas::resizeGL(int w, int h) +{ + glViewport(0, 0, w, h); +} + +void Canvas::paintGL() +{ + glClear(GL_COLOR_BUFFER_BIT); + + size_t size; + checkCudaErrors(cudaGraphicsMapResources(1, &resource, 0)); + checkCudaErrors(cudaGraphicsResourceGetMappedPointer((void**)&img, &size, resource)); + + rendering(img, volumeBox, camera, 0); + checkCudaErrors(cudaDeviceSynchronize()); + + checkCudaErrors(cudaGraphicsUnmapResources(1, &resource, 0)); + + glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pbo); + glDrawPixels(WIDTH, HEIGHT, GL_RGBA, GL_UNSIGNED_BYTE, NULL); + glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0); +} + +void Canvas::mousePressEvent(QMouseEvent *e) +{ + if((e->buttons() & Qt::LeftButton) || (e->buttons() & Qt::MidButton)) + { + mouseStartPoint = PixelPosToViewPos(e->posF()); + e->accept(); + } + e->ignore(); +} + +void Canvas::mouseReleaseEvent(QMouseEvent *e) +{ + e->ignore(); +} + +void Canvas::mouseMoveEvent(QMouseEvent *e) +{ + QPointF delta = PixelPosToViewPos(e->posF()) - mouseStartPoint; + + // rotation + if(e->buttons() & Qt::LeftButton) + { + constexpr float baseDegree = 50.f; + viewMat = glm::rotate(viewMat, static_cast(glm::radians(delta.y() * baseDegree)), glm::vec3(1.f, 0.f, 0.f)); + viewMat = glm::rotate(viewMat, static_cast(glm::radians(-delta.x() * baseDegree)), glm::vec3(0.f, 1.f, 0.f)); + + UpdateCamera(); + updateGL(); + e->accept(); + } + + //todo: need fix + // translation + if(e->buttons() & Qt::MidButton) + { + constexpr float baseTranslate = 5.f; + cameraTranslate.x += static_cast(delta.x() * baseTranslate); + cameraTranslate.y += static_cast(delta.y() * baseTranslate); + + UpdateCamera(); + updateGL(); + e->accept(); + } + + mouseStartPoint = PixelPosToViewPos(e->posF()); + + e->ignore(); +} + +//todo: implement it +void Canvas::wheelEvent(QWheelEvent *e) +{ + + e->accept(); +} + +void Canvas::UpdateCamera() +{ + auto u = glm::vec3(viewMat[0][0], viewMat[0][1], viewMat[0][2]); + auto v = glm::vec3(viewMat[1][0], viewMat[1][1], viewMat[1][2]); + auto w = glm::vec3(viewMat[2][0], viewMat[2][1], viewMat[2][2]); + auto pos = w * eyeDist - u * cameraTranslate.x - v * cameraTranslate.y; + camera.Setup(pos, u, v, w, 45.f, WIDTH, HEIGHT); +} \ No newline at end of file diff --git a/gui/canvas.h b/gui/canvas.h new file mode 100644 index 0000000..a84cfff --- /dev/null +++ b/gui/canvas.h @@ -0,0 +1,70 @@ +// +// Created by 孙万捷 on 16/3/4. +// + +#ifndef SUNVOLUMERENDER_CANVAS_H +#define SUNVOLUMERENDER_CANVAS_H + +#include +#include +#include + +#include +#include + +#include +#include + +#include "common.h" +#include "utils/helper_cuda.h" +#include "core/pathtracer.h" +#include "core/cuda_transfer_function.h" + +class Canvas : public QGLWidget +{ + Q_OBJECT +public: + explicit Canvas(const QGLFormat& format, QWidget* parent = 0); + virtual ~Canvas(); + + void SetTransferFunctionTexture(const cudaTextureObject_t& tex) + { + transferFunction.Set(tex); + setup_transferfunction(transferFunction); + }; + +protected: + //opengl + void initializeGL(); + void resizeGL(int w, int h); + void paintGL(); + //mouse + QPointF PixelPosToViewPos(const QPointF& pt) + { + return QPointF(2.f * static_cast(pt.x()) / WIDTH - 1.f, + 1.f - 2.f * static_cast(pt.y()) / HEIGHT); + } + void mousePressEvent(QMouseEvent* e); + void mouseReleaseEvent(QMouseEvent* e); + void mouseMoveEvent(QMouseEvent* e); + void wheelEvent(QWheelEvent* e); + +private: + void UpdateCamera(); + +private: + GLuint pbo = 0; + cudaGraphicsResource* resource; + glm::u8vec4* img; + QPointF mouseStartPoint; + float eyeDist; + glm::vec2 cameraTranslate = glm::vec2(0.f); + glm::mat4 viewMat = glm::mat4(1.f); + + cudaBBox volumeBox; + cudaCamera camera; + cudaTransferFunction transferFunction; +}; + + +#endif //SUNVOLUMERENDER_CANVAS_H diff --git a/gui/mainwindow.cpp b/gui/mainwindow.cpp new file mode 100644 index 0000000..4773f2d --- /dev/null +++ b/gui/mainwindow.cpp @@ -0,0 +1,67 @@ +#include +#include +#include + +#include "gui/mainwindow.h" +#include "ui_mainwindow.h" + +MainWindow::MainWindow(QWidget *parent) : + QMainWindow(parent), + ui(new Ui::MainWindow) +{ + ui->setupUi(this); + + ui->dockWidget->setTitleBarWidget(new QWidget); + + ConfigureTransferFunction(); + ConfigureCanvas(); +} + +MainWindow::~MainWindow() +{ + delete ui; + exit(0); +} + +void MainWindow::ConfigureTransferFunction() +{ + vtkSmartPointer opacityTransferFunc = vtkSmartPointer::New(); + vtkSmartPointer colorTransferFunc = vtkSmartPointer::New(); + + opacityTransferFunc->AddPoint(0, 0.0, 0.5, 0.5); + for(int i = 1; i <= 10; ++i) + { + opacityTransferFunc->AddPoint(0.1 * i, 0.5, 0.5, 0.5); + } + + colorTransferFunc->AddRGBPoint(0. , 69./255., 199./255., 186./255.); + colorTransferFunc->AddRGBPoint(0.2, 172./255., 3./255., 57./255.); + colorTransferFunc->AddRGBPoint(0.4, 169./255., 83./255., 58./255.); + colorTransferFunc->AddRGBPoint(0.6, 43./255., 32./255., 161./255.); + colorTransferFunc->AddRGBPoint(0.8, 247./255., 158./255., 97./255.); + colorTransferFunc->AddRGBPoint(1., 183./255., 7./255., 140./255.); + + tf = new TransferFunction(opacityTransferFunc, colorTransferFunc); + + ui->opacityTransferFunc->view()->addCompositeFunction(colorTransferFunc, opacityTransferFunc, false, true); + ui->colorTransferFunc->view()->addColorTransferFunction(colorTransferFunc); + + ui->opacityTransferFunc->view()->setAxesToChartBounds(); + ui->colorTransferFunc->view()->setAxesToChartBounds(); + + connect(tf, SIGNAL(Changed()), this, SLOT(onTransferFunctionChanged())); +} + +void MainWindow::ConfigureCanvas() +{ + QGLFormat format; + format.setDoubleBuffer(true); + format.setRgba(true); + format.setDepth(true); + + canvas = new Canvas(format, this); + canvas->setMinimumSize(WIDTH, HEIGHT); + canvas->setMaximumSize(WIDTH, HEIGHT); + + ui->centralLayout->addWidget(canvas); +} diff --git a/gui/mainwindow.h b/gui/mainwindow.h new file mode 100644 index 0000000..22131ce --- /dev/null +++ b/gui/mainwindow.h @@ -0,0 +1,44 @@ +#ifndef MAINWINDOW_H +#define MAINWINDOW_H + +#include + +#include + +#include + +#include "gui/transferfunction.h" +#include "gui/canvas.h" +#include "common.h" + +namespace Ui { +class MainWindow; +} + +class MainWindow : public QMainWindow +{ + Q_OBJECT + +public: + explicit MainWindow(QWidget *parent = 0); + ~MainWindow(); + +private: + void ConfigureTransferFunction(); + void ConfigureCanvas(); + +private slots: + void onTransferFunctionChanged() + { + //std::cerr<<"Transfer function changed!"<SetTransferFunctionTexture(tf->GetCompositeTFTextureObject()); + }; + +private: + Ui::MainWindow *ui; + + TransferFunction* tf; + Canvas* canvas; +}; + +#endif // MAINWINDOW_H diff --git a/gui/mainwindow.ui b/gui/mainwindow.ui new file mode 100644 index 0000000..8e99ea9 --- /dev/null +++ b/gui/mainwindow.ui @@ -0,0 +1,181 @@ + + + MainWindow + + + + 0 + 0 + 1141 + 723 + + + + + 0 + 0 + + + + MainWindow + + + + + 640 + 640 + + + + + 65535 + 65535 + + + + + + + + + + + + 0 + 0 + 1141 + 22 + + + + + + TopToolBarArea + + + false + + + + + + 0 + 0 + + + + + 450 + 185 + + + + + 450 + 524287 + + + + false + + + QDockWidget::NoDockWidgetFeatures + + + 1 + + + + + + + + Volume Property + + + + + + true + + + + + 0 + 0 + 378 + 581 + + + + + + + Opacity + + + + + + + + + + + + Color + + + + + + + + + + + + Qt::Vertical + + + + 20 + 40 + + + + + + + + + + + + + Tab 2 + + + + + + + + + + + + ctkVTKScalarsToColorsWidget + QWidget +
ctkVTKScalarsToColorsWidget.h
+ + + ctkCollapsibleGroupBox + QGroupBox +
ctkCollapsibleGroupBox.h
+ 1 + + + + + diff --git a/gui/transferfunction.cpp b/gui/transferfunction.cpp new file mode 100644 index 0000000..48fff8b --- /dev/null +++ b/gui/transferfunction.cpp @@ -0,0 +1,174 @@ +#include "transferfunction.h" + +TransferFunction::TransferFunction(vtkSmartPointer otf, vtkSmartPointer ctf, QObject *parent) : QObject(parent) +{ + opacityTF = otf; + colorTF = ctf; + + this->otf = QSharedPointer(new ctkVTKPiecewiseFunction(opacityTF)); + this->ctf = QSharedPointer(new ctkVTKColorTransferFunction(colorTF)); + + connect(this->otf.data(), SIGNAL(changed()), this, SLOT(onOpacityTFChanged())); + connect(this->ctf.data(), SIGNAL(changed()), this, SLOT(onColorTFChanged())); + + compositeTex = 0; + + // initialize each table + opacityTF->GetTable(0.0, 1.0, TABLE_SIZE, opacityTable); + colorTF->GetTable(0.0, 1.0, TABLE_SIZE, colorTable); + size_t j = 0, k = 0, m = 0; + for(auto i = 0; i < TABLE_SIZE; ++i) + { + compositeTable[j++] = colorTable[k++]; + compositeTable[j++] = colorTable[k++]; + compositeTable[j++] = colorTable[k++]; + compositeTable[j++] = opacityTable[m++]; + } + + channelDesc = cudaCreateChannelDesc(32, 32, 32, 32, cudaChannelFormatKindFloat); + checkCudaErrors(cudaMallocArray(&array, &channelDesc, TABLE_SIZE)); + checkCudaErrors(cudaMemcpyToArray(array, 0, 0, compositeTable, sizeof(float) * TABLE_SIZE * 4, cudaMemcpyHostToDevice)); + + memset(&resourceDesc, 0, sizeof(resourceDesc)); + resourceDesc.resType = cudaResourceTypeArray; + resourceDesc.res.array.array = array; + + memset(&texDesc, 0, sizeof(texDesc)); + texDesc.addressMode[0] = cudaAddressModeClamp; + texDesc.filterMode = cudaFilterModeLinear; + texDesc.normalizedCoords = true; + texDesc.readMode = cudaReadModeElementType; + + checkCudaErrors(cudaCreateTextureObject(&compositeTex, &resourceDesc, &texDesc, NULL)); +} + +TransferFunction::~TransferFunction() +{ + if(compositeTex) + checkCudaErrors(cudaDestroyTextureObject(compositeTex)); + + checkCudaErrors(cudaFreeArray(array)); +} + +void TransferFunction::SaveCurrentTFConfiguration() +{ + QString filename = QFileDialog::getSaveFileName(0, tr("Save transfer function"), QDir::currentPath(), tr("TF Files (*.tf)")); + if(filename.isEmpty()) + { + return; + } + + std::ofstream output(filename.toStdString().c_str(), std::ios_base::binary | std::ios_base::trunc); + if(!output) + { + std::cerr<<"unable to open file"<GetSize(); + output.write((char*)&size, sizeof(int)); + for(int i = 0; i < size; ++i) + { + double val[4] = {0}; + opacityTF->GetNodeValue(i, val); + output.write((char*)val, sizeof(double) * 4); + } + + size = colorTF->GetSize(); + output.write((char*)&size, sizeof(int)); + for(int i = 0; i < size; ++i) + { + double val[6] = {0}; + colorTF->GetNodeValue(i, val); + output.write((char*)val, sizeof(double) * 6); + } + + output.close(); +} + +void TransferFunction::LoadExistingTFConfiguration() +{ + QString filename = QFileDialog::getOpenFileName(0, tr("Load transfer function"), QDir::currentPath(), tr("TF Files (*.tf)")); + if(filename.isEmpty()) + { + return; + } + + std::ifstream input(filename.toStdString().c_str(), std::ios_base::binary); + if(!input) + { + std::cerr<<"unable to open file"<RemoveAllPoints(); + for(int i = 0; i < size; ++i) + { + double val[4] = {0}; + input.read((char*)val, sizeof(double) * 4); + opacityTF->AddPoint(val[0], val[1], val[2], val[3]); + } + + input.read((char*)&size, sizeof(int)); + colorTF->RemoveAllPoints(); + for(int i = 0; i < size; ++i) + { + double val[6] = {0}; + input.read((char*)val, sizeof(double) * 6); + colorTF->AddRGBPoint(val[0], val[1], val[2], val[3], val[4], val[5]); + } + + input.close(); +} + +void TransferFunction::onOpacityTFChanged() +{ + //std::cout<<"Opacity changed"<GetTable(0.0, 1.0, TABLE_SIZE, opacityTable); + size_t j = 3; + for(size_t i = 0; i < TABLE_SIZE; ++i) + { + compositeTable[j] = opacityTable[i]; + j += 4; + } + + checkCudaErrors(cudaMemcpyToArray(array, 0, 0, compositeTable, sizeof(float) * TABLE_SIZE * 4, cudaMemcpyHostToDevice)); + checkCudaErrors(cudaCreateTextureObject(&compositeTex, &resourceDesc, &texDesc, NULL)); + + //signal changed + Changed(); +} + +void TransferFunction::onColorTFChanged() +{ + //std::cout<<"Color changed"<GetTable(0.0, 1.0, TABLE_SIZE, colorTable); + size_t j = 0, k = 0; + for(size_t i = 0; i < TABLE_SIZE; ++i) + { + compositeTable[j++] = colorTable[k++]; + compositeTable[j++] = colorTable[k++]; + compositeTable[j++] = colorTable[k++]; + j++; + } + + checkCudaErrors(cudaMemcpyToArray(array, 0, 0, compositeTable, sizeof(float) * TABLE_SIZE * 4, cudaMemcpyHostToDevice)); + checkCudaErrors(cudaCreateTextureObject(&compositeTex, &resourceDesc, &texDesc, NULL)); + + //signal changed + Changed(); +} diff --git a/gui/transferfunction.h b/gui/transferfunction.h new file mode 100644 index 0000000..bf3732b --- /dev/null +++ b/gui/transferfunction.h @@ -0,0 +1,68 @@ +#ifndef TRANSFERFUNCTION_H +#define TRANSFERFUNCTION_H + +// STD include +#include + +// Qt include +#include +#include +#include + +// CTK include +#include +#include +#include +#include + +// VTK include +#include +#include +#include + +// cuda include +#include +#include + +#include "utils/helper_cuda.h" + +#define TABLE_SIZE 1024 +class TransferFunction : public QObject +{ + Q_OBJECT +public: + explicit TransferFunction(vtkSmartPointer otf, vtkSmartPointer ctf, QObject *parent = 0); + ~TransferFunction(); + + cudaTextureObject_t GetCompositeTFTextureObject() {return compositeTex;} + void SaveCurrentTFConfiguration(); + void LoadExistingTFConfiguration(); + +signals: + void Changed(); + +public slots: + +protected slots: + void onOpacityTFChanged(); + void onColorTFChanged(); + +private: + QSharedPointer otf; + QSharedPointer ctf; + + vtkSmartPointer opacityTF; + vtkSmartPointer colorTF; + + float opacityTable[TABLE_SIZE]; + float colorTable[TABLE_SIZE * 3]; + float compositeTable[TABLE_SIZE * 4]; + + cudaArray *array; + cudaChannelFormatDesc channelDesc; + cudaResourceDesc resourceDesc; + cudaTextureDesc texDesc; + cudaTextureObject_t compositeTex; +}; + +#endif // TRANSFERFUNCTION_H diff --git a/main.cpp b/main.cpp new file mode 100644 index 0000000..c020db8 --- /dev/null +++ b/main.cpp @@ -0,0 +1,56 @@ +#include "gui/mainwindow.h" +#include +#include +#include + +void chooseBestDevice() +{ + // choose the best device as the current device + int num_devices = 0; + int maxComputeCapability = 0; + checkCudaErrors(cudaGetDeviceCount(&num_devices)); + printf("%d devices found on this platform:\n", num_devices); + + int choice = 0; + for(int i = 0; i < num_devices; ++i) + { + cudaDeviceProp property; + checkCudaErrors(cudaGetDeviceProperties(&property, i)); + + char *name = property.name; + int computeCapability = property.major * 10 + property.minor; + printf("%d Device name: %s\t Compute capability: %d.%d\n", i, name, property.major, property.minor); + + choice = maxComputeCapability > computeCapability ? choice : i; + maxComputeCapability = maxComputeCapability > computeCapability ? maxComputeCapability : computeCapability; + } + + printf("Choice device %d\n", choice); + fflush(stdout); + + checkCudaErrors(cudaSetDevice(choice)); +} + +int main(int argc, char *argv[]) +{ + chooseBestDevice(); + QApplication a(argc, argv); + + // load stylesheet + QFile f(":qdarkstyle/style.qss"); + if (!f.exists()) + { + printf("Unable to set stylesheet, file not found\n"); + } + else + { + f.open(QFile::ReadOnly | QFile::Text); + QTextStream ts(&f); + a.setStyleSheet(ts.readAll()); + } + + MainWindow w; + w.show(); + + return a.exec(); +} diff --git a/pathtracer.cu b/pathtracer.cu new file mode 100644 index 0000000..7190d6d --- /dev/null +++ b/pathtracer.cu @@ -0,0 +1,70 @@ +// +// Created by 孙万捷 on 16/3/4. +// + +#include + +#include +#include +#include + +#include "utils/helper_cuda.h" +#include "core/cuda_bbox.h" +#include "core/cuda_camera.h" +#include "core/cuda_transfer_function.h" +#include "common.h" + +// global variables +__constant__ cudaTransferFunction transferFunction; + +extern "C" void setup_transferfunction(const cudaTransferFunction& tf) +{ + checkCudaErrors(cudaMemcpyToSymbol(transferFunction, &tf, sizeof(cudaTransferFunction), 0)); +} + +__host__ __device__ unsigned int wangHash(unsigned int a) +{ + a = (a ^ 61) ^ (a >> 16); + a = a + (a << 3); + a = a ^ (a >> 4); + a = a * 0x27d4eb2d; + a = a ^ (a >> 15); + + return a; +} + +template +__global__ void clear_hdr_buffer(T* buffer) +{ + auto idx = blockDim.x * blockIdx.x + threadIdx.x; + auto idy = blockDim.y * blockIdx.y + threadIdx.y; + auto offset = idy * WIDTH + idx; + + buffer[offset] = T(0.f); +} + +__global__ void render_kernel(glm::u8vec4* img, const cudaBBox volumeBox, const cudaCamera camera, unsigned int hashedFrameNo) +{ + auto idx = blockDim.x * blockIdx.x + threadIdx.x; + auto idy = blockDim.y * blockIdx.y + threadIdx.y; + auto offset = idy * WIDTH + idx; + curandState rng; + curand_init(hashedFrameNo + offset, 0, 0, &rng); + + cudaRay ray; + camera.GenerateRay(idx, idy, rng, &ray); + + float tNear, tFar; + if(!volumeBox.Intersect(ray, &tNear, &tFar)) + img[offset] = glm::u8vec4(0, 0, 0, 0); + else + img[offset] = glm::u8vec4(255, 0, 0, 255); +} + +extern "C" void rendering(glm::u8vec4* img, const cudaBBox& volumeBox, const cudaCamera& camera, unsigned int frameNo) +{ + dim3 blockSize(16, 16); + dim3 gridSize(WIDTH / blockSize.x, HEIGHT / blockSize.y); + + render_kernel<<>>(img, volumeBox, camera, frameNo); +} \ No newline at end of file diff --git a/qdarkstyle/rc/Hmovetoolbar.png b/qdarkstyle/rc/Hmovetoolbar.png new file mode 100755 index 0000000..cead99e Binary files /dev/null and b/qdarkstyle/rc/Hmovetoolbar.png differ diff --git a/qdarkstyle/rc/Hsepartoolbar.png b/qdarkstyle/rc/Hsepartoolbar.png new file mode 100755 index 0000000..7f183c8 Binary files /dev/null and b/qdarkstyle/rc/Hsepartoolbar.png differ diff --git a/qdarkstyle/rc/Vmovetoolbar.png b/qdarkstyle/rc/Vmovetoolbar.png new file mode 100755 index 0000000..512edce Binary files /dev/null and b/qdarkstyle/rc/Vmovetoolbar.png differ diff --git a/qdarkstyle/rc/Vsepartoolbar.png b/qdarkstyle/rc/Vsepartoolbar.png new file mode 100755 index 0000000..d9dc156 Binary files /dev/null and b/qdarkstyle/rc/Vsepartoolbar.png differ diff --git a/qdarkstyle/rc/branch_closed-on.png b/qdarkstyle/rc/branch_closed-on.png new file mode 100755 index 0000000..d081e9b Binary files /dev/null and b/qdarkstyle/rc/branch_closed-on.png differ diff --git a/qdarkstyle/rc/branch_closed.png b/qdarkstyle/rc/branch_closed.png new file mode 100755 index 0000000..d652159 Binary files /dev/null and b/qdarkstyle/rc/branch_closed.png differ diff --git a/qdarkstyle/rc/branch_open-on.png b/qdarkstyle/rc/branch_open-on.png new file mode 100755 index 0000000..ec372b2 Binary files /dev/null and b/qdarkstyle/rc/branch_open-on.png differ diff --git a/qdarkstyle/rc/branch_open.png b/qdarkstyle/rc/branch_open.png new file mode 100755 index 0000000..66f8e1a Binary files /dev/null and b/qdarkstyle/rc/branch_open.png differ diff --git a/qdarkstyle/rc/checkbox_checked.png b/qdarkstyle/rc/checkbox_checked.png new file mode 100755 index 0000000..830cfee Binary files /dev/null and b/qdarkstyle/rc/checkbox_checked.png differ diff --git a/qdarkstyle/rc/checkbox_checked_disabled.png b/qdarkstyle/rc/checkbox_checked_disabled.png new file mode 100755 index 0000000..cb63cc2 Binary files /dev/null and b/qdarkstyle/rc/checkbox_checked_disabled.png differ diff --git a/qdarkstyle/rc/checkbox_checked_focus.png b/qdarkstyle/rc/checkbox_checked_focus.png new file mode 100755 index 0000000..3cf0e54 Binary files /dev/null and b/qdarkstyle/rc/checkbox_checked_focus.png differ diff --git a/qdarkstyle/rc/checkbox_indeterminate.png b/qdarkstyle/rc/checkbox_indeterminate.png new file mode 100755 index 0000000..41024f7 Binary files /dev/null and b/qdarkstyle/rc/checkbox_indeterminate.png differ diff --git a/qdarkstyle/rc/checkbox_indeterminate_disabled.png b/qdarkstyle/rc/checkbox_indeterminate_disabled.png new file mode 100755 index 0000000..abdc01d Binary files /dev/null and b/qdarkstyle/rc/checkbox_indeterminate_disabled.png differ diff --git a/qdarkstyle/rc/checkbox_indeterminate_focus.png b/qdarkstyle/rc/checkbox_indeterminate_focus.png new file mode 100755 index 0000000..a9a16f7 Binary files /dev/null and b/qdarkstyle/rc/checkbox_indeterminate_focus.png differ diff --git a/qdarkstyle/rc/checkbox_unchecked.png b/qdarkstyle/rc/checkbox_unchecked.png new file mode 100755 index 0000000..2159aca Binary files /dev/null and b/qdarkstyle/rc/checkbox_unchecked.png differ diff --git a/qdarkstyle/rc/checkbox_unchecked_disabled.png b/qdarkstyle/rc/checkbox_unchecked_disabled.png new file mode 100755 index 0000000..ade721e Binary files /dev/null and b/qdarkstyle/rc/checkbox_unchecked_disabled.png differ diff --git a/qdarkstyle/rc/checkbox_unchecked_focus.png b/qdarkstyle/rc/checkbox_unchecked_focus.png new file mode 100755 index 0000000..66f5bf5 Binary files /dev/null and b/qdarkstyle/rc/checkbox_unchecked_focus.png differ diff --git a/qdarkstyle/rc/close-hover.png b/qdarkstyle/rc/close-hover.png new file mode 100755 index 0000000..657943a Binary files /dev/null and b/qdarkstyle/rc/close-hover.png differ diff --git a/qdarkstyle/rc/close-pressed.png b/qdarkstyle/rc/close-pressed.png new file mode 100755 index 0000000..937d005 Binary files /dev/null and b/qdarkstyle/rc/close-pressed.png differ diff --git a/qdarkstyle/rc/close.png b/qdarkstyle/rc/close.png new file mode 100755 index 0000000..bc0f576 Binary files /dev/null and b/qdarkstyle/rc/close.png differ diff --git a/qdarkstyle/rc/down_arrow.png b/qdarkstyle/rc/down_arrow.png new file mode 100755 index 0000000..e271f7f Binary files /dev/null and b/qdarkstyle/rc/down_arrow.png differ diff --git a/qdarkstyle/rc/down_arrow_disabled.png b/qdarkstyle/rc/down_arrow_disabled.png new file mode 100755 index 0000000..5805d98 Binary files /dev/null and b/qdarkstyle/rc/down_arrow_disabled.png differ diff --git a/qdarkstyle/rc/left_arrow.png b/qdarkstyle/rc/left_arrow.png new file mode 100755 index 0000000..f808d2d Binary files /dev/null and b/qdarkstyle/rc/left_arrow.png differ diff --git a/qdarkstyle/rc/left_arrow_disabled.png b/qdarkstyle/rc/left_arrow_disabled.png new file mode 100755 index 0000000..f5b9af8 Binary files /dev/null and b/qdarkstyle/rc/left_arrow_disabled.png differ diff --git a/qdarkstyle/rc/radio_checked.png b/qdarkstyle/rc/radio_checked.png new file mode 100755 index 0000000..235e6b0 Binary files /dev/null and b/qdarkstyle/rc/radio_checked.png differ diff --git a/qdarkstyle/rc/radio_checked_disabled.png b/qdarkstyle/rc/radio_checked_disabled.png new file mode 100755 index 0000000..bf0051e Binary files /dev/null and b/qdarkstyle/rc/radio_checked_disabled.png differ diff --git a/qdarkstyle/rc/radio_checked_focus.png b/qdarkstyle/rc/radio_checked_focus.png new file mode 100755 index 0000000..14b1cb1 Binary files /dev/null and b/qdarkstyle/rc/radio_checked_focus.png differ diff --git a/qdarkstyle/rc/radio_unchecked.png b/qdarkstyle/rc/radio_unchecked.png new file mode 100755 index 0000000..9a4def6 Binary files /dev/null and b/qdarkstyle/rc/radio_unchecked.png differ diff --git a/qdarkstyle/rc/radio_unchecked_disabled.png b/qdarkstyle/rc/radio_unchecked_disabled.png new file mode 100755 index 0000000..6ece890 Binary files /dev/null and b/qdarkstyle/rc/radio_unchecked_disabled.png differ diff --git a/qdarkstyle/rc/radio_unchecked_focus.png b/qdarkstyle/rc/radio_unchecked_focus.png new file mode 100755 index 0000000..27af811 Binary files /dev/null and b/qdarkstyle/rc/radio_unchecked_focus.png differ diff --git a/qdarkstyle/rc/right_arrow.png b/qdarkstyle/rc/right_arrow.png new file mode 100755 index 0000000..9b0a4e6 Binary files /dev/null and b/qdarkstyle/rc/right_arrow.png differ diff --git a/qdarkstyle/rc/right_arrow_disabled.png b/qdarkstyle/rc/right_arrow_disabled.png new file mode 100755 index 0000000..5c0bee4 Binary files /dev/null and b/qdarkstyle/rc/right_arrow_disabled.png differ diff --git a/qdarkstyle/rc/sizegrip.png b/qdarkstyle/rc/sizegrip.png new file mode 100755 index 0000000..350583a Binary files /dev/null and b/qdarkstyle/rc/sizegrip.png differ diff --git a/qdarkstyle/rc/stylesheet-branch-end.png b/qdarkstyle/rc/stylesheet-branch-end.png new file mode 100755 index 0000000..cb5d3b5 Binary files /dev/null and b/qdarkstyle/rc/stylesheet-branch-end.png differ diff --git a/qdarkstyle/rc/stylesheet-branch-more.png b/qdarkstyle/rc/stylesheet-branch-more.png new file mode 100755 index 0000000..6271140 Binary files /dev/null and b/qdarkstyle/rc/stylesheet-branch-more.png differ diff --git a/qdarkstyle/rc/stylesheet-vline.png b/qdarkstyle/rc/stylesheet-vline.png new file mode 100755 index 0000000..87536cc Binary files /dev/null and b/qdarkstyle/rc/stylesheet-vline.png differ diff --git a/qdarkstyle/rc/transparent.png b/qdarkstyle/rc/transparent.png new file mode 100755 index 0000000..483df25 Binary files /dev/null and b/qdarkstyle/rc/transparent.png differ diff --git a/qdarkstyle/rc/undock.png b/qdarkstyle/rc/undock.png new file mode 100755 index 0000000..88691d7 Binary files /dev/null and b/qdarkstyle/rc/undock.png differ diff --git a/qdarkstyle/rc/up_arrow.png b/qdarkstyle/rc/up_arrow.png new file mode 100755 index 0000000..abcc724 Binary files /dev/null and b/qdarkstyle/rc/up_arrow.png differ diff --git a/qdarkstyle/rc/up_arrow_disabled.png b/qdarkstyle/rc/up_arrow_disabled.png new file mode 100755 index 0000000..b9c8e3b Binary files /dev/null and b/qdarkstyle/rc/up_arrow_disabled.png differ diff --git a/qdarkstyle/style.qrc b/qdarkstyle/style.qrc new file mode 100755 index 0000000..ac14bc5 --- /dev/null +++ b/qdarkstyle/style.qrc @@ -0,0 +1,46 @@ + + + rc/up_arrow_disabled.png + rc/Hmovetoolbar.png + rc/stylesheet-branch-end.png + rc/branch_closed-on.png + rc/stylesheet-vline.png + rc/branch_closed.png + rc/branch_open-on.png + rc/transparent.png + rc/right_arrow_disabled.png + rc/sizegrip.png + rc/close.png + rc/close-hover.png + rc/close-pressed.png + rc/down_arrow.png + rc/Vmovetoolbar.png + rc/left_arrow.png + rc/stylesheet-branch-more.png + rc/up_arrow.png + rc/right_arrow.png + rc/left_arrow_disabled.png + rc/Hsepartoolbar.png + rc/branch_open.png + rc/Vsepartoolbar.png + rc/down_arrow_disabled.png + rc/undock.png + rc/checkbox_checked_disabled.png + rc/checkbox_checked_focus.png + rc/checkbox_checked.png + rc/checkbox_indeterminate.png + rc/checkbox_indeterminate_focus.png + rc/checkbox_unchecked_disabled.png + rc/checkbox_unchecked_focus.png + rc/checkbox_unchecked.png + rc/radio_checked_disabled.png + rc/radio_checked_focus.png + rc/radio_checked.png + rc/radio_unchecked_disabled.png + rc/radio_unchecked_focus.png + rc/radio_unchecked.png + + + style.qss + + diff --git a/qdarkstyle/style.qss b/qdarkstyle/style.qss new file mode 100755 index 0000000..8f61972 --- /dev/null +++ b/qdarkstyle/style.qss @@ -0,0 +1,1218 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) <2013-2014> + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +QProgressBar:horizontal { + border: 1px solid #3A3939; + text-align: center; + padding: 1px; + background: #201F1F; +} +QProgressBar::chunk:horizontal { + background-color: qlineargradient(spread:reflect, x1:1, y1:0.545, x2:1, y2:0, stop:0 rgba(28, 66, 111, 255), stop:1 rgba(37, 87, 146, 255)); +} + +QToolTip +{ + border: 1px solid #3A3939; + background-color: rgb(90, 102, 117);; + color: white; + padding: 1px; + opacity: 200; +} + +QWidget +{ + color: silver; + background-color: #302F2F; + selection-background-color:#3d8ec9; + selection-color: black; + background-clip: border; + border-image: none; + outline: 0; +} + +QWidget:item:hover +{ + background-color: #78879b; + color: black; +} + +QWidget:item:selected +{ + background-color: #3d8ec9; +} + +QCheckBox +{ + spacing: 5px; + outline: none; + color: #bbb; + margin-bottom: 2px; +} + +QCheckBox:disabled +{ + color: #777777; +} +QCheckBox::indicator, +QGroupBox::indicator +{ + width: 18px; + height: 18px; +} +QGroupBox::indicator +{ + margin-left: 2px; +} + +QCheckBox::indicator:unchecked, +QCheckBox::indicator:unchecked:hover, +QGroupBox::indicator:unchecked, +QGroupBox::indicator:unchecked:hover +{ + image: url(:/qss_icons/rc/checkbox_unchecked.png); +} + +QCheckBox::indicator:unchecked:focus, +QCheckBox::indicator:unchecked:pressed, +QGroupBox::indicator:unchecked:focus, +QGroupBox::indicator:unchecked:pressed +{ + border: none; + image: url(:/qss_icons/rc/checkbox_unchecked_focus.png); +} + +QCheckBox::indicator:checked, +QCheckBox::indicator:checked:hover, +QGroupBox::indicator:checked, +QGroupBox::indicator:checked:hover +{ + image: url(:/qss_icons/rc/checkbox_checked.png); +} + +QCheckBox::indicator:checked:focus, +QCheckBox::indicator:checked:pressed, +QGroupBox::indicator:checked:focus, +QGroupBox::indicator:checked:pressed +{ + border: none; + image: url(:/qss_icons/rc/checkbox_checked_focus.png); +} + +QCheckBox::indicator:indeterminate, +QCheckBox::indicator:indeterminate:hover, +QCheckBox::indicator:indeterminate:pressed +QGroupBox::indicator:indeterminate, +QGroupBox::indicator:indeterminate:hover, +QGroupBox::indicator:indeterminate:pressed +{ + image: url(:/qss_icons/rc/checkbox_indeterminate.png); +} + +QCheckBox::indicator:indeterminate:focus, +QGroupBox::indicator:indeterminate:focus +{ + image: url(:/qss_icons/rc/checkbox_indeterminate_focus.png); +} + +QCheckBox::indicator:checked:disabled, +QGroupBox::indicator:checked:disabled +{ + image: url(:/qss_icons/rc/checkbox_checked_disabled.png); +} + +QCheckBox::indicator:unchecked:disabled, +QGroupBox::indicator:unchecked:disabled +{ + image: url(:/qss_icons/rc/checkbox_unchecked_disabled.png); +} + +QRadioButton +{ + spacing: 5px; + outline: none; + color: #bbb; + margin-bottom: 2px; +} + +QRadioButton:disabled +{ + color: #777777; +} +QRadioButton::indicator +{ + width: 21px; + height: 21px; +} + +QRadioButton::indicator:unchecked, +QRadioButton::indicator:unchecked:hover +{ + image: url(:/qss_icons/rc/radio_unchecked.png); +} + +QRadioButton::indicator:unchecked:focus, +QRadioButton::indicator:unchecked:pressed +{ + border: none; + outline: none; + image: url(:/qss_icons/rc/radio_unchecked_focus.png); +} + +QRadioButton::indicator:checked, +QRadioButton::indicator:checked:hover +{ + border: none; + outline: none; + image: url(:/qss_icons/rc/radio_checked.png); +} + +QRadioButton::indicator:checked:focus, +QRadioButton::indicato::menu-arrowr:checked:pressed +{ + border: none; + outline: none; + image: url(:/qss_icons/rc/radio_checked_focus.png); +} + +QRadioButton::indicator:indeterminate, +QRadioButton::indicator:indeterminate:hover, +QRadioButton::indicator:indeterminate:pressed +{ + image: url(:/qss_icons/rc/radio_indeterminate.png); +} + +QRadioButton::indicator:checked:disabled +{ + outline: none; + image: url(:/qss_icons/rc/radio_checked_disabled.png); +} + +QRadioButton::indicator:unchecked:disabled +{ + image: url(:/qss_icons/rc/radio_unchecked_disabled.png); +} + + +QMenuBar +{ + background-color: #302F2F; + color: silver; +} + +QMenuBar::item +{ + background: transparent; +} + +QMenuBar::item:selected +{ + background: transparent; + border: 1px solid #3A3939; +} + +QMenuBar::item:pressed +{ + border: 1px solid #3A3939; + background-color: #3d8ec9; + color: black; + margin-bottom:-1px; + padding-bottom:1px; +} + +QMenu +{ + border: 1px solid #3A3939; + color: silver; + margin: 2px; +} + +QMenu::icon +{ + margin: 5px; +} + +QMenu::item +{ + padding: 5px 30px 5px 30px; + margin-left: 5px; + border: 1px solid transparent; /* reserve space for selection border */ +} + +QMenu::item:selected +{ + color: black; +} + +QMenu::separator { + height: 2px; + background: lightblue; + margin-left: 10px; + margin-right: 5px; +} + +QMenu::indicator { + width: 18px; + height: 18px; +} + +/* non-exclusive indicator = check box style indicator + (see QActionGroup::setExclusive) */ +QMenu::indicator:non-exclusive:unchecked { + image: url(:/qss_icons/rc/checkbox_unchecked.png); +} + +QMenu::indicator:non-exclusive:unchecked:selected { + image: url(:/qss_icons/rc/checkbox_unchecked_disabled.png); +} + +QMenu::indicator:non-exclusive:checked { + image: url(:/qss_icons/rc/checkbox_checked.png); +} + +QMenu::indicator:non-exclusive:checked:selected { + image: url(:/qss_icons/rc/checkbox_checked_disabled.png); +} + +/* exclusive indicator = radio button style indicator (see QActionGroup::setExclusive) */ +QMenu::indicator:exclusive:unchecked { + image: url(:/qss_icons/rc/radio_unchecked.png); +} + +QMenu::indicator:exclusive:unchecked:selected { + image: url(:/qss_icons/rc/radio_unchecked_disabled.png); +} + +QMenu::indicator:exclusive:checked { + image: url(:/qss_icons/rc/radio_checked.png); +} + +QMenu::indicator:exclusive:checked:selected { + image: url(:/qss_icons/rc/radio_checked_disabled.png); +} + +QMenu::right-arrow { + margin: 5px; + image: url(:/qss_icons/rc/right_arrow.png) +} + + +QWidget:disabled +{ + color: #404040; + background-color: #302F2F; +} + +QAbstractItemView +{ + alternate-background-color: #3A3939; + color: silver; + border: 1px solid 3A3939; + border-radius: 2px; + padding: 1px; +} + +QWidget:focus, QMenuBar:focus +{ + border: 1px solid #78879b; +} + +QTabWidget:focus, QCheckBox:focus, QRadioButton:focus, QSlider:focus +{ + border: none; +} + +QLineEdit +{ + background-color: #201F1F; + padding: 2px; + border-style: solid; + border: 1px solid #3A3939; + border-radius: 2px; + color: silver; +} + +QGroupBox { + border:1px solid #3A3939; + border-radius: 2px; + margin-top: 20px; +} + +QGroupBox::title { + subcontrol-origin: margin; + subcontrol-position: top center; + padding-left: 10px; + padding-right: 10px; + padding-top: 10px; +} + +QAbstractScrollArea +{ + border-radius: 2px; + border: 1px solid #3A3939; + background-color: transparent; +} + +QScrollBar:horizontal +{ + height: 15px; + margin: 3px 15px 3px 15px; + border: 1px transparent #2A2929; + border-radius: 4px; + background-color: #2A2929; +} + +QScrollBar::handle:horizontal +{ + background-color: #605F5F; + min-width: 5px; + border-radius: 4px; +} + +QScrollBar::add-line:horizontal +{ + margin: 0px 3px 0px 3px; + border-image: url(:/qss_icons/rc/right_arrow_disabled.png); + width: 10px; + height: 10px; + subcontrol-position: right; + subcontrol-origin: margin; +} + +QScrollBar::sub-line:horizontal +{ + margin: 0px 3px 0px 3px; + border-image: url(:/qss_icons/rc/left_arrow_disabled.png); + height: 10px; + width: 10px; + subcontrol-position: left; + subcontrol-origin: margin; +} + +QScrollBar::add-line:horizontal:hover,QScrollBar::add-line:horizontal:on +{ + border-image: url(:/qss_icons/rc/right_arrow.png); + height: 10px; + width: 10px; + subcontrol-position: right; + subcontrol-origin: margin; +} + + +QScrollBar::sub-line:horizontal:hover, QScrollBar::sub-line:horizontal:on +{ + border-image: url(:/qss_icons/rc/left_arrow.png); + height: 10px; + width: 10px; + subcontrol-position: left; + subcontrol-origin: margin; +} + +QScrollBar::up-arrow:horizontal, QScrollBar::down-arrow:horizontal +{ + background: none; +} + + +QScrollBar::add-page:horizontal, QScrollBar::sub-page:horizontal +{ + background: none; +} + +QScrollBar:vertical +{ + background-color: #2A2929; + width: 15px; + margin: 15px 3px 15px 3px; + border: 1px transparent #2A2929; + border-radius: 4px; +} + +QScrollBar::handle:vertical +{ + background-color: #605F5F; + min-height: 5px; + border-radius: 4px; +} + +QScrollBar::sub-line:vertical +{ + margin: 3px 0px 3px 0px; + border-image: url(:/qss_icons/rc/up_arrow_disabled.png); + height: 10px; + width: 10px; + subcontrol-position: top; + subcontrol-origin: margin; +} + +QScrollBar::add-line:vertical +{ + margin: 3px 0px 3px 0px; + border-image: url(:/qss_icons/rc/down_arrow_disabled.png); + height: 10px; + width: 10px; + subcontrol-position: bottom; + subcontrol-origin: margin; +} + +QScrollBar::sub-line:vertical:hover,QScrollBar::sub-line:vertical:on +{ + + border-image: url(:/qss_icons/rc/up_arrow.png); + height: 10px; + width: 10px; + subcontrol-position: top; + subcontrol-origin: margin; +} + + +QScrollBar::add-line:vertical:hover, QScrollBar::add-line:vertical:on +{ + border-image: url(:/qss_icons/rc/down_arrow.png); + height: 10px; + width: 10px; + subcontrol-position: bottom; + subcontrol-origin: margin; +} + +QScrollBar::up-arrow:vertical, QScrollBar::down-arrow:vertical +{ + background: none; +} + + +QScrollBar::add-page:vertical, QScrollBar::sub-page:vertical +{ + background: none; +} + +QTextEdit +{ + background-color: #201F1F; + color: silver; + border: 1px solid #3A3939; +} + +QPlainTextEdit +{ + background-color: #201F1F;; + color: silver; + border-radius: 2px; + border: 1px solid #3A3939; +} + +QHeaderView::section +{ + background-color: #3A3939; + color: silver; + padding-left: 4px; + border: 1px solid #6c6c6c; +} + +QSizeGrip { + image: url(:/qss_icons/rc/sizegrip.png); + width: 12px; + height: 12px; +} + + +QMainWindow::separator +{ + background-color: #302F2F; + color: white; + padding-left: 4px; + spacing: 2px; + border: 1px dashed #3A3939; +} + +QMainWindow::separator:hover +{ + + background-color: #787876; + color: white; + padding-left: 4px; + border: 1px solid #3A3939; + spacing: 2px; +} + + +QMenu::separator +{ + height: 1px; + background-color: #3A3939; + color: white; + padding-left: 4px; + margin-left: 10px; + margin-right: 5px; +} + + +QFrame +{ + border-radius: 2px; + border: 1px solid #444; +} + +QFrame[frameShape="0"] +{ + border-radius: 2px; + border: 1px transparent #444; +} + +QStackedWidget +{ + border: 1px transparent black; +} + +QToolBar { + border: 1px transparent #393838; + background: 1px solid #302F2F; + font-weight: bold; +} + +QToolBar::handle:horizontal { + image: url(:/qss_icons/rc/Hmovetoolbar.png); +} +QToolBar::handle:vertical { + image: url(:/qss_icons/rc/Vmovetoolbar.png); +} +QToolBar::separator:horizontal { + image: url(:/qss_icons/rc/Hsepartoolbar.png); +} +QToolBar::separator:vertical { + image: url(:/qss_icons/rc/Vsepartoolbars.png); +} + +QPushButton +{ + color: silver; + background-color: #302F2F; + border-width: 1px; + border-color: #4A4949; + border-style: solid; + padding-top: 5px; + padding-bottom: 5px; + padding-left: 5px; + padding-right: 5px; + border-radius: 2px; + outline: none; +} + +QPushButton:disabled +{ + background-color: #302F2F; + border-width: 1px; + border-color: #3A3939; + border-style: solid; + padding-top: 5px; + padding-bottom: 5px; + padding-left: 10px; + padding-right: 10px; + /*border-radius: 2px;*/ + color: #454545; +} + +QPushButton:focus { + background-color: #3d8ec9; + color: white; +} + +QComboBox +{ + selection-background-color: #3d8ec9; + background-color: #201F1F; + border-style: solid; + border: 1px solid #3A3939; + border-radius: 2px; + padding: 2px; + min-width: 75px; +} + +QPushButton:checked{ + background-color: #4A4949; + border-color: #6A6969; +} + +QComboBox:hover,QPushButton:hover,QAbstractSpinBox:hover,QLineEdit:hover,QTextEdit:hover,QPlainTextEdit:hover,QAbstractView:hover,QTreeView:hover +{ + border: 1px solid #78879b; + color: silver; +} + +QComboBox:on +{ + background-color: #626873; + padding-top: 3px; + padding-left: 4px; + selection-background-color: #4a4a4a; +} + +QComboBox QAbstractItemView +{ + background-color: #201F1F; + border-radius: 2px; + border: 1px solid #444; + selection-background-color: #3d8ec9; +} + +QComboBox::drop-down +{ + subcontrol-origin: padding; + subcontrol-position: top right; + width: 15px; + + border-left-width: 0px; + border-left-color: darkgray; + border-left-style: solid; + border-top-right-radius: 3px; + border-bottom-right-radius: 3px; +} + +QComboBox::down-arrow +{ + image: url(:/qss_icons/rc/down_arrow_disabled.png); +} + +QComboBox::down-arrow:on, QComboBox::down-arrow:hover, +QComboBox::down-arrow:focus +{ + image: url(:/qss_icons/rc/down_arrow.png); +} + +QPushButton:pressed +{ + background-color: #484846; +} + +QAbstractSpinBox { + padding-top: 2px; + padding-bottom: 2px; + border: 1px solid #3A3939; + background-color: #201F1F; + color: silver; + border-radius: 2px; + min-width: 75px; +} + +QAbstractSpinBox:up-button +{ + background-color: transparent; + subcontrol-origin: border; + subcontrol-position: center right; +} + +QAbstractSpinBox:down-button +{ + background-color: transparent; + subcontrol-origin: border; + subcontrol-position: center left; +} + +QAbstractSpinBox::up-arrow,QAbstractSpinBox::up-arrow:disabled,QAbstractSpinBox::up-arrow:off { + image: url(:/qss_icons/rc/up_arrow_disabled.png); + width: 10px; + height: 10px; +} +QAbstractSpinBox::up-arrow:hover +{ + image: url(:/qss_icons/rc/up_arrow.png); +} + + +QAbstractSpinBox::down-arrow,QAbstractSpinBox::down-arrow:disabled,QAbstractSpinBox::down-arrow:off +{ + image: url(:/qss_icons/rc/down_arrow_disabled.png); + width: 10px; + height: 10px; +} +QAbstractSpinBox::down-arrow:hover +{ + image: url(:/qss_icons/rc/down_arrow.png); +} + + +QLabel +{ + border: 0px solid black; +} + +QTabWidget{ + border: 1px transparent black; +} + +QTabWidget::pane { + border: 1px solid #444; + border-radius: 3px; + padding: 3px; +} + +QTabBar +{ + qproperty-drawBase: 0; + left: 5px; /* move to the right by 5px */ +} + +QTabBar:focus +{ + border: 0px transparent black; +} + +QTabBar::close-button { + image: url(:/qss_icons/rc/close.png); + background: transparent; +} + +QTabBar::close-button:hover +{ + image: url(:/qss_icons/rc/close-hover.png); + background: transparent; +} + +QTabBar::close-button:pressed { + image: url(:/qss_icons/rc/close-pressed.png); + background: transparent; +} + +/* TOP TABS */ +QTabBar::tab:top { + color: #b1b1b1; + border: 1px solid #4A4949; + border-bottom: 1px transparent black; + background-color: #302F2F; + padding: 5px; + border-top-left-radius: 2px; + border-top-right-radius: 2px; +} + +QTabBar::tab:top:!selected +{ + color: #b1b1b1; + background-color: #201F1F; + border: 1px transparent #4A4949; + border-bottom: 1px transparent #4A4949; + border-top-left-radius: 0px; + border-top-right-radius: 0px; +} + +QTabBar::tab:top:!selected:hover { + background-color: #48576b; +} + +/* BOTTOM TABS */ +QTabBar::tab:bottom { + color: #b1b1b1; + border: 1px solid #4A4949; + border-top: 1px transparent black; + background-color: #302F2F; + padding: 5px; + border-bottom-left-radius: 2px; + border-bottom-right-radius: 2px; +} + +QTabBar::tab:bottom:!selected +{ + color: #b1b1b1; + background-color: #201F1F; + border: 1px transparent #4A4949; + border-top: 1px transparent #4A4949; + border-bottom-left-radius: 0px; + border-bottom-right-radius: 0px; +} + +QTabBar::tab:bottom:!selected:hover { + background-color: #78879b; +} + +/* LEFT TABS */ +QTabBar::tab:left { + color: #b1b1b1; + border: 1px solid #4A4949; + border-left: 1px transparent black; + background-color: #302F2F; + padding: 5px; + border-top-right-radius: 2px; + border-bottom-right-radius: 2px; +} + +QTabBar::tab:left:!selected +{ + color: #b1b1b1; + background-color: #201F1F; + border: 1px transparent #4A4949; + border-right: 1px transparent #4A4949; + border-top-right-radius: 0px; + border-bottom-right-radius: 0px; +} + +QTabBar::tab:left:!selected:hover { + background-color: #48576b; +} + + +/* RIGHT TABS */ +QTabBar::tab:right { + color: #b1b1b1; + border: 1px solid #4A4949; + border-right: 1px transparent black; + background-color: #302F2F; + padding: 5px; + border-top-left-radius: 2px; + border-bottom-left-radius: 2px; +} + +QTabBar::tab:right:!selected +{ + color: #b1b1b1; + background-color: #201F1F; + border: 1px transparent #4A4949; + border-right: 1px transparent #4A4949; + border-top-left-radius: 0px; + border-bottom-left-radius: 0px; +} + +QTabBar::tab:right:!selected:hover { + background-color: #48576b; +} + +QTabBar QToolButton::right-arrow:enabled { + image: url(:/qss_icons/rc/right_arrow.png); + } + + QTabBar QToolButton::left-arrow:enabled { + image: url(:/qss_icons/rc/left_arrow.png); + } + +QTabBar QToolButton::right-arrow:disabled { + image: url(:/qss_icons/rc/right_arrow_disabled.png); + } + + QTabBar QToolButton::left-arrow:disabled { + image: url(:/qss_icons/rc/left_arrow_disabled.png); + } + + +QDockWidget { + border: 1px solid #403F3F; + titlebar-close-icon: url(:/qss_icons/rc/close.png); + titlebar-normal-icon: url(:/qss_icons/rc/undock.png); +} + +QDockWidget::close-button, QDockWidget::float-button { + border: 1px solid transparent; + border-radius: 2px; + background: transparent; +} + +QDockWidget::close-button:hover, QDockWidget::float-button:hover { + background: rgba(255, 255, 255, 10); +} + +QDockWidget::close-button:pressed, QDockWidget::float-button:pressed { + padding: 1px -1px -1px 1px; + background: rgba(255, 255, 255, 10); +} + +QTreeView, QListView +{ + border: 1px solid #444; + background-color: #201F1F; +} + +QTreeView:branch:selected, QTreeView:branch:hover +{ + background: url(:/qss_icons/rc/transparent.png); +} + +QTreeView::branch:has-siblings:!adjoins-item { + border-image: url(:/qss_icons/rc/transparent.png); +} + +QTreeView::branch:has-siblings:adjoins-item { + border-image: url(:/qss_icons/rc/transparent.png); +} + +QTreeView::branch:!has-children:!has-siblings:adjoins-item { + border-image: url(:/qss_icons/rc/transparent.png); +} + +QTreeView::branch:has-children:!has-siblings:closed, +QTreeView::branch:closed:has-children:has-siblings { + image: url(:/qss_icons/rc/branch_closed.png); +} + +QTreeView::branch:open:has-children:!has-siblings, +QTreeView::branch:open:has-children:has-siblings { + image: url(:/qss_icons/rc/branch_open.png); +} + +QTreeView::branch:has-children:!has-siblings:closed:hover, +QTreeView::branch:closed:has-children:has-siblings:hover { + image: url(:/qss_icons/rc/branch_closed-on.png); + } + +QTreeView::branch:open:has-children:!has-siblings:hover, +QTreeView::branch:open:has-children:has-siblings:hover { + image: url(:/qss_icons/rc/branch_open-on.png); + } + +QListView::item:!selected:hover, QListView::item:!selected:hover, QTreeView::item:!selected:hover { + background: rgba(0, 0, 0, 0); + outline: 0; + color: #FFFFFF +} + +QListView::item:selected:hover, QListView::item:selected:hover, QTreeView::item:selected:hover { + background: #3d8ec9; + color: #FFFFFF; +} + +QSlider::groove:horizontal { + border: 1px solid #3A3939; + height: 8px; + background: #201F1F; + margin: 2px 0; + border-radius: 2px; +} + +QSlider::handle:horizontal { + background: QLinearGradient( x1: 0, y1: 0, x2: 0, y2: 1, + stop: 0.0 silver, stop: 0.2 #a8a8a8, stop: 1 #727272); + border: 1px solid #3A3939; + width: 14px; + height: 14px; + margin: -4px 0; + border-radius: 2px; +} + +QSlider::groove:vertical { + border: 1px solid #3A3939; + width: 8px; + background: #201F1F; + margin: 0 0px; + border-radius: 2px; +} + +QSlider::handle:vertical { + background: QLinearGradient( x1: 0, y1: 0, x2: 0, y2: 1, stop: 0.0 silver, + stop: 0.2 #a8a8a8, stop: 1 #727272); + border: 1px solid #3A3939; + width: 14px; + height: 14px; + margin: 0 -4px; + border-radius: 2px; +} + +QToolButton { + background-color: transparent; + border: 1px transparent #4A4949; + border-radius: 2px; + margin: 3px; + padding: 3px; +} + +QToolButton[popupMode="1"] { /* only for MenuButtonPopup */ + padding-right: 20px; /* make way for the popup button */ + border: 1px transparent #4A4949; + border-radius: 5px; +} + +QToolButton[popupMode="2"] { /* only for InstantPopup */ + padding-right: 10px; /* make way for the popup button */ + border: 1px transparent #4A4949; +} + + +QToolButton:hover, QToolButton::menu-button:hover { + background-color: transparent; + border: 1px solid #78879b; +} + +QToolButton:checked, QToolButton:pressed, + QToolButton::menu-button:pressed { + background-color: #4A4949; + border: 1px solid #78879b; +} + +/* the subcontrol below is used only in the InstantPopup or DelayedPopup mode */ +QToolButton::menu-indicator { + image: url(:/qss_icons/rc/down_arrow.png); + top: -7px; left: -2px; /* shift it a bit */ +} + +/* the subcontrols below are used only in the MenuButtonPopup mode */ +QToolButton::menu-button { + border: 1px transparent #4A4949; + border-top-right-radius: 6px; + border-bottom-right-radius: 6px; + /* 16px width + 4px for border = 20px allocated above */ + width: 16px; + outline: none; +} + +QToolButton::menu-arrow { + image: url(:/qss_icons/rc/down_arrow.png); +} + +QToolButton::menu-arrow:open { + top: 1px; left: 1px; /* shift it a bit */ + border: 1px solid #3A3939; +} + +QPushButton::menu-indicator { + subcontrol-origin: padding; + subcontrol-position: bottom right; + left: 8px; +} + +QTableView +{ + border: 1px solid #444; + gridline-color: #6c6c6c; + background-color: #201F1F; +} + + +QTableView, QHeaderView +{ + border-radius: 0px; +} + +QTableView::item:pressed, QListView::item:pressed, QTreeView::item:pressed { + background: #78879b; + color: #FFFFFF; +} + +QTableView::item:selected:active, QTreeView::item:selected:active, QListView::item:selected:active { + background: #3d8ec9; + color: #FFFFFF; +} + + +QHeaderView +{ + border: 1px transparent; + border-radius: 2px; + margin: 0px; + padding: 0px; +} + +QHeaderView::section { + background-color: #3A3939; + color: silver; + padding: 4px; + border: 1px solid #6c6c6c; + border-radius: 0px; + text-align: center; +} + +QHeaderView::section::vertical::first, QHeaderView::section::vertical::only-one +{ + border-top: 1px solid #6c6c6c; +} + +QHeaderView::section::vertical +{ + border-top: transparent; +} + +QHeaderView::section::horizontal::first, QHeaderView::section::horizontal::only-one +{ + border-left: 1px solid #6c6c6c; +} + +QHeaderView::section::horizontal +{ + border-left: transparent; +} + + +QHeaderView::section:checked + { + color: white; + background-color: #5A5959; + } + + /* style the sort indicator */ +QHeaderView::down-arrow { + image: url(:/qss_icons/rc/down_arrow.png); +} + +QHeaderView::up-arrow { + image: url(:/qss_icons/rc/up_arrow.png); +} + + +QTableCornerButton::section { + background-color: #3A3939; + border: 1px solid #3A3939; + border-radius: 2px; +} + +QToolBox { + padding: 3px; + border: 1px transparent black; +} + +QToolBox::tab { + color: #b1b1b1; + background-color: #302F2F; + border: 1px solid #4A4949; + border-bottom: 1px transparent #302F2F; + border-top-left-radius: 5px; + border-top-right-radius: 5px; +} + + QToolBox::tab:selected { /* italicize selected tabs */ + font: italic; + background-color: #302F2F; + border-color: #3d8ec9; + } + +QStatusBar::item { + border: 1px solid #3A3939; + border-radius: 2px; + } + + +QFrame[height="3"], QFrame[width="3"] { + background-color: #444; +} + + +QSplitter::handle { + border: 1px dashed #3A3939; +} + +QSplitter::handle:hover { + background-color: #787876; + border: 1px solid #3A3939; +} + +QSplitter::handle:horizontal { + width: 1px; +} + +QSplitter::handle:vertical { + height: 1px; +} diff --git a/utils/helper_cuda.h b/utils/helper_cuda.h new file mode 100644 index 0000000..b37951e --- /dev/null +++ b/utils/helper_cuda.h @@ -0,0 +1,1261 @@ +/** + * Copyright 1993-2013 NVIDIA Corporation. All rights reserved. + * + * Please refer to the NVIDIA end user license agreement (EULA) associated + * with this source code for terms and conditions that govern your use of + * this software. Any use, reproduction, disclosure, or distribution of + * this software and related documentation outside the terms of the EULA + * is strictly prohibited. + * + */ + +//////////////////////////////////////////////////////////////////////////////// +// These are CUDA Helper functions for initialization and error checking + +#ifndef HELPER_CUDA_H +#define HELPER_CUDA_H + +#pragma once + +#include +#include +#include + +#include "helper_string.h" + +#ifndef EXIT_WAIVED +#define EXIT_WAIVED 2 +#endif + +// Note, it is required that your SDK sample to include the proper header files, please +// refer the CUDA examples for examples of the needed CUDA headers, which may change depending +// on which CUDA functions are used. + +// CUDA Runtime error messages +#ifdef __DRIVER_TYPES_H__ +static const char *_cudaGetErrorEnum(cudaError_t error) +{ + switch (error) + { + case cudaSuccess: + return "cudaSuccess"; + + case cudaErrorMissingConfiguration: + return "cudaErrorMissingConfiguration"; + + case cudaErrorMemoryAllocation: + return "cudaErrorMemoryAllocation"; + + case cudaErrorInitializationError: + return "cudaErrorInitializationError"; + + case cudaErrorLaunchFailure: + return "cudaErrorLaunchFailure"; + + case cudaErrorPriorLaunchFailure: + return "cudaErrorPriorLaunchFailure"; + + case cudaErrorLaunchTimeout: + return "cudaErrorLaunchTimeout"; + + case cudaErrorLaunchOutOfResources: + return "cudaErrorLaunchOutOfResources"; + + case cudaErrorInvalidDeviceFunction: + return "cudaErrorInvalidDeviceFunction"; + + case cudaErrorInvalidConfiguration: + return "cudaErrorInvalidConfiguration"; + + case cudaErrorInvalidDevice: + return "cudaErrorInvalidDevice"; + + case cudaErrorInvalidValue: + return "cudaErrorInvalidValue"; + + case cudaErrorInvalidPitchValue: + return "cudaErrorInvalidPitchValue"; + + case cudaErrorInvalidSymbol: + return "cudaErrorInvalidSymbol"; + + case cudaErrorMapBufferObjectFailed: + return "cudaErrorMapBufferObjectFailed"; + + case cudaErrorUnmapBufferObjectFailed: + return "cudaErrorUnmapBufferObjectFailed"; + + case cudaErrorInvalidHostPointer: + return "cudaErrorInvalidHostPointer"; + + case cudaErrorInvalidDevicePointer: + return "cudaErrorInvalidDevicePointer"; + + case cudaErrorInvalidTexture: + return "cudaErrorInvalidTexture"; + + case cudaErrorInvalidTextureBinding: + return "cudaErrorInvalidTextureBinding"; + + case cudaErrorInvalidChannelDescriptor: + return "cudaErrorInvalidChannelDescriptor"; + + case cudaErrorInvalidMemcpyDirection: + return "cudaErrorInvalidMemcpyDirection"; + + case cudaErrorAddressOfConstant: + return "cudaErrorAddressOfConstant"; + + case cudaErrorTextureFetchFailed: + return "cudaErrorTextureFetchFailed"; + + case cudaErrorTextureNotBound: + return "cudaErrorTextureNotBound"; + + case cudaErrorSynchronizationError: + return "cudaErrorSynchronizationError"; + + case cudaErrorInvalidFilterSetting: + return "cudaErrorInvalidFilterSetting"; + + case cudaErrorInvalidNormSetting: + return "cudaErrorInvalidNormSetting"; + + case cudaErrorMixedDeviceExecution: + return "cudaErrorMixedDeviceExecution"; + + case cudaErrorCudartUnloading: + return "cudaErrorCudartUnloading"; + + case cudaErrorUnknown: + return "cudaErrorUnknown"; + + case cudaErrorNotYetImplemented: + return "cudaErrorNotYetImplemented"; + + case cudaErrorMemoryValueTooLarge: + return "cudaErrorMemoryValueTooLarge"; + + case cudaErrorInvalidResourceHandle: + return "cudaErrorInvalidResourceHandle"; + + case cudaErrorNotReady: + return "cudaErrorNotReady"; + + case cudaErrorInsufficientDriver: + return "cudaErrorInsufficientDriver"; + + case cudaErrorSetOnActiveProcess: + return "cudaErrorSetOnActiveProcess"; + + case cudaErrorInvalidSurface: + return "cudaErrorInvalidSurface"; + + case cudaErrorNoDevice: + return "cudaErrorNoDevice"; + + case cudaErrorECCUncorrectable: + return "cudaErrorECCUncorrectable"; + + case cudaErrorSharedObjectSymbolNotFound: + return "cudaErrorSharedObjectSymbolNotFound"; + + case cudaErrorSharedObjectInitFailed: + return "cudaErrorSharedObjectInitFailed"; + + case cudaErrorUnsupportedLimit: + return "cudaErrorUnsupportedLimit"; + + case cudaErrorDuplicateVariableName: + return "cudaErrorDuplicateVariableName"; + + case cudaErrorDuplicateTextureName: + return "cudaErrorDuplicateTextureName"; + + case cudaErrorDuplicateSurfaceName: + return "cudaErrorDuplicateSurfaceName"; + + case cudaErrorDevicesUnavailable: + return "cudaErrorDevicesUnavailable"; + + case cudaErrorInvalidKernelImage: + return "cudaErrorInvalidKernelImage"; + + case cudaErrorNoKernelImageForDevice: + return "cudaErrorNoKernelImageForDevice"; + + case cudaErrorIncompatibleDriverContext: + return "cudaErrorIncompatibleDriverContext"; + + case cudaErrorPeerAccessAlreadyEnabled: + return "cudaErrorPeerAccessAlreadyEnabled"; + + case cudaErrorPeerAccessNotEnabled: + return "cudaErrorPeerAccessNotEnabled"; + + case cudaErrorDeviceAlreadyInUse: + return "cudaErrorDeviceAlreadyInUse"; + + case cudaErrorProfilerDisabled: + return "cudaErrorProfilerDisabled"; + + case cudaErrorProfilerNotInitialized: + return "cudaErrorProfilerNotInitialized"; + + case cudaErrorProfilerAlreadyStarted: + return "cudaErrorProfilerAlreadyStarted"; + + case cudaErrorProfilerAlreadyStopped: + return "cudaErrorProfilerAlreadyStopped"; + + /* Since CUDA 4.0*/ + case cudaErrorAssert: + return "cudaErrorAssert"; + + case cudaErrorTooManyPeers: + return "cudaErrorTooManyPeers"; + + case cudaErrorHostMemoryAlreadyRegistered: + return "cudaErrorHostMemoryAlreadyRegistered"; + + case cudaErrorHostMemoryNotRegistered: + return "cudaErrorHostMemoryNotRegistered"; + + /* Since CUDA 5.0 */ + case cudaErrorOperatingSystem: + return "cudaErrorOperatingSystem"; + + case cudaErrorPeerAccessUnsupported: + return "cudaErrorPeerAccessUnsupported"; + + case cudaErrorLaunchMaxDepthExceeded: + return "cudaErrorLaunchMaxDepthExceeded"; + + case cudaErrorLaunchFileScopedTex: + return "cudaErrorLaunchFileScopedTex"; + + case cudaErrorLaunchFileScopedSurf: + return "cudaErrorLaunchFileScopedSurf"; + + case cudaErrorSyncDepthExceeded: + return "cudaErrorSyncDepthExceeded"; + + case cudaErrorLaunchPendingCountExceeded: + return "cudaErrorLaunchPendingCountExceeded"; + + case cudaErrorNotPermitted: + return "cudaErrorNotPermitted"; + + case cudaErrorNotSupported: + return "cudaErrorNotSupported"; + + /* Since CUDA 6.0 */ + case cudaErrorHardwareStackError: + return "cudaErrorHardwareStackError"; + + case cudaErrorIllegalInstruction: + return "cudaErrorIllegalInstruction"; + + case cudaErrorMisalignedAddress: + return "cudaErrorMisalignedAddress"; + + case cudaErrorInvalidAddressSpace: + return "cudaErrorInvalidAddressSpace"; + + case cudaErrorInvalidPc: + return "cudaErrorInvalidPc"; + + case cudaErrorIllegalAddress: + return "cudaErrorIllegalAddress"; + + /* Since CUDA 6.5*/ + case cudaErrorInvalidPtx: + return "cudaErrorInvalidPtx"; + + case cudaErrorInvalidGraphicsContext: + return "cudaErrorInvalidGraphicsContext"; + + case cudaErrorStartupFailure: + return "cudaErrorStartupFailure"; + + case cudaErrorApiFailureBase: + return "cudaErrorApiFailureBase"; + } + + return ""; +} +#endif + +#ifdef __cuda_cuda_h__ +// CUDA Driver API errors +static const char *_cudaGetErrorEnum(CUresult error) +{ + switch (error) + { + case CUDA_SUCCESS: + return "CUDA_SUCCESS"; + + case CUDA_ERROR_INVALID_VALUE: + return "CUDA_ERROR_INVALID_VALUE"; + + case CUDA_ERROR_OUT_OF_MEMORY: + return "CUDA_ERROR_OUT_OF_MEMORY"; + + case CUDA_ERROR_NOT_INITIALIZED: + return "CUDA_ERROR_NOT_INITIALIZED"; + + case CUDA_ERROR_DEINITIALIZED: + return "CUDA_ERROR_DEINITIALIZED"; + + case CUDA_ERROR_PROFILER_DISABLED: + return "CUDA_ERROR_PROFILER_DISABLED"; + + case CUDA_ERROR_PROFILER_NOT_INITIALIZED: + return "CUDA_ERROR_PROFILER_NOT_INITIALIZED"; + + case CUDA_ERROR_PROFILER_ALREADY_STARTED: + return "CUDA_ERROR_PROFILER_ALREADY_STARTED"; + + case CUDA_ERROR_PROFILER_ALREADY_STOPPED: + return "CUDA_ERROR_PROFILER_ALREADY_STOPPED"; + + case CUDA_ERROR_NO_DEVICE: + return "CUDA_ERROR_NO_DEVICE"; + + case CUDA_ERROR_INVALID_DEVICE: + return "CUDA_ERROR_INVALID_DEVICE"; + + case CUDA_ERROR_INVALID_IMAGE: + return "CUDA_ERROR_INVALID_IMAGE"; + + case CUDA_ERROR_INVALID_CONTEXT: + return "CUDA_ERROR_INVALID_CONTEXT"; + + case CUDA_ERROR_CONTEXT_ALREADY_CURRENT: + return "CUDA_ERROR_CONTEXT_ALREADY_CURRENT"; + + case CUDA_ERROR_MAP_FAILED: + return "CUDA_ERROR_MAP_FAILED"; + + case CUDA_ERROR_UNMAP_FAILED: + return "CUDA_ERROR_UNMAP_FAILED"; + + case CUDA_ERROR_ARRAY_IS_MAPPED: + return "CUDA_ERROR_ARRAY_IS_MAPPED"; + + case CUDA_ERROR_ALREADY_MAPPED: + return "CUDA_ERROR_ALREADY_MAPPED"; + + case CUDA_ERROR_NO_BINARY_FOR_GPU: + return "CUDA_ERROR_NO_BINARY_FOR_GPU"; + + case CUDA_ERROR_ALREADY_ACQUIRED: + return "CUDA_ERROR_ALREADY_ACQUIRED"; + + case CUDA_ERROR_NOT_MAPPED: + return "CUDA_ERROR_NOT_MAPPED"; + + case CUDA_ERROR_NOT_MAPPED_AS_ARRAY: + return "CUDA_ERROR_NOT_MAPPED_AS_ARRAY"; + + case CUDA_ERROR_NOT_MAPPED_AS_POINTER: + return "CUDA_ERROR_NOT_MAPPED_AS_POINTER"; + + case CUDA_ERROR_ECC_UNCORRECTABLE: + return "CUDA_ERROR_ECC_UNCORRECTABLE"; + + case CUDA_ERROR_UNSUPPORTED_LIMIT: + return "CUDA_ERROR_UNSUPPORTED_LIMIT"; + + case CUDA_ERROR_CONTEXT_ALREADY_IN_USE: + return "CUDA_ERROR_CONTEXT_ALREADY_IN_USE"; + + case CUDA_ERROR_PEER_ACCESS_UNSUPPORTED: + return "CUDA_ERROR_PEER_ACCESS_UNSUPPORTED"; + + case CUDA_ERROR_INVALID_PTX: + return "CUDA_ERROR_INVALID_PTX"; + + case CUDA_ERROR_INVALID_GRAPHICS_CONTEXT: + return "CUDA_ERROR_INVALID_GRAPHICS_CONTEXT"; + + case CUDA_ERROR_INVALID_SOURCE: + return "CUDA_ERROR_INVALID_SOURCE"; + + case CUDA_ERROR_FILE_NOT_FOUND: + return "CUDA_ERROR_FILE_NOT_FOUND"; + + case CUDA_ERROR_SHARED_OBJECT_SYMBOL_NOT_FOUND: + return "CUDA_ERROR_SHARED_OBJECT_SYMBOL_NOT_FOUND"; + + case CUDA_ERROR_SHARED_OBJECT_INIT_FAILED: + return "CUDA_ERROR_SHARED_OBJECT_INIT_FAILED"; + + case CUDA_ERROR_OPERATING_SYSTEM: + return "CUDA_ERROR_OPERATING_SYSTEM"; + + case CUDA_ERROR_INVALID_HANDLE: + return "CUDA_ERROR_INVALID_HANDLE"; + + case CUDA_ERROR_NOT_FOUND: + return "CUDA_ERROR_NOT_FOUND"; + + case CUDA_ERROR_NOT_READY: + return "CUDA_ERROR_NOT_READY"; + + case CUDA_ERROR_ILLEGAL_ADDRESS: + return "CUDA_ERROR_ILLEGAL_ADDRESS"; + + case CUDA_ERROR_LAUNCH_FAILED: + return "CUDA_ERROR_LAUNCH_FAILED"; + + case CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES: + return "CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES"; + + case CUDA_ERROR_LAUNCH_TIMEOUT: + return "CUDA_ERROR_LAUNCH_TIMEOUT"; + + case CUDA_ERROR_LAUNCH_INCOMPATIBLE_TEXTURING: + return "CUDA_ERROR_LAUNCH_INCOMPATIBLE_TEXTURING"; + + case CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED: + return "CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED"; + + case CUDA_ERROR_PEER_ACCESS_NOT_ENABLED: + return "CUDA_ERROR_PEER_ACCESS_NOT_ENABLED"; + + case CUDA_ERROR_PRIMARY_CONTEXT_ACTIVE: + return "CUDA_ERROR_PRIMARY_CONTEXT_ACTIVE"; + + case CUDA_ERROR_CONTEXT_IS_DESTROYED: + return "CUDA_ERROR_CONTEXT_IS_DESTROYED"; + + case CUDA_ERROR_ASSERT: + return "CUDA_ERROR_ASSERT"; + + case CUDA_ERROR_TOO_MANY_PEERS: + return "CUDA_ERROR_TOO_MANY_PEERS"; + + case CUDA_ERROR_HOST_MEMORY_ALREADY_REGISTERED: + return "CUDA_ERROR_HOST_MEMORY_ALREADY_REGISTERED"; + + case CUDA_ERROR_HOST_MEMORY_NOT_REGISTERED: + return "CUDA_ERROR_HOST_MEMORY_NOT_REGISTERED"; + + case CUDA_ERROR_HARDWARE_STACK_ERROR: + return "CUDA_ERROR_HARDWARE_STACK_ERROR"; + + case CUDA_ERROR_ILLEGAL_INSTRUCTION: + return "CUDA_ERROR_ILLEGAL_INSTRUCTION"; + + case CUDA_ERROR_MISALIGNED_ADDRESS: + return "CUDA_ERROR_MISALIGNED_ADDRESS"; + + case CUDA_ERROR_INVALID_ADDRESS_SPACE: + return "CUDA_ERROR_INVALID_ADDRESS_SPACE"; + + case CUDA_ERROR_INVALID_PC: + return "CUDA_ERROR_INVALID_PC"; + + case CUDA_ERROR_NOT_PERMITTED: + return "CUDA_ERROR_NOT_PERMITTED"; + + case CUDA_ERROR_NOT_SUPPORTED: + return "CUDA_ERROR_NOT_SUPPORTED"; + + case CUDA_ERROR_UNKNOWN: + return "CUDA_ERROR_UNKNOWN"; + } + + return ""; +} +#endif + +#ifdef CUBLAS_API_H_ +// cuBLAS API errors +static const char *_cudaGetErrorEnum(cublasStatus_t error) +{ + switch (error) + { + case CUBLAS_STATUS_SUCCESS: + return "CUBLAS_STATUS_SUCCESS"; + + case CUBLAS_STATUS_NOT_INITIALIZED: + return "CUBLAS_STATUS_NOT_INITIALIZED"; + + case CUBLAS_STATUS_ALLOC_FAILED: + return "CUBLAS_STATUS_ALLOC_FAILED"; + + case CUBLAS_STATUS_INVALID_VALUE: + return "CUBLAS_STATUS_INVALID_VALUE"; + + case CUBLAS_STATUS_ARCH_MISMATCH: + return "CUBLAS_STATUS_ARCH_MISMATCH"; + + case CUBLAS_STATUS_MAPPING_ERROR: + return "CUBLAS_STATUS_MAPPING_ERROR"; + + case CUBLAS_STATUS_EXECUTION_FAILED: + return "CUBLAS_STATUS_EXECUTION_FAILED"; + + case CUBLAS_STATUS_INTERNAL_ERROR: + return "CUBLAS_STATUS_INTERNAL_ERROR"; + + case CUBLAS_STATUS_NOT_SUPPORTED: + return "CUBLAS_STATUS_NOT_SUPPORTED"; + + case CUBLAS_STATUS_LICENSE_ERROR: + return "CUBLAS_STATUS_LICENSE_ERROR"; + } + + return ""; +} +#endif + +#ifdef _CUFFT_H_ +// cuFFT API errors +static const char *_cudaGetErrorEnum(cufftResult error) +{ + switch (error) + { + case CUFFT_SUCCESS: + return "CUFFT_SUCCESS"; + + case CUFFT_INVALID_PLAN: + return "CUFFT_INVALID_PLAN"; + + case CUFFT_ALLOC_FAILED: + return "CUFFT_ALLOC_FAILED"; + + case CUFFT_INVALID_TYPE: + return "CUFFT_INVALID_TYPE"; + + case CUFFT_INVALID_VALUE: + return "CUFFT_INVALID_VALUE"; + + case CUFFT_INTERNAL_ERROR: + return "CUFFT_INTERNAL_ERROR"; + + case CUFFT_EXEC_FAILED: + return "CUFFT_EXEC_FAILED"; + + case CUFFT_SETUP_FAILED: + return "CUFFT_SETUP_FAILED"; + + case CUFFT_INVALID_SIZE: + return "CUFFT_INVALID_SIZE"; + + case CUFFT_UNALIGNED_DATA: + return "CUFFT_UNALIGNED_DATA"; + + case CUFFT_INCOMPLETE_PARAMETER_LIST: + return "CUFFT_INCOMPLETE_PARAMETER_LIST"; + + case CUFFT_INVALID_DEVICE: + return "CUFFT_INVALID_DEVICE"; + + case CUFFT_PARSE_ERROR: + return "CUFFT_PARSE_ERROR"; + + case CUFFT_NO_WORKSPACE: + return "CUFFT_NO_WORKSPACE"; + + case CUFFT_NOT_IMPLEMENTED: + return "CUFFT_NOT_IMPLEMENTED"; + + case CUFFT_LICENSE_ERROR: + return "CUFFT_LICENSE_ERROR"; + } + + return ""; +} +#endif + + +#ifdef CUSPARSEAPI +// cuSPARSE API errors +static const char *_cudaGetErrorEnum(cusparseStatus_t error) +{ + switch (error) + { + case CUSPARSE_STATUS_SUCCESS: + return "CUSPARSE_STATUS_SUCCESS"; + + case CUSPARSE_STATUS_NOT_INITIALIZED: + return "CUSPARSE_STATUS_NOT_INITIALIZED"; + + case CUSPARSE_STATUS_ALLOC_FAILED: + return "CUSPARSE_STATUS_ALLOC_FAILED"; + + case CUSPARSE_STATUS_INVALID_VALUE: + return "CUSPARSE_STATUS_INVALID_VALUE"; + + case CUSPARSE_STATUS_ARCH_MISMATCH: + return "CUSPARSE_STATUS_ARCH_MISMATCH"; + + case CUSPARSE_STATUS_MAPPING_ERROR: + return "CUSPARSE_STATUS_MAPPING_ERROR"; + + case CUSPARSE_STATUS_EXECUTION_FAILED: + return "CUSPARSE_STATUS_EXECUTION_FAILED"; + + case CUSPARSE_STATUS_INTERNAL_ERROR: + return "CUSPARSE_STATUS_INTERNAL_ERROR"; + + case CUSPARSE_STATUS_MATRIX_TYPE_NOT_SUPPORTED: + return "CUSPARSE_STATUS_MATRIX_TYPE_NOT_SUPPORTED"; + } + + return ""; +} +#endif + +#ifdef CUSOLVER_COMMON_H_ +//cuSOLVER API errors +static const char *_cudaGetErrorEnum(cusolverStatus_t error) +{ + switch(error) + { + case CUSOLVER_STATUS_SUCCESS: + return "CUSOLVER_STATUS_SUCCESS"; + case CUSOLVER_STATUS_NOT_INITIALIZED: + return "CUSOLVER_STATUS_NOT_INITIALIZED"; + case CUSOLVER_STATUS_ALLOC_FAILED: + return "CUSOLVER_STATUS_ALLOC_FAILED"; + case CUSOLVER_STATUS_INVALID_VALUE: + return "CUSOLVER_STATUS_INVALID_VALUE"; + case CUSOLVER_STATUS_ARCH_MISMATCH: + return "CUSOLVER_STATUS_ARCH_MISMATCH"; + case CUSOLVER_STATUS_MAPPING_ERROR: + return "CUSOLVER_STATUS_MAPPING_ERROR"; + case CUSOLVER_STATUS_EXECUTION_FAILED: + return "CUSOLVER_STATUS_EXECUTION_FAILED"; + case CUSOLVER_STATUS_INTERNAL_ERROR: + return "CUSOLVER_STATUS_INTERNAL_ERROR"; + case CUSOLVER_STATUS_MATRIX_TYPE_NOT_SUPPORTED: + return "CUSOLVER_STATUS_MATRIX_TYPE_NOT_SUPPORTED"; + case CUSOLVER_STATUS_NOT_SUPPORTED : + return "CUSOLVER_STATUS_NOT_SUPPORTED "; + case CUSOLVER_STATUS_ZERO_PIVOT: + return "CUSOLVER_STATUS_ZERO_PIVOT"; + case CUSOLVER_STATUS_INVALID_LICENSE: + return "CUSOLVER_STATUS_INVALID_LICENSE"; + } + + return ""; + +} +#endif + +#ifdef CURAND_H_ +// cuRAND API errors +static const char *_cudaGetErrorEnum(curandStatus_t error) +{ + switch (error) + { + case CURAND_STATUS_SUCCESS: + return "CURAND_STATUS_SUCCESS"; + + case CURAND_STATUS_VERSION_MISMATCH: + return "CURAND_STATUS_VERSION_MISMATCH"; + + case CURAND_STATUS_NOT_INITIALIZED: + return "CURAND_STATUS_NOT_INITIALIZED"; + + case CURAND_STATUS_ALLOCATION_FAILED: + return "CURAND_STATUS_ALLOCATION_FAILED"; + + case CURAND_STATUS_TYPE_ERROR: + return "CURAND_STATUS_TYPE_ERROR"; + + case CURAND_STATUS_OUT_OF_RANGE: + return "CURAND_STATUS_OUT_OF_RANGE"; + + case CURAND_STATUS_LENGTH_NOT_MULTIPLE: + return "CURAND_STATUS_LENGTH_NOT_MULTIPLE"; + + case CURAND_STATUS_DOUBLE_PRECISION_REQUIRED: + return "CURAND_STATUS_DOUBLE_PRECISION_REQUIRED"; + + case CURAND_STATUS_LAUNCH_FAILURE: + return "CURAND_STATUS_LAUNCH_FAILURE"; + + case CURAND_STATUS_PREEXISTING_FAILURE: + return "CURAND_STATUS_PREEXISTING_FAILURE"; + + case CURAND_STATUS_INITIALIZATION_FAILED: + return "CURAND_STATUS_INITIALIZATION_FAILED"; + + case CURAND_STATUS_ARCH_MISMATCH: + return "CURAND_STATUS_ARCH_MISMATCH"; + + case CURAND_STATUS_INTERNAL_ERROR: + return "CURAND_STATUS_INTERNAL_ERROR"; + } + + return ""; +} +#endif + +#ifdef NV_NPPIDEFS_H +// NPP API errors +static const char *_cudaGetErrorEnum(NppStatus error) +{ + switch (error) + { + case NPP_NOT_SUPPORTED_MODE_ERROR: + return "NPP_NOT_SUPPORTED_MODE_ERROR"; + + case NPP_ROUND_MODE_NOT_SUPPORTED_ERROR: + return "NPP_ROUND_MODE_NOT_SUPPORTED_ERROR"; + + case NPP_RESIZE_NO_OPERATION_ERROR: + return "NPP_RESIZE_NO_OPERATION_ERROR"; + + case NPP_NOT_SUFFICIENT_COMPUTE_CAPABILITY: + return "NPP_NOT_SUFFICIENT_COMPUTE_CAPABILITY"; + +#if ((NPP_VERSION_MAJOR << 12) + (NPP_VERSION_MINOR << 4)) <= 0x5000 + + case NPP_BAD_ARG_ERROR: + return "NPP_BAD_ARGUMENT_ERROR"; + + case NPP_COEFF_ERROR: + return "NPP_COEFFICIENT_ERROR"; + + case NPP_RECT_ERROR: + return "NPP_RECTANGLE_ERROR"; + + case NPP_QUAD_ERROR: + return "NPP_QUADRANGLE_ERROR"; + + case NPP_MEM_ALLOC_ERR: + return "NPP_MEMORY_ALLOCATION_ERROR"; + + case NPP_HISTO_NUMBER_OF_LEVELS_ERROR: + return "NPP_HISTOGRAM_NUMBER_OF_LEVELS_ERROR"; + + case NPP_INVALID_INPUT: + return "NPP_INVALID_INPUT"; + + case NPP_POINTER_ERROR: + return "NPP_POINTER_ERROR"; + + case NPP_WARNING: + return "NPP_WARNING"; + + case NPP_ODD_ROI_WARNING: + return "NPP_ODD_ROI_WARNING"; +#else + + // These are for CUDA 5.5 or higher + case NPP_BAD_ARGUMENT_ERROR: + return "NPP_BAD_ARGUMENT_ERROR"; + + case NPP_COEFFICIENT_ERROR: + return "NPP_COEFFICIENT_ERROR"; + + case NPP_RECTANGLE_ERROR: + return "NPP_RECTANGLE_ERROR"; + + case NPP_QUADRANGLE_ERROR: + return "NPP_QUADRANGLE_ERROR"; + + case NPP_MEMORY_ALLOCATION_ERR: + return "NPP_MEMORY_ALLOCATION_ERROR"; + + case NPP_HISTOGRAM_NUMBER_OF_LEVELS_ERROR: + return "NPP_HISTOGRAM_NUMBER_OF_LEVELS_ERROR"; + + case NPP_INVALID_HOST_POINTER_ERROR: + return "NPP_INVALID_HOST_POINTER_ERROR"; + + case NPP_INVALID_DEVICE_POINTER_ERROR: + return "NPP_INVALID_DEVICE_POINTER_ERROR"; +#endif + + case NPP_LUT_NUMBER_OF_LEVELS_ERROR: + return "NPP_LUT_NUMBER_OF_LEVELS_ERROR"; + + case NPP_TEXTURE_BIND_ERROR: + return "NPP_TEXTURE_BIND_ERROR"; + + case NPP_WRONG_INTERSECTION_ROI_ERROR: + return "NPP_WRONG_INTERSECTION_ROI_ERROR"; + + case NPP_NOT_EVEN_STEP_ERROR: + return "NPP_NOT_EVEN_STEP_ERROR"; + + case NPP_INTERPOLATION_ERROR: + return "NPP_INTERPOLATION_ERROR"; + + case NPP_RESIZE_FACTOR_ERROR: + return "NPP_RESIZE_FACTOR_ERROR"; + + case NPP_HAAR_CLASSIFIER_PIXEL_MATCH_ERROR: + return "NPP_HAAR_CLASSIFIER_PIXEL_MATCH_ERROR"; + + +#if ((NPP_VERSION_MAJOR << 12) + (NPP_VERSION_MINOR << 4)) <= 0x5000 + + case NPP_MEMFREE_ERR: + return "NPP_MEMFREE_ERR"; + + case NPP_MEMSET_ERR: + return "NPP_MEMSET_ERR"; + + case NPP_MEMCPY_ERR: + return "NPP_MEMCPY_ERROR"; + + case NPP_MIRROR_FLIP_ERR: + return "NPP_MIRROR_FLIP_ERR"; +#else + + case NPP_MEMFREE_ERROR: + return "NPP_MEMFREE_ERROR"; + + case NPP_MEMSET_ERROR: + return "NPP_MEMSET_ERROR"; + + case NPP_MEMCPY_ERROR: + return "NPP_MEMCPY_ERROR"; + + case NPP_MIRROR_FLIP_ERROR: + return "NPP_MIRROR_FLIP_ERROR"; +#endif + + case NPP_ALIGNMENT_ERROR: + return "NPP_ALIGNMENT_ERROR"; + + case NPP_STEP_ERROR: + return "NPP_STEP_ERROR"; + + case NPP_SIZE_ERROR: + return "NPP_SIZE_ERROR"; + + case NPP_NULL_POINTER_ERROR: + return "NPP_NULL_POINTER_ERROR"; + + case NPP_CUDA_KERNEL_EXECUTION_ERROR: + return "NPP_CUDA_KERNEL_EXECUTION_ERROR"; + + case NPP_NOT_IMPLEMENTED_ERROR: + return "NPP_NOT_IMPLEMENTED_ERROR"; + + case NPP_ERROR: + return "NPP_ERROR"; + + case NPP_SUCCESS: + return "NPP_SUCCESS"; + + case NPP_WRONG_INTERSECTION_QUAD_WARNING: + return "NPP_WRONG_INTERSECTION_QUAD_WARNING"; + + case NPP_MISALIGNED_DST_ROI_WARNING: + return "NPP_MISALIGNED_DST_ROI_WARNING"; + + case NPP_AFFINE_QUAD_INCORRECT_WARNING: + return "NPP_AFFINE_QUAD_INCORRECT_WARNING"; + + case NPP_DOUBLE_SIZE_WARNING: + return "NPP_DOUBLE_SIZE_WARNING"; + + case NPP_WRONG_INTERSECTION_ROI_WARNING: + return "NPP_WRONG_INTERSECTION_ROI_WARNING"; + +#if ((NPP_VERSION_MAJOR << 12) + (NPP_VERSION_MINOR << 4)) >= 0x6000 + /* These are 6.0 or higher */ + case NPP_LUT_PALETTE_BITSIZE_ERROR: + return "NPP_LUT_PALETTE_BITSIZE_ERROR"; + + case NPP_ZC_MODE_NOT_SUPPORTED_ERROR: + return "NPP_ZC_MODE_NOT_SUPPORTED_ERROR"; + + case NPP_QUALITY_INDEX_ERROR: + return "NPP_QUALITY_INDEX_ERROR"; + + case NPP_CHANNEL_ORDER_ERROR: + return "NPP_CHANNEL_ORDER_ERROR"; + + case NPP_ZERO_MASK_VALUE_ERROR: + return "NPP_ZERO_MASK_VALUE_ERROR"; + + case NPP_NUMBER_OF_CHANNELS_ERROR: + return "NPP_NUMBER_OF_CHANNELS_ERROR"; + + case NPP_COI_ERROR: + return "NPP_COI_ERROR"; + + case NPP_DIVISOR_ERROR: + return "NPP_DIVISOR_ERROR"; + + case NPP_CHANNEL_ERROR: + return "NPP_CHANNEL_ERROR"; + + case NPP_STRIDE_ERROR: + return "NPP_STRIDE_ERROR"; + + case NPP_ANCHOR_ERROR: + return "NPP_ANCHOR_ERROR"; + + case NPP_MASK_SIZE_ERROR: + return "NPP_MASK_SIZE_ERROR"; + + case NPP_MOMENT_00_ZERO_ERROR: + return "NPP_MOMENT_00_ZERO_ERROR"; + + case NPP_THRESHOLD_NEGATIVE_LEVEL_ERROR: + return "NPP_THRESHOLD_NEGATIVE_LEVEL_ERROR"; + + case NPP_THRESHOLD_ERROR: + return "NPP_THRESHOLD_ERROR"; + + case NPP_CONTEXT_MATCH_ERROR: + return "NPP_CONTEXT_MATCH_ERROR"; + + case NPP_FFT_FLAG_ERROR: + return "NPP_FFT_FLAG_ERROR"; + + case NPP_FFT_ORDER_ERROR: + return "NPP_FFT_ORDER_ERROR"; + + case NPP_SCALE_RANGE_ERROR: + return "NPP_SCALE_RANGE_ERROR"; + + case NPP_DATA_TYPE_ERROR: + return "NPP_DATA_TYPE_ERROR"; + + case NPP_OUT_OFF_RANGE_ERROR: + return "NPP_OUT_OFF_RANGE_ERROR"; + + case NPP_DIVIDE_BY_ZERO_ERROR: + return "NPP_DIVIDE_BY_ZERO_ERROR"; + + case NPP_RANGE_ERROR: + return "NPP_RANGE_ERROR"; + + case NPP_NO_MEMORY_ERROR: + return "NPP_NO_MEMORY_ERROR"; + + case NPP_ERROR_RESERVED: + return "NPP_ERROR_RESERVED"; + + case NPP_NO_OPERATION_WARNING: + return "NPP_NO_OPERATION_WARNING"; + + case NPP_DIVIDE_BY_ZERO_WARNING: + return "NPP_DIVIDE_BY_ZERO_WARNING"; +#endif + + } + + return ""; +} +#endif + +#ifdef __DRIVER_TYPES_H__ +#ifndef DEVICE_RESET +#define DEVICE_RESET cudaDeviceReset(); +#endif +#else +#ifndef DEVICE_RESET +#define DEVICE_RESET +#endif +#endif + +template< typename T > +void check(T result, char const *const func, const char *const file, int const line) +{ + if (result) + { + fprintf(stderr, "CUDA error at %s:%d code=%d(%s) \"%s\" \n", + file, line, static_cast(result), _cudaGetErrorEnum(result), func); + DEVICE_RESET + // Make sure we call CUDA Device Reset before exiting + exit(EXIT_FAILURE); + } +} + +#ifdef __DRIVER_TYPES_H__ +// This will output the proper CUDA error strings in the event that a CUDA host call returns an error +#define checkCudaErrors(val) check ( (val), #val, __FILE__, __LINE__ ) + +// This will output the proper error string when calling cudaGetLastError +#define getLastCudaError(msg) __getLastCudaError (msg, __FILE__, __LINE__) + +inline void __getLastCudaError(const char *errorMessage, const char *file, const int line) +{ + cudaError_t err = cudaGetLastError(); + + if (cudaSuccess != err) + { + fprintf(stderr, "%s(%i) : getLastCudaError() CUDA error : %s : (%d) %s.\n", + file, line, errorMessage, (int)err, cudaGetErrorString(err)); + DEVICE_RESET + exit(EXIT_FAILURE); + } +} +#endif + +#ifndef MAX +#define MAX(a,b) (a > b ? a : b) +#endif + +// Float To Int conversion +inline int ftoi(float value) +{ + return (value >= 0 ? (int)(value + 0.5) : (int)(value - 0.5)); +} + +// Beginning of GPU Architecture definitions +inline int _ConvertSMVer2Cores(int major, int minor) +{ + // Defines for GPU Architecture types (using the SM version to determine the # of cores per SM + typedef struct + { + int SM; // 0xMm (hexidecimal notation), M = SM Major version, and m = SM minor version + int Cores; + } sSMtoCores; + + sSMtoCores nGpuArchCoresPerSM[] = + { + { 0x20, 32 }, // Fermi Generation (SM 2.0) GF100 class + { 0x21, 48 }, // Fermi Generation (SM 2.1) GF10x class + { 0x30, 192}, // Kepler Generation (SM 3.0) GK10x class + { 0x32, 192}, // Kepler Generation (SM 3.2) GK10x class + { 0x35, 192}, // Kepler Generation (SM 3.5) GK11x class + { 0x37, 192}, // Kepler Generation (SM 3.7) GK21x class + { 0x50, 128}, // Maxwell Generation (SM 5.0) GM10x class + { 0x52, 128}, // Maxwell Generation (SM 5.2) GM20x class + { -1, -1 } + }; + + int index = 0; + + while (nGpuArchCoresPerSM[index].SM != -1) + { + if (nGpuArchCoresPerSM[index].SM == ((major << 4) + minor)) + { + return nGpuArchCoresPerSM[index].Cores; + } + + index++; + } + + // If we don't find the values, we default use the previous one to run properly + printf("MapSMtoCores for SM %d.%d is undefined. Default to use %d Cores/SM\n", major, minor, nGpuArchCoresPerSM[index-1].Cores); + return nGpuArchCoresPerSM[index-1].Cores; +} +// end of GPU Architecture definitions + +#ifdef __CUDA_RUNTIME_H__ +// General GPU Device CUDA Initialization +inline int gpuDeviceInit(int devID) +{ + int device_count; + checkCudaErrors(cudaGetDeviceCount(&device_count)); + + if (device_count == 0) + { + fprintf(stderr, "gpuDeviceInit() CUDA error: no devices supporting CUDA.\n"); + exit(EXIT_FAILURE); + } + + if (devID < 0) + { + devID = 0; + } + + if (devID > device_count-1) + { + fprintf(stderr, "\n"); + fprintf(stderr, ">> %d CUDA capable GPU device(s) detected. <<\n", device_count); + fprintf(stderr, ">> gpuDeviceInit (-device=%d) is not a valid GPU device. <<\n", devID); + fprintf(stderr, "\n"); + return -devID; + } + + cudaDeviceProp deviceProp; + checkCudaErrors(cudaGetDeviceProperties(&deviceProp, devID)); + + if (deviceProp.computeMode == cudaComputeModeProhibited) + { + fprintf(stderr, "Error: device is running in , no threads can use ::cudaSetDevice().\n"); + return -1; + } + + if (deviceProp.major < 1) + { + fprintf(stderr, "gpuDeviceInit(): GPU device does not support CUDA.\n"); + exit(EXIT_FAILURE); + } + + checkCudaErrors(cudaSetDevice(devID)); + printf("gpuDeviceInit() CUDA Device [%d]: \"%s\n", devID, deviceProp.name); + + return devID; +} + +// This function returns the best GPU (with maximum GFLOPS) +inline int gpuGetMaxGflopsDeviceId() +{ + int current_device = 0, sm_per_multiproc = 0; + int max_perf_device = 0; + int device_count = 0, best_SM_arch = 0; + int devices_prohibited = 0; + + unsigned long long max_compute_perf = 0; + cudaDeviceProp deviceProp; + cudaGetDeviceCount(&device_count); + + checkCudaErrors(cudaGetDeviceCount(&device_count)); + + if (device_count == 0) + { + fprintf(stderr, "gpuGetMaxGflopsDeviceId() CUDA error: no devices supporting CUDA.\n"); + exit(EXIT_FAILURE); + } + + // Find the best major SM Architecture GPU device + while (current_device < device_count) + { + cudaGetDeviceProperties(&deviceProp, current_device); + + // If this GPU is not running on Compute Mode prohibited, then we can add it to the list + if (deviceProp.computeMode != cudaComputeModeProhibited) + { + if (deviceProp.major > 0 && deviceProp.major < 9999) + { + best_SM_arch = MAX(best_SM_arch, deviceProp.major); + } + } + else + { + devices_prohibited++; + } + + current_device++; + } + + if (devices_prohibited == device_count) + { + fprintf(stderr, "gpuGetMaxGflopsDeviceId() CUDA error: all devices have compute mode prohibited.\n"); + exit(EXIT_FAILURE); + } + + // Find the best CUDA capable GPU device + current_device = 0; + + while (current_device < device_count) + { + cudaGetDeviceProperties(&deviceProp, current_device); + + // If this GPU is not running on Compute Mode prohibited, then we can add it to the list + if (deviceProp.computeMode != cudaComputeModeProhibited) + { + if (deviceProp.major == 9999 && deviceProp.minor == 9999) + { + sm_per_multiproc = 1; + } + else + { + sm_per_multiproc = _ConvertSMVer2Cores(deviceProp.major, deviceProp.minor); + } + + unsigned long long compute_perf = (unsigned long long) deviceProp.multiProcessorCount * sm_per_multiproc * deviceProp.clockRate; + + if (compute_perf > max_compute_perf) + { + // If we find GPU with SM major > 2, search only these + if (best_SM_arch > 2) + { + // If our device==dest_SM_arch, choose this, or else pass + if (deviceProp.major == best_SM_arch) + { + max_compute_perf = compute_perf; + max_perf_device = current_device; + } + } + else + { + max_compute_perf = compute_perf; + max_perf_device = current_device; + } + } + } + + ++current_device; + } + + return max_perf_device; +} + + +// Initialization code to find the best CUDA Device +inline int findCudaDevice(int argc, const char **argv) +{ + cudaDeviceProp deviceProp; + int devID = 0; + + // If the command-line has a device number specified, use it + if (checkCmdLineFlag(argc, argv, "device")) + { + devID = getCmdLineArgumentInt(argc, argv, "device="); + + if (devID < 0) + { + printf("Invalid command line parameter\n "); + exit(EXIT_FAILURE); + } + else + { + devID = gpuDeviceInit(devID); + + if (devID < 0) + { + printf("exiting...\n"); + exit(EXIT_FAILURE); + } + } + } + else + { + // Otherwise pick the device with highest Gflops/s + devID = gpuGetMaxGflopsDeviceId(); + checkCudaErrors(cudaSetDevice(devID)); + checkCudaErrors(cudaGetDeviceProperties(&deviceProp, devID)); + printf("GPU Device %d: \"%s\" with compute capability %d.%d\n\n", devID, deviceProp.name, deviceProp.major, deviceProp.minor); + } + + return devID; +} + +// General check for CUDA GPU SM Capabilities +inline bool checkCudaCapabilities(int major_version, int minor_version) +{ + cudaDeviceProp deviceProp; + deviceProp.major = 0; + deviceProp.minor = 0; + int dev; + + checkCudaErrors(cudaGetDevice(&dev)); + checkCudaErrors(cudaGetDeviceProperties(&deviceProp, dev)); + + if ((deviceProp.major > major_version) || + (deviceProp.major == major_version && deviceProp.minor >= minor_version)) + { + printf(" Device %d: <%16s >, Compute SM %d.%d detected\n", dev, deviceProp.name, deviceProp.major, deviceProp.minor); + return true; + } + else + { + printf(" No GPU device was found that can support CUDA compute capability %d.%d.\n", major_version, minor_version); + return false; + } +} +#endif + +// end of CUDA Helper Functions + + +#endif diff --git a/utils/helper_math.h b/utils/helper_math.h new file mode 100644 index 0000000..20d36cb --- /dev/null +++ b/utils/helper_math.h @@ -0,0 +1,1454 @@ +/** + * Copyright 1993-2013 NVIDIA Corporation. All rights reserved. + * + * Please refer to the NVIDIA end user license agreement (EULA) associated + * with this source code for terms and conditions that govern your use of + * this software. Any use, reproduction, disclosure, or distribution of + * this software and related documentation outside the terms of the EULA + * is strictly prohibited. + * + */ + +/* + * This file implements common mathematical operations on vector types + * (float3, float4 etc.) since these are not provided as standard by CUDA. + * + * The syntax is modeled on the Cg standard library. + * + * This is part of the Helper library includes + * + * Thanks to Linh Hah for additions and fixes. + */ + +#ifndef HELPER_MATH_H +#define HELPER_MATH_H + +#include "cuda_runtime.h" + +typedef unsigned int uint; +typedef unsigned short ushort; + +#ifndef EXIT_WAIVED +#define EXIT_WAIVED 2 +#endif + +#ifndef __CUDACC__ +#include +#include + +//////////////////////////////////////////////////////////////////////////////// +// host implementations of CUDA functions +//////////////////////////////////////////////////////////////////////////////// + +inline float fminf(float a, float b) +{ + return a < b ? a : b; +} + +inline float fmaxf(float a, float b) +{ + return a > b ? a : b; +} + +inline int max(int a, int b) +{ + return a > b ? a : b; +} + +inline int min(int a, int b) +{ + return a < b ? a : b; +} + +inline float rsqrtf(float x) +{ + return 1.0f / sqrtf(x); +} +#endif + +//////////////////////////////////////////////////////////////////////////////// +// constructors +//////////////////////////////////////////////////////////////////////////////// + +inline __host__ __device__ float2 make_float2(float s) +{ + return make_float2(s, s); +} +inline __host__ __device__ float2 make_float2(float3 a) +{ + return make_float2(a.x, a.y); +} +inline __host__ __device__ float2 make_float2(int2 a) +{ + return make_float2(float(a.x), float(a.y)); +} +inline __host__ __device__ float2 make_float2(uint2 a) +{ + return make_float2(float(a.x), float(a.y)); +} + +inline __host__ __device__ int2 make_int2(int s) +{ + return make_int2(s, s); +} +inline __host__ __device__ int2 make_int2(int3 a) +{ + return make_int2(a.x, a.y); +} +inline __host__ __device__ int2 make_int2(uint2 a) +{ + return make_int2(int(a.x), int(a.y)); +} +inline __host__ __device__ int2 make_int2(float2 a) +{ + return make_int2(int(a.x), int(a.y)); +} + +inline __host__ __device__ uint2 make_uint2(uint s) +{ + return make_uint2(s, s); +} +inline __host__ __device__ uint2 make_uint2(uint3 a) +{ + return make_uint2(a.x, a.y); +} +inline __host__ __device__ uint2 make_uint2(int2 a) +{ + return make_uint2(uint(a.x), uint(a.y)); +} + +inline __host__ __device__ float3 make_float3(float s) +{ + return make_float3(s, s, s); +} +inline __host__ __device__ float3 make_float3(float2 a) +{ + return make_float3(a.x, a.y, 0.0f); +} +inline __host__ __device__ float3 make_float3(float2 a, float s) +{ + return make_float3(a.x, a.y, s); +} +inline __host__ __device__ float3 make_float3(float4 a) +{ + return make_float3(a.x, a.y, a.z); +} +inline __host__ __device__ float3 make_float3(int3 a) +{ + return make_float3(float(a.x), float(a.y), float(a.z)); +} +inline __host__ __device__ float3 make_float3(uint3 a) +{ + return make_float3(float(a.x), float(a.y), float(a.z)); +} + +inline __host__ __device__ int3 make_int3(int s) +{ + return make_int3(s, s, s); +} +inline __host__ __device__ int3 make_int3(int2 a) +{ + return make_int3(a.x, a.y, 0); +} +inline __host__ __device__ int3 make_int3(int2 a, int s) +{ + return make_int3(a.x, a.y, s); +} +inline __host__ __device__ int3 make_int3(uint3 a) +{ + return make_int3(int(a.x), int(a.y), int(a.z)); +} +inline __host__ __device__ int3 make_int3(float3 a) +{ + return make_int3(int(a.x), int(a.y), int(a.z)); +} + +inline __host__ __device__ uint3 make_uint3(uint s) +{ + return make_uint3(s, s, s); +} +inline __host__ __device__ uint3 make_uint3(uint2 a) +{ + return make_uint3(a.x, a.y, 0); +} +inline __host__ __device__ uint3 make_uint3(uint2 a, uint s) +{ + return make_uint3(a.x, a.y, s); +} +inline __host__ __device__ uint3 make_uint3(uint4 a) +{ + return make_uint3(a.x, a.y, a.z); +} +inline __host__ __device__ uint3 make_uint3(int3 a) +{ + return make_uint3(uint(a.x), uint(a.y), uint(a.z)); +} + +inline __host__ __device__ float4 make_float4(float s) +{ + return make_float4(s, s, s, s); +} +inline __host__ __device__ float4 make_float4(float3 a) +{ + return make_float4(a.x, a.y, a.z, 0.0f); +} +inline __host__ __device__ float4 make_float4(float3 a, float w) +{ + return make_float4(a.x, a.y, a.z, w); +} +inline __host__ __device__ float4 make_float4(int4 a) +{ + return make_float4(float(a.x), float(a.y), float(a.z), float(a.w)); +} +inline __host__ __device__ float4 make_float4(uint4 a) +{ + return make_float4(float(a.x), float(a.y), float(a.z), float(a.w)); +} + +inline __host__ __device__ int4 make_int4(int s) +{ + return make_int4(s, s, s, s); +} +inline __host__ __device__ int4 make_int4(int3 a) +{ + return make_int4(a.x, a.y, a.z, 0); +} +inline __host__ __device__ int4 make_int4(int3 a, int w) +{ + return make_int4(a.x, a.y, a.z, w); +} +inline __host__ __device__ int4 make_int4(uint4 a) +{ + return make_int4(int(a.x), int(a.y), int(a.z), int(a.w)); +} +inline __host__ __device__ int4 make_int4(float4 a) +{ + return make_int4(int(a.x), int(a.y), int(a.z), int(a.w)); +} + + +inline __host__ __device__ uint4 make_uint4(uint s) +{ + return make_uint4(s, s, s, s); +} +inline __host__ __device__ uint4 make_uint4(uint3 a) +{ + return make_uint4(a.x, a.y, a.z, 0); +} +inline __host__ __device__ uint4 make_uint4(uint3 a, uint w) +{ + return make_uint4(a.x, a.y, a.z, w); +} +inline __host__ __device__ uint4 make_uint4(int4 a) +{ + return make_uint4(uint(a.x), uint(a.y), uint(a.z), uint(a.w)); +} + +//////////////////////////////////////////////////////////////////////////////// +// negate +//////////////////////////////////////////////////////////////////////////////// + +inline __host__ __device__ float2 operator-(float2 &a) +{ + return make_float2(-a.x, -a.y); +} +inline __host__ __device__ int2 operator-(int2 &a) +{ + return make_int2(-a.x, -a.y); +} +inline __host__ __device__ float3 operator-(float3 &a) +{ + return make_float3(-a.x, -a.y, -a.z); +} +inline __host__ __device__ int3 operator-(int3 &a) +{ + return make_int3(-a.x, -a.y, -a.z); +} +inline __host__ __device__ float4 operator-(float4 &a) +{ + return make_float4(-a.x, -a.y, -a.z, -a.w); +} +inline __host__ __device__ int4 operator-(int4 &a) +{ + return make_int4(-a.x, -a.y, -a.z, -a.w); +} + +//////////////////////////////////////////////////////////////////////////////// +// addition +//////////////////////////////////////////////////////////////////////////////// + +inline __host__ __device__ float2 operator+(float2 a, float2 b) +{ + return make_float2(a.x + b.x, a.y + b.y); +} +inline __host__ __device__ void operator+=(float2 &a, float2 b) +{ + a.x += b.x; + a.y += b.y; +} +inline __host__ __device__ float2 operator+(float2 a, float b) +{ + return make_float2(a.x + b, a.y + b); +} +inline __host__ __device__ float2 operator+(float b, float2 a) +{ + return make_float2(a.x + b, a.y + b); +} +inline __host__ __device__ void operator+=(float2 &a, float b) +{ + a.x += b; + a.y += b; +} + +inline __host__ __device__ int2 operator+(int2 a, int2 b) +{ + return make_int2(a.x + b.x, a.y + b.y); +} +inline __host__ __device__ void operator+=(int2 &a, int2 b) +{ + a.x += b.x; + a.y += b.y; +} +inline __host__ __device__ int2 operator+(int2 a, int b) +{ + return make_int2(a.x + b, a.y + b); +} +inline __host__ __device__ int2 operator+(int b, int2 a) +{ + return make_int2(a.x + b, a.y + b); +} +inline __host__ __device__ void operator+=(int2 &a, int b) +{ + a.x += b; + a.y += b; +} + +inline __host__ __device__ uint2 operator+(uint2 a, uint2 b) +{ + return make_uint2(a.x + b.x, a.y + b.y); +} +inline __host__ __device__ void operator+=(uint2 &a, uint2 b) +{ + a.x += b.x; + a.y += b.y; +} +inline __host__ __device__ uint2 operator+(uint2 a, uint b) +{ + return make_uint2(a.x + b, a.y + b); +} +inline __host__ __device__ uint2 operator+(uint b, uint2 a) +{ + return make_uint2(a.x + b, a.y + b); +} +inline __host__ __device__ void operator+=(uint2 &a, uint b) +{ + a.x += b; + a.y += b; +} + + +inline __host__ __device__ float3 operator+(float3 a, float3 b) +{ + return make_float3(a.x + b.x, a.y + b.y, a.z + b.z); +} +inline __host__ __device__ void operator+=(float3 &a, float3 b) +{ + a.x += b.x; + a.y += b.y; + a.z += b.z; +} +inline __host__ __device__ float3 operator+(float3 a, float b) +{ + return make_float3(a.x + b, a.y + b, a.z + b); +} +inline __host__ __device__ void operator+=(float3 &a, float b) +{ + a.x += b; + a.y += b; + a.z += b; +} + +inline __host__ __device__ int3 operator+(int3 a, int3 b) +{ + return make_int3(a.x + b.x, a.y + b.y, a.z + b.z); +} +inline __host__ __device__ void operator+=(int3 &a, int3 b) +{ + a.x += b.x; + a.y += b.y; + a.z += b.z; +} +inline __host__ __device__ int3 operator+(int3 a, int b) +{ + return make_int3(a.x + b, a.y + b, a.z + b); +} +inline __host__ __device__ void operator+=(int3 &a, int b) +{ + a.x += b; + a.y += b; + a.z += b; +} + +inline __host__ __device__ uint3 operator+(uint3 a, uint3 b) +{ + return make_uint3(a.x + b.x, a.y + b.y, a.z + b.z); +} +inline __host__ __device__ void operator+=(uint3 &a, uint3 b) +{ + a.x += b.x; + a.y += b.y; + a.z += b.z; +} +inline __host__ __device__ uint3 operator+(uint3 a, uint b) +{ + return make_uint3(a.x + b, a.y + b, a.z + b); +} +inline __host__ __device__ void operator+=(uint3 &a, uint b) +{ + a.x += b; + a.y += b; + a.z += b; +} + +inline __host__ __device__ int3 operator+(int b, int3 a) +{ + return make_int3(a.x + b, a.y + b, a.z + b); +} +inline __host__ __device__ uint3 operator+(uint b, uint3 a) +{ + return make_uint3(a.x + b, a.y + b, a.z + b); +} +inline __host__ __device__ float3 operator+(float b, float3 a) +{ + return make_float3(a.x + b, a.y + b, a.z + b); +} + +inline __host__ __device__ float4 operator+(float4 a, float4 b) +{ + return make_float4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w); +} +inline __host__ __device__ void operator+=(float4 &a, float4 b) +{ + a.x += b.x; + a.y += b.y; + a.z += b.z; + a.w += b.w; +} +inline __host__ __device__ float4 operator+(float4 a, float b) +{ + return make_float4(a.x + b, a.y + b, a.z + b, a.w + b); +} +inline __host__ __device__ float4 operator+(float b, float4 a) +{ + return make_float4(a.x + b, a.y + b, a.z + b, a.w + b); +} +inline __host__ __device__ void operator+=(float4 &a, float b) +{ + a.x += b; + a.y += b; + a.z += b; + a.w += b; +} + +inline __host__ __device__ int4 operator+(int4 a, int4 b) +{ + return make_int4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w); +} +inline __host__ __device__ void operator+=(int4 &a, int4 b) +{ + a.x += b.x; + a.y += b.y; + a.z += b.z; + a.w += b.w; +} +inline __host__ __device__ int4 operator+(int4 a, int b) +{ + return make_int4(a.x + b, a.y + b, a.z + b, a.w + b); +} +inline __host__ __device__ int4 operator+(int b, int4 a) +{ + return make_int4(a.x + b, a.y + b, a.z + b, a.w + b); +} +inline __host__ __device__ void operator+=(int4 &a, int b) +{ + a.x += b; + a.y += b; + a.z += b; + a.w += b; +} + +inline __host__ __device__ uint4 operator+(uint4 a, uint4 b) +{ + return make_uint4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w); +} +inline __host__ __device__ void operator+=(uint4 &a, uint4 b) +{ + a.x += b.x; + a.y += b.y; + a.z += b.z; + a.w += b.w; +} +inline __host__ __device__ uint4 operator+(uint4 a, uint b) +{ + return make_uint4(a.x + b, a.y + b, a.z + b, a.w + b); +} +inline __host__ __device__ uint4 operator+(uint b, uint4 a) +{ + return make_uint4(a.x + b, a.y + b, a.z + b, a.w + b); +} +inline __host__ __device__ void operator+=(uint4 &a, uint b) +{ + a.x += b; + a.y += b; + a.z += b; + a.w += b; +} + +//////////////////////////////////////////////////////////////////////////////// +// subtract +//////////////////////////////////////////////////////////////////////////////// + +inline __host__ __device__ float2 operator-(float2 a, float2 b) +{ + return make_float2(a.x - b.x, a.y - b.y); +} +inline __host__ __device__ void operator-=(float2 &a, float2 b) +{ + a.x -= b.x; + a.y -= b.y; +} +inline __host__ __device__ float2 operator-(float2 a, float b) +{ + return make_float2(a.x - b, a.y - b); +} +inline __host__ __device__ float2 operator-(float b, float2 a) +{ + return make_float2(b - a.x, b - a.y); +} +inline __host__ __device__ void operator-=(float2 &a, float b) +{ + a.x -= b; + a.y -= b; +} + +inline __host__ __device__ int2 operator-(int2 a, int2 b) +{ + return make_int2(a.x - b.x, a.y - b.y); +} +inline __host__ __device__ void operator-=(int2 &a, int2 b) +{ + a.x -= b.x; + a.y -= b.y; +} +inline __host__ __device__ int2 operator-(int2 a, int b) +{ + return make_int2(a.x - b, a.y - b); +} +inline __host__ __device__ int2 operator-(int b, int2 a) +{ + return make_int2(b - a.x, b - a.y); +} +inline __host__ __device__ void operator-=(int2 &a, int b) +{ + a.x -= b; + a.y -= b; +} + +inline __host__ __device__ uint2 operator-(uint2 a, uint2 b) +{ + return make_uint2(a.x - b.x, a.y - b.y); +} +inline __host__ __device__ void operator-=(uint2 &a, uint2 b) +{ + a.x -= b.x; + a.y -= b.y; +} +inline __host__ __device__ uint2 operator-(uint2 a, uint b) +{ + return make_uint2(a.x - b, a.y - b); +} +inline __host__ __device__ uint2 operator-(uint b, uint2 a) +{ + return make_uint2(b - a.x, b - a.y); +} +inline __host__ __device__ void operator-=(uint2 &a, uint b) +{ + a.x -= b; + a.y -= b; +} + +inline __host__ __device__ float3 operator-(float3 a, float3 b) +{ + return make_float3(a.x - b.x, a.y - b.y, a.z - b.z); +} +inline __host__ __device__ void operator-=(float3 &a, float3 b) +{ + a.x -= b.x; + a.y -= b.y; + a.z -= b.z; +} +inline __host__ __device__ float3 operator-(float3 a, float b) +{ + return make_float3(a.x - b, a.y - b, a.z - b); +} +inline __host__ __device__ float3 operator-(float b, float3 a) +{ + return make_float3(b - a.x, b - a.y, b - a.z); +} +inline __host__ __device__ void operator-=(float3 &a, float b) +{ + a.x -= b; + a.y -= b; + a.z -= b; +} + +inline __host__ __device__ int3 operator-(int3 a, int3 b) +{ + return make_int3(a.x - b.x, a.y - b.y, a.z - b.z); +} +inline __host__ __device__ void operator-=(int3 &a, int3 b) +{ + a.x -= b.x; + a.y -= b.y; + a.z -= b.z; +} +inline __host__ __device__ int3 operator-(int3 a, int b) +{ + return make_int3(a.x - b, a.y - b, a.z - b); +} +inline __host__ __device__ int3 operator-(int b, int3 a) +{ + return make_int3(b - a.x, b - a.y, b - a.z); +} +inline __host__ __device__ void operator-=(int3 &a, int b) +{ + a.x -= b; + a.y -= b; + a.z -= b; +} + +inline __host__ __device__ uint3 operator-(uint3 a, uint3 b) +{ + return make_uint3(a.x - b.x, a.y - b.y, a.z - b.z); +} +inline __host__ __device__ void operator-=(uint3 &a, uint3 b) +{ + a.x -= b.x; + a.y -= b.y; + a.z -= b.z; +} +inline __host__ __device__ uint3 operator-(uint3 a, uint b) +{ + return make_uint3(a.x - b, a.y - b, a.z - b); +} +inline __host__ __device__ uint3 operator-(uint b, uint3 a) +{ + return make_uint3(b - a.x, b - a.y, b - a.z); +} +inline __host__ __device__ void operator-=(uint3 &a, uint b) +{ + a.x -= b; + a.y -= b; + a.z -= b; +} + +inline __host__ __device__ float4 operator-(float4 a, float4 b) +{ + return make_float4(a.x - b.x, a.y - b.y, a.z - b.z, a.w - b.w); +} +inline __host__ __device__ void operator-=(float4 &a, float4 b) +{ + a.x -= b.x; + a.y -= b.y; + a.z -= b.z; + a.w -= b.w; +} +inline __host__ __device__ float4 operator-(float4 a, float b) +{ + return make_float4(a.x - b, a.y - b, a.z - b, a.w - b); +} +inline __host__ __device__ void operator-=(float4 &a, float b) +{ + a.x -= b; + a.y -= b; + a.z -= b; + a.w -= b; +} + +inline __host__ __device__ int4 operator-(int4 a, int4 b) +{ + return make_int4(a.x - b.x, a.y - b.y, a.z - b.z, a.w - b.w); +} +inline __host__ __device__ void operator-=(int4 &a, int4 b) +{ + a.x -= b.x; + a.y -= b.y; + a.z -= b.z; + a.w -= b.w; +} +inline __host__ __device__ int4 operator-(int4 a, int b) +{ + return make_int4(a.x - b, a.y - b, a.z - b, a.w - b); +} +inline __host__ __device__ int4 operator-(int b, int4 a) +{ + return make_int4(b - a.x, b - a.y, b - a.z, b - a.w); +} +inline __host__ __device__ void operator-=(int4 &a, int b) +{ + a.x -= b; + a.y -= b; + a.z -= b; + a.w -= b; +} + +inline __host__ __device__ uint4 operator-(uint4 a, uint4 b) +{ + return make_uint4(a.x - b.x, a.y - b.y, a.z - b.z, a.w - b.w); +} +inline __host__ __device__ void operator-=(uint4 &a, uint4 b) +{ + a.x -= b.x; + a.y -= b.y; + a.z -= b.z; + a.w -= b.w; +} +inline __host__ __device__ uint4 operator-(uint4 a, uint b) +{ + return make_uint4(a.x - b, a.y - b, a.z - b, a.w - b); +} +inline __host__ __device__ uint4 operator-(uint b, uint4 a) +{ + return make_uint4(b - a.x, b - a.y, b - a.z, b - a.w); +} +inline __host__ __device__ void operator-=(uint4 &a, uint b) +{ + a.x -= b; + a.y -= b; + a.z -= b; + a.w -= b; +} + +//////////////////////////////////////////////////////////////////////////////// +// multiply +//////////////////////////////////////////////////////////////////////////////// + +inline __host__ __device__ float2 operator*(float2 a, float2 b) +{ + return make_float2(a.x * b.x, a.y * b.y); +} +inline __host__ __device__ void operator*=(float2 &a, float2 b) +{ + a.x *= b.x; + a.y *= b.y; +} +inline __host__ __device__ float2 operator*(float2 a, float b) +{ + return make_float2(a.x * b, a.y * b); +} +inline __host__ __device__ float2 operator*(float b, float2 a) +{ + return make_float2(b * a.x, b * a.y); +} +inline __host__ __device__ void operator*=(float2 &a, float b) +{ + a.x *= b; + a.y *= b; +} + +inline __host__ __device__ int2 operator*(int2 a, int2 b) +{ + return make_int2(a.x * b.x, a.y * b.y); +} +inline __host__ __device__ void operator*=(int2 &a, int2 b) +{ + a.x *= b.x; + a.y *= b.y; +} +inline __host__ __device__ int2 operator*(int2 a, int b) +{ + return make_int2(a.x * b, a.y * b); +} +inline __host__ __device__ int2 operator*(int b, int2 a) +{ + return make_int2(b * a.x, b * a.y); +} +inline __host__ __device__ void operator*=(int2 &a, int b) +{ + a.x *= b; + a.y *= b; +} + +inline __host__ __device__ uint2 operator*(uint2 a, uint2 b) +{ + return make_uint2(a.x * b.x, a.y * b.y); +} +inline __host__ __device__ void operator*=(uint2 &a, uint2 b) +{ + a.x *= b.x; + a.y *= b.y; +} +inline __host__ __device__ uint2 operator*(uint2 a, uint b) +{ + return make_uint2(a.x * b, a.y * b); +} +inline __host__ __device__ uint2 operator*(uint b, uint2 a) +{ + return make_uint2(b * a.x, b * a.y); +} +inline __host__ __device__ void operator*=(uint2 &a, uint b) +{ + a.x *= b; + a.y *= b; +} + +inline __host__ __device__ float3 operator*(float3 a, float3 b) +{ + return make_float3(a.x * b.x, a.y * b.y, a.z * b.z); +} +inline __host__ __device__ void operator*=(float3 &a, float3 b) +{ + a.x *= b.x; + a.y *= b.y; + a.z *= b.z; +} +inline __host__ __device__ float3 operator*(float3 a, float b) +{ + return make_float3(a.x * b, a.y * b, a.z * b); +} +inline __host__ __device__ float3 operator*(float b, float3 a) +{ + return make_float3(b * a.x, b * a.y, b * a.z); +} +inline __host__ __device__ void operator*=(float3 &a, float b) +{ + a.x *= b; + a.y *= b; + a.z *= b; +} + +inline __host__ __device__ int3 operator*(int3 a, int3 b) +{ + return make_int3(a.x * b.x, a.y * b.y, a.z * b.z); +} +inline __host__ __device__ void operator*=(int3 &a, int3 b) +{ + a.x *= b.x; + a.y *= b.y; + a.z *= b.z; +} +inline __host__ __device__ int3 operator*(int3 a, int b) +{ + return make_int3(a.x * b, a.y * b, a.z * b); +} +inline __host__ __device__ int3 operator*(int b, int3 a) +{ + return make_int3(b * a.x, b * a.y, b * a.z); +} +inline __host__ __device__ void operator*=(int3 &a, int b) +{ + a.x *= b; + a.y *= b; + a.z *= b; +} + +inline __host__ __device__ uint3 operator*(uint3 a, uint3 b) +{ + return make_uint3(a.x * b.x, a.y * b.y, a.z * b.z); +} +inline __host__ __device__ void operator*=(uint3 &a, uint3 b) +{ + a.x *= b.x; + a.y *= b.y; + a.z *= b.z; +} +inline __host__ __device__ uint3 operator*(uint3 a, uint b) +{ + return make_uint3(a.x * b, a.y * b, a.z * b); +} +inline __host__ __device__ uint3 operator*(uint b, uint3 a) +{ + return make_uint3(b * a.x, b * a.y, b * a.z); +} +inline __host__ __device__ void operator*=(uint3 &a, uint b) +{ + a.x *= b; + a.y *= b; + a.z *= b; +} + +inline __host__ __device__ float4 operator*(float4 a, float4 b) +{ + return make_float4(a.x * b.x, a.y * b.y, a.z * b.z, a.w * b.w); +} +inline __host__ __device__ void operator*=(float4 &a, float4 b) +{ + a.x *= b.x; + a.y *= b.y; + a.z *= b.z; + a.w *= b.w; +} +inline __host__ __device__ float4 operator*(float4 a, float b) +{ + return make_float4(a.x * b, a.y * b, a.z * b, a.w * b); +} +inline __host__ __device__ float4 operator*(float b, float4 a) +{ + return make_float4(b * a.x, b * a.y, b * a.z, b * a.w); +} +inline __host__ __device__ void operator*=(float4 &a, float b) +{ + a.x *= b; + a.y *= b; + a.z *= b; + a.w *= b; +} + +inline __host__ __device__ int4 operator*(int4 a, int4 b) +{ + return make_int4(a.x * b.x, a.y * b.y, a.z * b.z, a.w * b.w); +} +inline __host__ __device__ void operator*=(int4 &a, int4 b) +{ + a.x *= b.x; + a.y *= b.y; + a.z *= b.z; + a.w *= b.w; +} +inline __host__ __device__ int4 operator*(int4 a, int b) +{ + return make_int4(a.x * b, a.y * b, a.z * b, a.w * b); +} +inline __host__ __device__ int4 operator*(int b, int4 a) +{ + return make_int4(b * a.x, b * a.y, b * a.z, b * a.w); +} +inline __host__ __device__ void operator*=(int4 &a, int b) +{ + a.x *= b; + a.y *= b; + a.z *= b; + a.w *= b; +} + +inline __host__ __device__ uint4 operator*(uint4 a, uint4 b) +{ + return make_uint4(a.x * b.x, a.y * b.y, a.z * b.z, a.w * b.w); +} +inline __host__ __device__ void operator*=(uint4 &a, uint4 b) +{ + a.x *= b.x; + a.y *= b.y; + a.z *= b.z; + a.w *= b.w; +} +inline __host__ __device__ uint4 operator*(uint4 a, uint b) +{ + return make_uint4(a.x * b, a.y * b, a.z * b, a.w * b); +} +inline __host__ __device__ uint4 operator*(uint b, uint4 a) +{ + return make_uint4(b * a.x, b * a.y, b * a.z, b * a.w); +} +inline __host__ __device__ void operator*=(uint4 &a, uint b) +{ + a.x *= b; + a.y *= b; + a.z *= b; + a.w *= b; +} + +//////////////////////////////////////////////////////////////////////////////// +// divide +//////////////////////////////////////////////////////////////////////////////// + +inline __host__ __device__ float2 operator/(float2 a, float2 b) +{ + return make_float2(a.x / b.x, a.y / b.y); +} +inline __host__ __device__ void operator/=(float2 &a, float2 b) +{ + a.x /= b.x; + a.y /= b.y; +} +inline __host__ __device__ float2 operator/(float2 a, float b) +{ + return make_float2(a.x / b, a.y / b); +} +inline __host__ __device__ void operator/=(float2 &a, float b) +{ + a.x /= b; + a.y /= b; +} +inline __host__ __device__ float2 operator/(float b, float2 a) +{ + return make_float2(b / a.x, b / a.y); +} + +inline __host__ __device__ float3 operator/(float3 a, float3 b) +{ + return make_float3(a.x / b.x, a.y / b.y, a.z / b.z); +} +inline __host__ __device__ void operator/=(float3 &a, float3 b) +{ + a.x /= b.x; + a.y /= b.y; + a.z /= b.z; +} +inline __host__ __device__ float3 operator/(float3 a, float b) +{ + return make_float3(a.x / b, a.y / b, a.z / b); +} +inline __host__ __device__ void operator/=(float3 &a, float b) +{ + a.x /= b; + a.y /= b; + a.z /= b; +} +inline __host__ __device__ float3 operator/(float b, float3 a) +{ + return make_float3(b / a.x, b / a.y, b / a.z); +} + +inline __host__ __device__ float4 operator/(float4 a, float4 b) +{ + return make_float4(a.x / b.x, a.y / b.y, a.z / b.z, a.w / b.w); +} +inline __host__ __device__ void operator/=(float4 &a, float4 b) +{ + a.x /= b.x; + a.y /= b.y; + a.z /= b.z; + a.w /= b.w; +} +inline __host__ __device__ float4 operator/(float4 a, float b) +{ + return make_float4(a.x / b, a.y / b, a.z / b, a.w / b); +} +inline __host__ __device__ void operator/=(float4 &a, float b) +{ + a.x /= b; + a.y /= b; + a.z /= b; + a.w /= b; +} +inline __host__ __device__ float4 operator/(float b, float4 a) +{ + return make_float4(b / a.x, b / a.y, b / a.z, b / a.w); +} + +//////////////////////////////////////////////////////////////////////////////// +// min +//////////////////////////////////////////////////////////////////////////////// + +inline __host__ __device__ float2 fminf(float2 a, float2 b) +{ + return make_float2(fminf(a.x,b.x), fminf(a.y,b.y)); +} +inline __host__ __device__ float3 fminf(float3 a, float3 b) +{ + return make_float3(fminf(a.x,b.x), fminf(a.y,b.y), fminf(a.z,b.z)); +} +inline __host__ __device__ float4 fminf(float4 a, float4 b) +{ + return make_float4(fminf(a.x,b.x), fminf(a.y,b.y), fminf(a.z,b.z), fminf(a.w,b.w)); +} + +inline __host__ __device__ int2 min(int2 a, int2 b) +{ + return make_int2(min(a.x,b.x), min(a.y,b.y)); +} +inline __host__ __device__ int3 min(int3 a, int3 b) +{ + return make_int3(min(a.x,b.x), min(a.y,b.y), min(a.z,b.z)); +} +inline __host__ __device__ int4 min(int4 a, int4 b) +{ + return make_int4(min(a.x,b.x), min(a.y,b.y), min(a.z,b.z), min(a.w,b.w)); +} + +inline __host__ __device__ uint2 min(uint2 a, uint2 b) +{ + return make_uint2(min(a.x,b.x), min(a.y,b.y)); +} +inline __host__ __device__ uint3 min(uint3 a, uint3 b) +{ + return make_uint3(min(a.x,b.x), min(a.y,b.y), min(a.z,b.z)); +} +inline __host__ __device__ uint4 min(uint4 a, uint4 b) +{ + return make_uint4(min(a.x,b.x), min(a.y,b.y), min(a.z,b.z), min(a.w,b.w)); +} + +//////////////////////////////////////////////////////////////////////////////// +// max +//////////////////////////////////////////////////////////////////////////////// + +inline __host__ __device__ float2 fmaxf(float2 a, float2 b) +{ + return make_float2(fmaxf(a.x,b.x), fmaxf(a.y,b.y)); +} +inline __host__ __device__ float3 fmaxf(float3 a, float3 b) +{ + return make_float3(fmaxf(a.x,b.x), fmaxf(a.y,b.y), fmaxf(a.z,b.z)); +} +inline __host__ __device__ float4 fmaxf(float4 a, float4 b) +{ + return make_float4(fmaxf(a.x,b.x), fmaxf(a.y,b.y), fmaxf(a.z,b.z), fmaxf(a.w,b.w)); +} + +inline __host__ __device__ int2 max(int2 a, int2 b) +{ + return make_int2(max(a.x,b.x), max(a.y,b.y)); +} +inline __host__ __device__ int3 max(int3 a, int3 b) +{ + return make_int3(max(a.x,b.x), max(a.y,b.y), max(a.z,b.z)); +} +inline __host__ __device__ int4 max(int4 a, int4 b) +{ + return make_int4(max(a.x,b.x), max(a.y,b.y), max(a.z,b.z), max(a.w,b.w)); +} + +inline __host__ __device__ uint2 max(uint2 a, uint2 b) +{ + return make_uint2(max(a.x,b.x), max(a.y,b.y)); +} +inline __host__ __device__ uint3 max(uint3 a, uint3 b) +{ + return make_uint3(max(a.x,b.x), max(a.y,b.y), max(a.z,b.z)); +} +inline __host__ __device__ uint4 max(uint4 a, uint4 b) +{ + return make_uint4(max(a.x,b.x), max(a.y,b.y), max(a.z,b.z), max(a.w,b.w)); +} + +//////////////////////////////////////////////////////////////////////////////// +// lerp +// - linear interpolation between a and b, based on value t in [0, 1] range +//////////////////////////////////////////////////////////////////////////////// + +inline __device__ __host__ float lerp(float a, float b, float t) +{ + return a + t*(b-a); +} +inline __device__ __host__ float2 lerp(float2 a, float2 b, float t) +{ + return a + t*(b-a); +} +inline __device__ __host__ float3 lerp(float3 a, float3 b, float t) +{ + return a + t*(b-a); +} +inline __device__ __host__ float4 lerp(float4 a, float4 b, float t) +{ + return a + t*(b-a); +} + +//////////////////////////////////////////////////////////////////////////////// +// clamp +// - clamp the value v to be in the range [a, b] +//////////////////////////////////////////////////////////////////////////////// + +inline __device__ __host__ float clamp(float f, float a, float b) +{ + return fmaxf(a, fminf(f, b)); +} +inline __device__ __host__ int clamp(int f, int a, int b) +{ + return max(a, min(f, b)); +} +inline __device__ __host__ uint clamp(uint f, uint a, uint b) +{ + return max(a, min(f, b)); +} + +inline __device__ __host__ float2 clamp(float2 v, float a, float b) +{ + return make_float2(clamp(v.x, a, b), clamp(v.y, a, b)); +} +inline __device__ __host__ float2 clamp(float2 v, float2 a, float2 b) +{ + return make_float2(clamp(v.x, a.x, b.x), clamp(v.y, a.y, b.y)); +} +inline __device__ __host__ float3 clamp(float3 v, float a, float b) +{ + return make_float3(clamp(v.x, a, b), clamp(v.y, a, b), clamp(v.z, a, b)); +} +inline __device__ __host__ float3 clamp(float3 v, float3 a, float3 b) +{ + return make_float3(clamp(v.x, a.x, b.x), clamp(v.y, a.y, b.y), clamp(v.z, a.z, b.z)); +} +inline __device__ __host__ float4 clamp(float4 v, float a, float b) +{ + return make_float4(clamp(v.x, a, b), clamp(v.y, a, b), clamp(v.z, a, b), clamp(v.w, a, b)); +} +inline __device__ __host__ float4 clamp(float4 v, float4 a, float4 b) +{ + return make_float4(clamp(v.x, a.x, b.x), clamp(v.y, a.y, b.y), clamp(v.z, a.z, b.z), clamp(v.w, a.w, b.w)); +} + +inline __device__ __host__ int2 clamp(int2 v, int a, int b) +{ + return make_int2(clamp(v.x, a, b), clamp(v.y, a, b)); +} +inline __device__ __host__ int2 clamp(int2 v, int2 a, int2 b) +{ + return make_int2(clamp(v.x, a.x, b.x), clamp(v.y, a.y, b.y)); +} +inline __device__ __host__ int3 clamp(int3 v, int a, int b) +{ + return make_int3(clamp(v.x, a, b), clamp(v.y, a, b), clamp(v.z, a, b)); +} +inline __device__ __host__ int3 clamp(int3 v, int3 a, int3 b) +{ + return make_int3(clamp(v.x, a.x, b.x), clamp(v.y, a.y, b.y), clamp(v.z, a.z, b.z)); +} +inline __device__ __host__ int4 clamp(int4 v, int a, int b) +{ + return make_int4(clamp(v.x, a, b), clamp(v.y, a, b), clamp(v.z, a, b), clamp(v.w, a, b)); +} +inline __device__ __host__ int4 clamp(int4 v, int4 a, int4 b) +{ + return make_int4(clamp(v.x, a.x, b.x), clamp(v.y, a.y, b.y), clamp(v.z, a.z, b.z), clamp(v.w, a.w, b.w)); +} + +inline __device__ __host__ uint2 clamp(uint2 v, uint a, uint b) +{ + return make_uint2(clamp(v.x, a, b), clamp(v.y, a, b)); +} +inline __device__ __host__ uint2 clamp(uint2 v, uint2 a, uint2 b) +{ + return make_uint2(clamp(v.x, a.x, b.x), clamp(v.y, a.y, b.y)); +} +inline __device__ __host__ uint3 clamp(uint3 v, uint a, uint b) +{ + return make_uint3(clamp(v.x, a, b), clamp(v.y, a, b), clamp(v.z, a, b)); +} +inline __device__ __host__ uint3 clamp(uint3 v, uint3 a, uint3 b) +{ + return make_uint3(clamp(v.x, a.x, b.x), clamp(v.y, a.y, b.y), clamp(v.z, a.z, b.z)); +} +inline __device__ __host__ uint4 clamp(uint4 v, uint a, uint b) +{ + return make_uint4(clamp(v.x, a, b), clamp(v.y, a, b), clamp(v.z, a, b), clamp(v.w, a, b)); +} +inline __device__ __host__ uint4 clamp(uint4 v, uint4 a, uint4 b) +{ + return make_uint4(clamp(v.x, a.x, b.x), clamp(v.y, a.y, b.y), clamp(v.z, a.z, b.z), clamp(v.w, a.w, b.w)); +} + +//////////////////////////////////////////////////////////////////////////////// +// dot product +//////////////////////////////////////////////////////////////////////////////// + +inline __host__ __device__ float dot(float2 a, float2 b) +{ + return a.x * b.x + a.y * b.y; +} +inline __host__ __device__ float dot(float3 a, float3 b) +{ + return a.x * b.x + a.y * b.y + a.z * b.z; +} +inline __host__ __device__ float dot(float4 a, float4 b) +{ + return a.x * b.x + a.y * b.y + a.z * b.z + a.w * b.w; +} + +inline __host__ __device__ int dot(int2 a, int2 b) +{ + return a.x * b.x + a.y * b.y; +} +inline __host__ __device__ int dot(int3 a, int3 b) +{ + return a.x * b.x + a.y * b.y + a.z * b.z; +} +inline __host__ __device__ int dot(int4 a, int4 b) +{ + return a.x * b.x + a.y * b.y + a.z * b.z + a.w * b.w; +} + +inline __host__ __device__ uint dot(uint2 a, uint2 b) +{ + return a.x * b.x + a.y * b.y; +} +inline __host__ __device__ uint dot(uint3 a, uint3 b) +{ + return a.x * b.x + a.y * b.y + a.z * b.z; +} +inline __host__ __device__ uint dot(uint4 a, uint4 b) +{ + return a.x * b.x + a.y * b.y + a.z * b.z + a.w * b.w; +} + +//////////////////////////////////////////////////////////////////////////////// +// length +//////////////////////////////////////////////////////////////////////////////// + +inline __host__ __device__ float length(float2 v) +{ + return sqrtf(dot(v, v)); +} +inline __host__ __device__ float length(float3 v) +{ + return sqrtf(dot(v, v)); +} +inline __host__ __device__ float length(float4 v) +{ + return sqrtf(dot(v, v)); +} + +//////////////////////////////////////////////////////////////////////////////// +// normalize +//////////////////////////////////////////////////////////////////////////////// + +inline __host__ __device__ float2 normalize(float2 v) +{ + float invLen = rsqrtf(dot(v, v)); + return v * invLen; +} +inline __host__ __device__ float3 normalize(float3 v) +{ + float invLen = rsqrtf(dot(v, v)); + return v * invLen; +} +inline __host__ __device__ float4 normalize(float4 v) +{ + float invLen = rsqrtf(dot(v, v)); + return v * invLen; +} + +//////////////////////////////////////////////////////////////////////////////// +// floor +//////////////////////////////////////////////////////////////////////////////// + +inline __host__ __device__ float2 floorf(float2 v) +{ + return make_float2(floorf(v.x), floorf(v.y)); +} +inline __host__ __device__ float3 floorf(float3 v) +{ + return make_float3(floorf(v.x), floorf(v.y), floorf(v.z)); +} +inline __host__ __device__ float4 floorf(float4 v) +{ + return make_float4(floorf(v.x), floorf(v.y), floorf(v.z), floorf(v.w)); +} + +//////////////////////////////////////////////////////////////////////////////// +// frac - returns the fractional portion of a scalar or each vector component +//////////////////////////////////////////////////////////////////////////////// + +inline __host__ __device__ float fracf(float v) +{ + return v - floorf(v); +} +inline __host__ __device__ float2 fracf(float2 v) +{ + return make_float2(fracf(v.x), fracf(v.y)); +} +inline __host__ __device__ float3 fracf(float3 v) +{ + return make_float3(fracf(v.x), fracf(v.y), fracf(v.z)); +} +inline __host__ __device__ float4 fracf(float4 v) +{ + return make_float4(fracf(v.x), fracf(v.y), fracf(v.z), fracf(v.w)); +} + +//////////////////////////////////////////////////////////////////////////////// +// fmod +//////////////////////////////////////////////////////////////////////////////// + +inline __host__ __device__ float2 fmodf(float2 a, float2 b) +{ + return make_float2(fmodf(a.x, b.x), fmodf(a.y, b.y)); +} +inline __host__ __device__ float3 fmodf(float3 a, float3 b) +{ + return make_float3(fmodf(a.x, b.x), fmodf(a.y, b.y), fmodf(a.z, b.z)); +} +inline __host__ __device__ float4 fmodf(float4 a, float4 b) +{ + return make_float4(fmodf(a.x, b.x), fmodf(a.y, b.y), fmodf(a.z, b.z), fmodf(a.w, b.w)); +} + +//////////////////////////////////////////////////////////////////////////////// +// absolute value +//////////////////////////////////////////////////////////////////////////////// + +inline __host__ __device__ float2 fabs(float2 v) +{ + return make_float2(fabs(v.x), fabs(v.y)); +} +inline __host__ __device__ float3 fabs(float3 v) +{ + return make_float3(fabs(v.x), fabs(v.y), fabs(v.z)); +} +inline __host__ __device__ float4 fabs(float4 v) +{ + return make_float4(fabs(v.x), fabs(v.y), fabs(v.z), fabs(v.w)); +} + +inline __host__ __device__ int2 abs(int2 v) +{ + return make_int2(abs(v.x), abs(v.y)); +} +inline __host__ __device__ int3 abs(int3 v) +{ + return make_int3(abs(v.x), abs(v.y), abs(v.z)); +} +inline __host__ __device__ int4 abs(int4 v) +{ + return make_int4(abs(v.x), abs(v.y), abs(v.z), abs(v.w)); +} + +//////////////////////////////////////////////////////////////////////////////// +// reflect +// - returns reflection of incident ray I around surface normal N +// - N should be normalized, reflected vector's length is equal to length of I +//////////////////////////////////////////////////////////////////////////////// + +inline __host__ __device__ float3 reflect(float3 i, float3 n) +{ + return i - 2.0f * n * dot(n,i); +} + +//////////////////////////////////////////////////////////////////////////////// +// cross product +//////////////////////////////////////////////////////////////////////////////// + +inline __host__ __device__ float3 cross(float3 a, float3 b) +{ + return make_float3(a.y*b.z - a.z*b.y, a.z*b.x - a.x*b.z, a.x*b.y - a.y*b.x); +} + +//////////////////////////////////////////////////////////////////////////////// +// smoothstep +// - returns 0 if x < a +// - returns 1 if x > b +// - otherwise returns smooth interpolation between 0 and 1 based on x +//////////////////////////////////////////////////////////////////////////////// + +inline __device__ __host__ float smoothstep(float a, float b, float x) +{ + float y = clamp((x - a) / (b - a), 0.0f, 1.0f); + return (y*y*(3.0f - (2.0f*y))); +} +inline __device__ __host__ float2 smoothstep(float2 a, float2 b, float2 x) +{ + float2 y = clamp((x - a) / (b - a), 0.0f, 1.0f); + return (y*y*(make_float2(3.0f) - (make_float2(2.0f)*y))); +} +inline __device__ __host__ float3 smoothstep(float3 a, float3 b, float3 x) +{ + float3 y = clamp((x - a) / (b - a), 0.0f, 1.0f); + return (y*y*(make_float3(3.0f) - (make_float3(2.0f)*y))); +} +inline __device__ __host__ float4 smoothstep(float4 a, float4 b, float4 x) +{ + float4 y = clamp((x - a) / (b - a), 0.0f, 1.0f); + return (y*y*(make_float4(3.0f) - (make_float4(2.0f)*y))); +} + +#endif diff --git a/utils/helper_string.h b/utils/helper_string.h new file mode 100644 index 0000000..9b68cc7 --- /dev/null +++ b/utils/helper_string.h @@ -0,0 +1,526 @@ +/** + * Copyright 1993-2013 NVIDIA Corporation. All rights reserved. + * + * Please refer to the NVIDIA end user license agreement (EULA) associated + * with this source code for terms and conditions that govern your use of + * this software. Any use, reproduction, disclosure, or distribution of + * this software and related documentation outside the terms of the EULA + * is strictly prohibited. + * + */ + +// These are helper functions for the SDK samples (string parsing, timers, etc) +#ifndef STRING_HELPER_H +#define STRING_HELPER_H + +#include +#include +#include +#include + +#if defined(WIN32) || defined(_WIN32) || defined(WIN64) || defined(_WIN64) +#ifndef _CRT_SECURE_NO_DEPRECATE +#define _CRT_SECURE_NO_DEPRECATE +#endif +#ifndef STRCASECMP +#define STRCASECMP _stricmp +#endif +#ifndef STRNCASECMP +#define STRNCASECMP _strnicmp +#endif +#ifndef STRCPY +#define STRCPY(sFilePath, nLength, sPath) strcpy_s(sFilePath, nLength, sPath) +#endif + +#ifndef FOPEN +#define FOPEN(fHandle,filename,mode) fopen_s(&fHandle, filename, mode) +#endif +#ifndef FOPEN_FAIL +#define FOPEN_FAIL(result) (result != 0) +#endif +#ifndef SSCANF +#define SSCANF sscanf_s +#endif +#ifndef SPRINTF +#define SPRINTF sprintf_s +#endif +#else // Linux Includes +#include +#include + +#ifndef STRCASECMP +#define STRCASECMP strcasecmp +#endif +#ifndef STRNCASECMP +#define STRNCASECMP strncasecmp +#endif +#ifndef STRCPY +#define STRCPY(sFilePath, nLength, sPath) strcpy(sFilePath, sPath) +#endif + +#ifndef FOPEN +#define FOPEN(fHandle,filename,mode) (fHandle = fopen(filename, mode)) +#endif +#ifndef FOPEN_FAIL +#define FOPEN_FAIL(result) (result == NULL) +#endif +#ifndef SSCANF +#define SSCANF sscanf +#endif +#ifndef SPRINTF +#define SPRINTF sprintf +#endif +#endif + +#ifndef EXIT_WAIVED +#define EXIT_WAIVED 2 +#endif + +// CUDA Utility Helper Functions +inline int stringRemoveDelimiter(char delimiter, const char *string) +{ + int string_start = 0; + + while (string[string_start] == delimiter) + { + string_start++; + } + + if (string_start >= (int)strlen(string)-1) + { + return 0; + } + + return string_start; +} + +inline int getFileExtension(char *filename, char **extension) +{ + int string_length = (int)strlen(filename); + + while (filename[string_length--] != '.') + { + if (string_length == 0) + break; + } + + if (string_length > 0) string_length += 2; + + if (string_length == 0) + *extension = NULL; + else + *extension = &filename[string_length]; + + return string_length; +} + + +inline bool checkCmdLineFlag(const int argc, const char **argv, const char *string_ref) +{ + bool bFound = false; + + if (argc >= 1) + { + for (int i=1; i < argc; i++) + { + int string_start = stringRemoveDelimiter('-', argv[i]); + const char *string_argv = &argv[i][string_start]; + + const char *equal_pos = strchr(string_argv, '='); + int argv_length = (int)(equal_pos == 0 ? strlen(string_argv) : equal_pos - string_argv); + + int length = (int)strlen(string_ref); + + if (length == argv_length && !STRNCASECMP(string_argv, string_ref, length)) + { + bFound = true; + continue; + } + } + } + + return bFound; +} + +// This function wraps the CUDA Driver API into a template function +template +inline bool getCmdLineArgumentValue(const int argc, const char **argv, const char *string_ref, T *value) +{ + bool bFound = false; + + if (argc >= 1) + { + for (int i=1; i < argc; i++) + { + int string_start = stringRemoveDelimiter('-', argv[i]); + const char *string_argv = &argv[i][string_start]; + int length = (int)strlen(string_ref); + + if (!STRNCASECMP(string_argv, string_ref, length)) + { + if (length+1 <= (int)strlen(string_argv)) + { + int auto_inc = (string_argv[length] == '=') ? 1 : 0; + *value = (T)atoi(&string_argv[length + auto_inc]); + } + + bFound = true; + i=argc; + } + } + } + + return bFound; +} + +inline int getCmdLineArgumentInt(const int argc, const char **argv, const char *string_ref) +{ + bool bFound = false; + int value = -1; + + if (argc >= 1) + { + for (int i=1; i < argc; i++) + { + int string_start = stringRemoveDelimiter('-', argv[i]); + const char *string_argv = &argv[i][string_start]; + int length = (int)strlen(string_ref); + + if (!STRNCASECMP(string_argv, string_ref, length)) + { + if (length+1 <= (int)strlen(string_argv)) + { + int auto_inc = (string_argv[length] == '=') ? 1 : 0; + value = atoi(&string_argv[length + auto_inc]); + } + else + { + value = 0; + } + + bFound = true; + continue; + } + } + } + + if (bFound) + { + return value; + } + else + { + return 0; + } +} + +inline float getCmdLineArgumentFloat(const int argc, const char **argv, const char *string_ref) +{ + bool bFound = false; + float value = -1; + + if (argc >= 1) + { + for (int i=1; i < argc; i++) + { + int string_start = stringRemoveDelimiter('-', argv[i]); + const char *string_argv = &argv[i][string_start]; + int length = (int)strlen(string_ref); + + if (!STRNCASECMP(string_argv, string_ref, length)) + { + if (length+1 <= (int)strlen(string_argv)) + { + int auto_inc = (string_argv[length] == '=') ? 1 : 0; + value = (float)atof(&string_argv[length + auto_inc]); + } + else + { + value = 0.f; + } + + bFound = true; + continue; + } + } + } + + if (bFound) + { + return value; + } + else + { + return 0; + } +} + +inline bool getCmdLineArgumentString(const int argc, const char **argv, + const char *string_ref, char **string_retval) +{ + bool bFound = false; + + if (argc >= 1) + { + for (int i=1; i < argc; i++) + { + int string_start = stringRemoveDelimiter('-', argv[i]); + char *string_argv = (char *)&argv[i][string_start]; + int length = (int)strlen(string_ref); + + if (!STRNCASECMP(string_argv, string_ref, length)) + { + *string_retval = &string_argv[length+1]; + bFound = true; + continue; + } + } + } + + if (!bFound) + { + *string_retval = NULL; + } + + return bFound; +} + +////////////////////////////////////////////////////////////////////////////// +//! Find the path for a file assuming that +//! files are found in the searchPath. +//! +//! @return the path if succeeded, otherwise 0 +//! @param filename name of the file +//! @param executable_path optional absolute path of the executable +////////////////////////////////////////////////////////////////////////////// +inline char *sdkFindFilePath(const char *filename, const char *executable_path) +{ + // defines a variable that is replaced with the name of the executable + + // Typical relative search paths to locate needed companion files (e.g. sample input data, or JIT source files) + // The origin for the relative search may be the .exe file, a .bat file launching an .exe, a browser .exe launching the .exe or .bat, etc + const char *searchPath[] = + { + "./", // same dir + "./common/", // "/common/" subdir + "./common/data/", // "/common/data/" subdir + "./data/", // "/data/" subdir + "./src/", // "/src/" subdir + "./src//data/", // "/src//data/" subdir + "./inc/", // "/inc/" subdir + "./0_Simple/", // "/0_Simple/" subdir + "./1_Utilities/", // "/1_Utilities/" subdir + "./2_Graphics/", // "/2_Graphics/" subdir + "./3_Imaging/", // "/3_Imaging/" subdir + "./4_Finance/", // "/4_Finance/" subdir + "./5_Simulations/", // "/5_Simulations/" subdir + "./6_Advanced/", // "/6_Advanced/" subdir + "./7_CUDALibraries/", // "/7_CUDALibraries/" subdir + "./8_Android/", // "/8_Android/" subdir + "./samples/", // "/samples/" subdir + + "./0_Simple//data/", // "/0_Simple//data/" subdir + "./1_Utilities//data/", // "/1_Utilities//data/" subdir + "./2_Graphics//data/", // "/2_Graphics//data/" subdir + "./3_Imaging//data/", // "/3_Imaging//data/" subdir + "./4_Finance//data/", // "/4_Finance//data/" subdir + "./5_Simulations//data/", // "/5_Simulations//data/" subdir + "./6_Advanced//data/", // "/6_Advanced//data/" subdir + "./7_CUDALibraries//", // "/7_CUDALibraries//" subdir + "./7_CUDALibraries//data/", // "/7_CUDALibraries//data/" subdir + + "../", // up 1 in tree + "../common/", // up 1 in tree, "/common/" subdir + "../common/data/", // up 1 in tree, "/common/data/" subdir + "../data/", // up 1 in tree, "/data/" subdir + "../src/", // up 1 in tree, "/src/" subdir + "../inc/", // up 1 in tree, "/inc/" subdir + + "../0_Simple//data/", // up 1 in tree, "/0_Simple//" subdir + "../1_Utilities//data/", // up 1 in tree, "/1_Utilities//" subdir + "../2_Graphics//data/", // up 1 in tree, "/2_Graphics//" subdir + "../3_Imaging//data/", // up 1 in tree, "/3_Imaging//" subdir + "../4_Finance//data/", // up 1 in tree, "/4_Finance//" subdir + "../5_Simulations//data/", // up 1 in tree, "/5_Simulations//" subdir + "../6_Advanced//data/", // up 1 in tree, "/6_Advanced//" subdir + "../7_CUDALibraries//data/",// up 1 in tree, "/7_CUDALibraries//" subdir + "../8_Android//data/", // up 1 in tree, "/8_Android//" subdir + "../samples//data/", // up 1 in tree, "/samples//" subdir + "../../", // up 2 in tree + "../../common/", // up 2 in tree, "/common/" subdir + "../../common/data/", // up 2 in tree, "/common/data/" subdir + "../../data/", // up 2 in tree, "/data/" subdir + "../../src/", // up 2 in tree, "/src/" subdir + "../../inc/", // up 2 in tree, "/inc/" subdir + "../../sandbox//data/", // up 2 in tree, "/sandbox//" subdir + "../../0_Simple//data/", // up 2 in tree, "/0_Simple//" subdir + "../../1_Utilities//data/", // up 2 in tree, "/1_Utilities//" subdir + "../../2_Graphics//data/", // up 2 in tree, "/2_Graphics//" subdir + "../../3_Imaging//data/", // up 2 in tree, "/3_Imaging//" subdir + "../../4_Finance//data/", // up 2 in tree, "/4_Finance//" subdir + "../../5_Simulations//data/", // up 2 in tree, "/5_Simulations//" subdir + "../../6_Advanced//data/", // up 2 in tree, "/6_Advanced//" subdir + "../../7_CUDALibraries//data/", // up 2 in tree, "/7_CUDALibraries//" subdir + "../../8_Android//data/", // up 2 in tree, "/8_Android//" subdir + "../../samples//data/", // up 2 in tree, "/samples//" subdir + "../../../", // up 3 in tree + "../../../src//", // up 3 in tree, "/src//" subdir + "../../../src//data/", // up 3 in tree, "/src//data/" subdir + "../../../src//src/", // up 3 in tree, "/src//src/" subdir + "../../../src//inc/", // up 3 in tree, "/src//inc/" subdir + "../../../sandbox//", // up 3 in tree, "/sandbox//" subdir + "../../../sandbox//data/", // up 3 in tree, "/sandbox//data/" subdir + "../../../sandbox//src/", // up 3 in tree, "/sandbox//src/" subdir + "../../../sandbox//inc/", // up 3 in tree, "/sandbox//inc/" subdir + "../../../0_Simple//data/", // up 3 in tree, "/0_Simple//" subdir + "../../../1_Utilities//data/", // up 3 in tree, "/1_Utilities//" subdir + "../../../2_Graphics//data/", // up 3 in tree, "/2_Graphics//" subdir + "../../../3_Imaging//data/", // up 3 in tree, "/3_Imaging//" subdir + "../../../4_Finance//data/", // up 3 in tree, "/4_Finance//" subdir + "../../../5_Simulations//data/", // up 3 in tree, "/5_Simulations//" subdir + "../../../6_Advanced//data/", // up 3 in tree, "/6_Advanced//" subdir + "../../../7_CUDALibraries//data/", // up 3 in tree, "/7_CUDALibraries//" subdir + "../../../8_Android//data/", // up 3 in tree, "/8_Android//" subdir + "../../../0_Simple//", // up 3 in tree, "/0_Simple//" subdir + "../../../1_Utilities//", // up 3 in tree, "/1_Utilities//" subdir + "../../../2_Graphics//", // up 3 in tree, "/2_Graphics//" subdir + "../../../3_Imaging//", // up 3 in tree, "/3_Imaging//" subdir + "../../../4_Finance//", // up 3 in tree, "/4_Finance//" subdir + "../../../5_Simulations//", // up 3 in tree, "/5_Simulations//" subdir + "../../../6_Advanced//", // up 3 in tree, "/6_Advanced//" subdir + "../../../7_CUDALibraries//", // up 3 in tree, "/7_CUDALibraries//" subdir + "../../../8_Android//", // up 3 in tree, "/8_Android//" subdir + "../../../samples//data/", // up 3 in tree, "/samples//" subdir + "../../../common/", // up 3 in tree, "../../../common/" subdir + "../../../common/data/", // up 3 in tree, "../../../common/data/" subdir + "../../../data/", // up 3 in tree, "../../../data/" subdir + "../../../../", // up 4 in tree + "../../../../src//", // up 4 in tree, "/src//" subdir + "../../../../src//data/", // up 4 in tree, "/src//data/" subdir + "../../../../src//src/", // up 4 in tree, "/src//src/" subdir + "../../../../src//inc/", // up 4 in tree, "/src//inc/" subdir + "../../../../sandbox//", // up 4 in tree, "/sandbox//" subdir + "../../../../sandbox//data/", // up 4 in tree, "/sandbox//data/" subdir + "../../../../sandbox//src/", // up 4 in tree, "/sandbox//src/" subdir + "../../../../sandbox//inc/", // up 4 in tree, "/sandbox//inc/" subdir + "../../../../0_Simple//data/", // up 4 in tree, "/0_Simple//" subdir + "../../../../1_Utilities//data/", // up 4 in tree, "/1_Utilities//" subdir + "../../../../2_Graphics//data/", // up 4 in tree, "/2_Graphics//" subdir + "../../../../3_Imaging//data/", // up 4 in tree, "/3_Imaging//" subdir + "../../../../4_Finance//data/", // up 4 in tree, "/4_Finance//" subdir + "../../../../5_Simulations//data/",// up 4 in tree, "/5_Simulations//" subdir + "../../../../6_Advanced//data/", // up 4 in tree, "/6_Advanced//" subdir + "../../../../7_CUDALibraries//data/", // up 4 in tree, "/7_CUDALibraries//" subdir + "../../../../8_Android//data/", // up 4 in tree, "/8_Android//" subdir + "../../../../0_Simple//", // up 4 in tree, "/0_Simple//" subdir + "../../../../1_Utilities//", // up 4 in tree, "/1_Utilities//" subdir + "../../../../2_Graphics//", // up 4 in tree, "/2_Graphics//" subdir + "../../../../3_Imaging//", // up 4 in tree, "/3_Imaging//" subdir + "../../../../4_Finance//", // up 4 in tree, "/4_Finance//" subdir + "../../../../5_Simulations//",// up 4 in tree, "/5_Simulations//" subdir + "../../../../6_Advanced//", // up 4 in tree, "/6_Advanced//" subdir + "../../../../7_CUDALibraries//", // up 4 in tree, "/7_CUDALibraries//" subdir + "../../../../8_Android//", // up 4 in tree, "/8_Android//" subdir + "../../../../samples//data/", // up 4 in tree, "/samples//" subdir + "../../../../common/", // up 4 in tree, "../../../common/" subdir + "../../../../common/data/", // up 4 in tree, "../../../common/data/" subdir + "../../../../data/", // up 4 in tree, "../../../data/" subdir + "../../../../../", // up 5 in tree + "../../../../../src//", // up 5 in tree, "/src//" subdir + "../../../../../src//data/", // up 5 in tree, "/src//data/" subdir + "../../../../../src//src/", // up 5 in tree, "/src//src/" subdir + "../../../../../src//inc/", // up 5 in tree, "/src//inc/" subdir + "../../../../../sandbox//", // up 5 in tree, "/sandbox//" subdir + "../../../../../sandbox//data/", // up 5 in tree, "/sandbox//data/" subdir + "../../../../../sandbox//src/", // up 5 in tree, "/sandbox//src/" subdir + "../../../../../sandbox//inc/", // up 5 in tree, "/sandbox//inc/" subdir + "../../../../../0_Simple//data/", // up 5 in tree, "/0_Simple//" subdir + "../../../../../1_Utilities//data/", // up 5 in tree, "/1_Utilities//" subdir + "../../../../../2_Graphics//data/", // up 5 in tree, "/2_Graphics//" subdir + "../../../../../3_Imaging//data/", // up 5 in tree, "/3_Imaging//" subdir + "../../../../../4_Finance//data/", // up 5 in tree, "/4_Finance//" subdir + "../../../../../5_Simulations//data/",// up 5 in tree, "/5_Simulations//" subdir + "../../../../../6_Advanced//data/", // up 5 in tree, "/6_Advanced//" subdir + "../../../../../7_CUDALibraries//data/", // up 5 in tree, "/7_CUDALibraries//" subdir + "../../../../../8_Android//data/", // up 5 in tree, "/8_Android//" subdir + "../../../../../samples//data/", // up 5 in tree, "/samples//" subdir + "../../../../../common/", // up 5 in tree, "../../../common/" subdir + "../../../../../common/data/", // up 5 in tree, "../../../common/data/" subdir + }; + + // Extract the executable name + std::string executable_name; + + if (executable_path != 0) + { + executable_name = std::string(executable_path); + +#if defined(WIN32) || defined(_WIN32) || defined(WIN64) || defined(_WIN64) + // Windows path delimiter + size_t delimiter_pos = executable_name.find_last_of('\\'); + executable_name.erase(0, delimiter_pos + 1); + + if (executable_name.rfind(".exe") != std::string::npos) + { + // we strip .exe, only if the .exe is found + executable_name.resize(executable_name.size() - 4); + } + +#else + // Linux & OSX path delimiter + size_t delimiter_pos = executable_name.find_last_of('/'); + executable_name.erase(0,delimiter_pos+1); +#endif + } + + // Loop over all search paths and return the first hit + for (unsigned int i = 0; i < sizeof(searchPath)/sizeof(char *); ++i) + { + std::string path(searchPath[i]); + size_t executable_name_pos = path.find(""); + + // If there is executable_name variable in the searchPath + // replace it with the value + if (executable_name_pos != std::string::npos) + { + if (executable_path != 0) + { + path.replace(executable_name_pos, strlen(""), executable_name); + } + else + { + // Skip this path entry if no executable argument is given + continue; + } + } + +#ifdef _DEBUG + printf("sdkFindFilePath <%s> in %s\n", filename, path.c_str()); +#endif + + // Test if the file exists + path.append(filename); + FILE *fp; + FOPEN(fp, path.c_str(), "rb"); + + if (fp != NULL) + { + fclose(fp); + // File found + // returning an allocated array here for backwards compatibility reasons + char *file_path = (char *) malloc(path.length() + 1); + STRCPY(file_path, path.length() + 1, path.c_str()); + return file_path; + } + + if (fp) + { + fclose(fp); + } + } + + // File not found + return 0; +} + +#endif