ini

09/01_texture/01/CMakeLists.txt

@@ -0,0 +1,11 @@
+cmake_minimum_required(VERSION 3.10)
+
+set(CMAKE_CXX_STANDARD 17)
+set(CMAKE_BUILD_TYPE Release)
+
+project(hellocuda LANGUAGES CXX CUDA)
+
+add_executable(main main.cu)
+target_include_directories(main PUBLIC . ../../include)
+target_compile_options(main PUBLIC $<$<COMPILE_LANGUAGE:CUDA>:--extended-lambda>)
+target_compile_options(main PUBLIC $<$<COMPILE_LANGUAGE:CUDA>:--expt-relaxed-constexpr>)

09/01_texture/01/main.cu

@@ -0,0 +1,124 @@
+#include <cstdio>
+#include <vector>
+#include <memory>
+#include <cuda_runtime.h>
+#include "helper_cuda.h"
+#include "CudaAllocator.h"
+#include "ticktock.h"
+
+struct DisableCopy {
+    DisableCopy() = default;
+    DisableCopy(DisableCopy const &) = delete;
+    DisableCopy &operator=(DisableCopy const &) = delete;
+    DisableCopy(DisableCopy &&) = delete;
+    DisableCopy &operator=(DisableCopy &&) = delete;
+};
+
+template <class T>
+struct CudaArray : DisableCopy {
+    cudaArray *m_cuArray{};
+    std::array<unsigned int, 3> m_dim{};
+
+    struct BuildArgs {
+        std::array<unsigned int, 3> const _dim;
+        int flags = 0; // or cudaArraySurfaceLoadStore
+    };
+
+    explicit CudaArray(BuildArgs _args) : m_dim(_args.dim) {
+        cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<T>();  // or cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindUnsigned)
+        checkCudaErrors(cudaMalloc3DArray(&m_cuArray, &channelDesc, make_cudaExtent(m_dim[0], m_dim[1], m_dim[2]), _args.flags));
+    }
+
+    CudaArray &assign(T *_data) {
+        cudaMemcpy3DParms copy3DParams{};
+        copy3DParams.srcPtr = make_cudaPitchedPtr((void *)_data, m_dim[0] * sizeof(T), m_dim[1], m_dim[2]);
+        copy3DParams.dstArray = m_cuArray;
+        copy3DParams.extent = make_cudaExtent(m_dim[0], m_dim[1], m_dim[2]);
+        copy3DParams.kind = cudaMemcpyHostToDevice;
+        checkCudaErrors(cudaMemcpy3D(&copy3DParams));
+        return *this;
+    }
+
+    ~CudaArray() {
+        checkCudaErrors(cudaFreeArray(m_cuArray));
+    }
+
+    operator cudaArray *() const {
+        return m_cuArray;
+    }
+};
+
+template <class T>
+struct CudaSurface : DisableCopy {
+    cudaSurfaceObject_t m_cuSuf{};
+    CudaArray<T> m_cuarr;
+
+    explicit CudaSurface(typename CudaArray<T>::BuildArgs _cuarrArgs) : m_cuarr(_cuarrArgs) {
+        cudaResourceDesc resDesc{};
+        resDesc.resType = cudaResourceTypeArray;
+
+        resDesc.res.array.array = m_cuarr.m_cuArray;
+        cudaCreateSurfaceObject(&m_cuSuf, &resDesc);
+    }
+
+    ~CudaTexture() {
+        checkCudaErrors(cudaDestroySurfaceObject(m_cuSuf));
+    }
+
+    CudaArray<T> &array() {
+        return m_cuarr;
+    }
+
+    operator cudaSurfaceObject_t() const {
+        return m_cuSuf;
+    }
+};
+
+template <class T>
+struct CudaTexture : DisableCopy {
+    cudaTextureObject_t m_cuTex{};
+    CudaArray<T> m_cuarr;
+
+    explicit CudaTexture(typename CudaArray<T>::BuildArgs _cuarrArgs) : m_cuarr(_cuarrArgs) {
+        cudaResourceDesc resDesc{};
+        resDesc.resType = cudaResourceTypeArray;
+        resDesc.res.array.array = m_cuarr.m_cuArray;
+
+        cudaTextureDesc texDesc{};
+        texDesc.addressMode[0] = cudaAddressModeClamp; // or cudaAddressModeWrap
+        texDesc.addressMode[1] = cudaAddressModeClamp; // or cudaAddressModeWrap
+        texDesc.addressMode[2] = cudaAddressModeClamp; // or cudaAddressModeWrap
+        texDesc.filterMode = cudaFilterModePoint;      // or cudaFilterModeLinear
+        texDesc.readMode = cudaReadModeElementType;    // or cudaReadModeNormalizedFloat
+        texDesc.normalizedCoords = false;              // or true
+
+        checkCudaErrors(cudaCreateTextureObject(&m_cuTex, &resDesc, &texDesc, NULL));
+    }
+
+    ~CudaTexture() {
+        checkCudaErrors(cudaDestroyTextureObject(m_cuTex));
+    }
+
+    CudaArray<T> &array() {
+        return m_cuarr;
+    }
+
+    operator cudaTextureObject_t() const {
+        return m_cuTex;
+    }
+};
+
+__global__ void kernel(cudaSurfaceObject_t out, cudaTextureObject_t in) {
+    int x = 0, y = 0;
+    float fx = 0, fy = 0, fz = 0;
+    float value = tex3D<float>(in, fx, fy, fz);
+    value += 1;
+    surf2Dwrite(value, out, x, y);
+    // or cudaBoundaryModeTrap, cudaBoundaryModeClamp
+}
+
+int main() {
+    CudaSurface<float> out({{1, 1, 1}, cudaArraySurfaceLoadStore});
+    CudaTexture<float> in({{1, 1, 1}, 0});
+    return 0;
+}

09/include/CudaAllocator.h

@@ -0,0 +1,31 @@
+#pragma once
+
+#include <cstddef>
+#include <utility>
+#include <cuda_runtime.h>
+#include "helper_cuda.h"
+
+template <class T>
+struct CudaAllocator {
+    using value_type = T;
+
+    T *allocate(size_t size) {
+        T *ptr = nullptr;
+        checkCudaErrors(cudaMallocManaged(&ptr, size * sizeof(T)));
+        return ptr;
+    }
+
+    void deallocate(T *ptr, size_t size = 0) {
+        checkCudaErrors(cudaFree(ptr));
+    }
+
+    template <class ...Args>
+    void construct(T *p, Args &&...args) {
+        if constexpr (!(sizeof...(Args) == 0 && std::is_pod_v<T>))
+            ::new((void *)p) T(std::forward<Args>(args)...);
+    }
+
+    constexpr bool operator==(CudaAllocator<T> const &other) const {
+        return this == &other;
+    }
+};

09/include/CudaTexture.h

@@ -0,0 +1,39 @@
+#pragma once
+
+
+#include "helper_cuda.h"
+#include <cuda_runtime.h>
+
+
+struct CudaTexture {
+    cudaTextureObject_t tex;
+
+    CudaTexture(CudaTexture const &) = delete;
+    CudaTexture(CudaTexture &&) = default;
+    CudaTexture &operator=(CudaTexture const &) = delete;
+    CudaTexture &operator=(CudaTexture &&) = default;
+
+    template <class T>
+    CudaTexture(T *dataDev, int width, int height) {
+        cudaTextureObject_t tex;
+        cudaResourceDesc resDesc;
+        memset(&resDesc, 0, sizeof(resDesc));
+        resDesc.resType = cudaResourceTypePitch2D;
+        resDesc.res.pitch2D.devPtr = dataDev;
+        resDesc.res.pitch2D.width = width;
+        resDesc.res.pitch2D.height = height;
+        resDesc.res.pitch2D.desc = cudaCreateChannelDesc<T>();
+        resDesc.res.pitch2D.pitchInBytes = width * sizeof(T);
+        cudaTextureDesc texDesc;
+        memset(&texDesc, 0, sizeof(texDesc));
+        checkCudaErrors(cudaCreateTextureObject(&tex, &resDesc, &texDesc, NULL));
+    }
+
+    ~CudaTexture() {
+        checkCudaErrors(cudaDestroyTextureObject(tex));
+    }
+
+    constexpr operator cudaTextureObject_t() const {
+        return tex;
+    }
+};