#include <Volume.h>
#include <common_device.hpp>

void Volume::createTexture(unsigned char* data) {
    unsigned char* importanceValues;
    CUDA::malloc(importanceValues, xBound * yBound * zBound);
    unsigned char* dataDevice;
    CUDA::malloc(dataDevice, xBound * yBound * zBound);
    CUDA::memcopyToDevice(dataDevice, data, xBound * yBound * zBound);

    generateCacheImportanceSDF<<<CUDA::getBlockCount(xBound * yBound * zBound), CUDA::THREADS_PER_BLOCK>>>(
            dataDevice, make_int3(xBound, yBound, zBound), importanceValues);


    auto extend = make_cudaExtent(xBound, yBound, zBound);
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<unsigned char>();
    CUDA::checkCudaError(cudaMalloc3DArray(&d_volumeArray, &channelDesc, extend));

    cudaMemcpy3DParms copyParams = {nullptr};
    copyParams.srcPtr = make_cudaPitchedPtr(importanceValues, xBound * sizeof(unsigned char), xBound, yBound);
    copyParams.dstArray = d_volumeArray;
    copyParams.extent = extend;
    copyParams.kind = cudaMemcpyDeviceToDevice;
    CUDA::checkCudaError(cudaMemcpy3D(&copyParams));

    cudaResourceDesc resDescArray{};
    memset(&resDescArray, 0, sizeof(cudaResourceDesc));
    resDescArray.resType = cudaResourceTypeArray;
    resDescArray.res.array.array = d_volumeArray;

    cudaTextureDesc texDescArray{};
    memset(&texDescArray, 0, sizeof(cudaTextureDesc));
    texDescArray.addressMode[0] = cudaAddressModeClamp;
    texDescArray.addressMode[1] = cudaAddressModeClamp;
    texDescArray.addressMode[2] = cudaAddressModeClamp;
    texDescArray.normalizedCoords = true;
    texDescArray.readMode = cudaReadModeNormalizedFloat;
    texDescArray.filterMode = cudaFilterModeLinear;

    cudaCreateTextureObject(&texture, &resDescArray, &texDescArray, nullptr);

    CUDA::free(importanceValues);
    CUDA::free(dataDevice);
}

void Volume::createTexture(unsigned short* data) {
    unsigned short* importanceValues;
    CUDA::malloc(importanceValues, xBound * yBound * zBound);
    unsigned short* dataDevice;
    CUDA::malloc(dataDevice, xBound * yBound * zBound);
    CUDA::memcopyToDevice(dataDevice, data, xBound * yBound * zBound);

    generateCacheImportanceSDF<<<CUDA::getBlockCount(xBound * yBound * zBound), CUDA::THREADS_PER_BLOCK>>>(
            dataDevice, make_int3(xBound, yBound, zBound), importanceValues);


    auto extend = make_cudaExtent(xBound, yBound, zBound);
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<unsigned short>();
    CUDA::checkCudaError(cudaMalloc3DArray(&d_volumeArray, &channelDesc, extend));

    cudaMemcpy3DParms copyParams = {nullptr};
    copyParams.srcPtr = make_cudaPitchedPtr(importanceValues, xBound * sizeof(unsigned short), xBound, yBound);
    copyParams.dstArray = d_volumeArray;
    copyParams.extent = extend;
    copyParams.kind = cudaMemcpyDeviceToDevice;
    CUDA::checkCudaError(cudaMemcpy3D(&copyParams));

    cudaResourceDesc resDescArray{};
    memset(&resDescArray, 0, sizeof(cudaResourceDesc));
    resDescArray.resType = cudaResourceTypeArray;
    resDescArray.res.array.array = d_volumeArray;

    cudaTextureDesc texDescArray{};
    memset(&texDescArray, 0, sizeof(cudaTextureDesc));
    texDescArray.addressMode[0] = cudaAddressModeClamp;
    texDescArray.addressMode[1] = cudaAddressModeClamp;
    texDescArray.addressMode[2] = cudaAddressModeClamp;
    texDescArray.normalizedCoords = true;
    texDescArray.readMode = cudaReadModeNormalizedFloat;
    texDescArray.filterMode = cudaFilterModeLinear;
    // texDescArray.filterMode = cudaFilterModePoint;

    cudaCreateTextureObject(&texture, &resDescArray, &texDescArray, nullptr);

    CUDA::free(importanceValues);
    CUDA::free(dataDevice);
}

void Volume::createTexture(short* data) {
    short* importanceValues;
    CUDA::malloc(importanceValues, xBound * yBound * zBound);
    short* dataDevice;
    CUDA::malloc(dataDevice, xBound * yBound * zBound);
    CUDA::memcopyToDevice(dataDevice, data, xBound * yBound * zBound);

    generateCacheImportanceSDF<<<CUDA::getBlockCount(xBound * yBound * zBound), CUDA::THREADS_PER_BLOCK>>>(
            dataDevice, make_int3(xBound, yBound, zBound), importanceValues);


    auto extend = make_cudaExtent(xBound, yBound, zBound);
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<short>();
    CUDA::checkCudaError(cudaMalloc3DArray(&d_volumeArray, &channelDesc, extend));

    cudaMemcpy3DParms copyParams = {nullptr};
    copyParams.srcPtr = make_cudaPitchedPtr(importanceValues, xBound * sizeof(short), xBound, yBound);
    copyParams.dstArray = d_volumeArray;
    copyParams.extent = extend;
    copyParams.kind = cudaMemcpyDeviceToDevice;
    CUDA::checkCudaError(cudaMemcpy3D(&copyParams));

    cudaResourceDesc resDescArray{};
    memset(&resDescArray, 0, sizeof(cudaResourceDesc));
    resDescArray.resType = cudaResourceTypeArray;
    resDescArray.res.array.array = d_volumeArray;

    cudaTextureDesc texDescArray{};
    memset(&texDescArray, 0, sizeof(cudaTextureDesc));
    texDescArray.addressMode[0] = cudaAddressModeClamp;
    texDescArray.addressMode[1] = cudaAddressModeClamp;
    texDescArray.addressMode[2] = cudaAddressModeClamp;
    texDescArray.normalizedCoords = true;
    texDescArray.readMode = cudaReadModeNormalizedFloat;
    texDescArray.filterMode = cudaFilterModeLinear;

    cudaCreateTextureObject(&texture, &resDescArray, &texDescArray, nullptr);

    CUDA::free(importanceValues);
    CUDA::free(dataDevice);
}
