src/lib/cuda_j2k_frame_encoder.cc

   1 /*
   2     Copyright (C) 2022 Carl Hetherington <cth@carlh.net>
   3
   4     This file is part of DCP-o-matic.
   5
   6     DCP-o-matic is free software; you can redistribute it and/or modify
   7     it under the terms of the GNU General Public License as published by
   8     the Free Software Foundation; either version 2 of the License, or
   9     (at your option) any later version.
  10
  11     DCP-o-matic is distributed in the hope that it will be useful,
  12     but WITHOUT ANY WARRANTY; without even the implied warranty of
  13     MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
  14     GNU General Public License for more details.
  15
  16     You should have received a copy of the GNU General Public License
  17     along with DCP-o-matic.  If not, see <http://www.gnu.org/licenses/>.
  18
  19 */
  20
  21
  22 #include "cross.h"
  23 #include "cuda_j2k_frame_encoder.h"
  24 #include "dcpomatic_log.h"
  25 #include "dcp_video.h"
  26 #include "exceptions.h"
  27 #include "player_video.h"
  28 #include <dcp/openjpeg_image.h>
  29 #include <nvjpeg2k.h>
  30 #include <vector>
  31
  32
  33 using std::make_pair;
  34 using std::vector;
  35 using boost::optional;
  36
  37
  38 CUDAJ2KFrameEncoder::CUDAJ2KFrameEncoder()
  39 {
  40         nvjpeg2kEncoderCreateSimple(&_encoder_handle);
  41         nvjpeg2kEncodeStateCreate(_encoder_handle, &_encoder_state);
  42         nvjpeg2kEncodeParamsCreate(&_encoder_params);
  43
  44         cudaStreamCreateWithFlags(&_stream, cudaStreamNonBlocking);
  45 }
  46
  47
  48 CUDAJ2KFrameEncoder::Input::Input(DCPVideo const& vf, cudaStream_t stream)
  49         : _index(vf.index())
  50         , _eyes(vf.eyes())
  51 {
  52         _xyz = convert_to_xyz(vf.frame(), boost::bind(&Log::dcp_log, dcpomatic_log.get(), _1, _2));
  53
  54         for (int i = 0; i < 3; ++i) {
  55                 _pixel_data_h[i] = reinterpret_cast<uint8_t*>(_xyz->data(i));
  56         }
  57
  58         auto const pitch = _xyz->size().width * 2;
  59
  60         for (int i = 0; i < 3; ++i) {
  61                 _pitch_in_bytes[i] = pitch;
  62                 auto status = cudaMallocPitch(
  63                         reinterpret_cast<void**>(&_pixel_data_d[i]),
  64                         &_pitch_in_bytes[i],
  65                         pitch,
  66                         _xyz->size().height
  67                         );
  68
  69                 if (status != cudaSuccess) {
  70                         throw CUDAError("cudaMallocPitch", status);
  71                 }
  72
  73                 status = cudaMemcpy2DAsync(
  74                         _pixel_data_d[i],
  75                         _pitch_in_bytes[i],
  76                         _pixel_data_h[i],
  77                         _pitch_in_bytes[i],
  78                         pitch,
  79                         _xyz->size().height,
  80                         cudaMemcpyHostToDevice,
  81                         stream
  82                         );
  83
  84                 if (status != cudaSuccess) {
  85                         throw CUDAError("cudaMemcpy2D", status);
  86                 }
  87         }
  88
  89         _device_image.num_components = 3;
  90         _device_image.pixel_data = reinterpret_cast<void**>(_pixel_data_d);
  91         _device_image.pixel_type = NVJPEG2K_UINT16;
  92         _device_image.pitch_in_bytes = reinterpret_cast<size_t*>(_pitch_in_bytes);
  93 }
  94
  95
  96 CUDAJ2KFrameEncoder::Input::Input(Input&& other)
  97         : _index(other._index)
  98         , _eyes(other._eyes)
  99 {
 100         for (int i = 0; i < 3; ++i) {
 101                 _pixel_data_d[i] = other._pixel_data_d[i];
 102                 other._pixel_data_d[i] = nullptr;
 103                 _pitch_in_bytes[i] = other._pitch_in_bytes[i];
 104         }
 105
 106         _device_image.num_components = other._device_image.num_components;
 107         _device_image.pixel_data = reinterpret_cast<void**>(_pixel_data_d);
 108         _device_image.pixel_type = NVJPEG2K_UINT16;
 109         _device_image.pitch_in_bytes = reinterpret_cast<size_t*>(_pitch_in_bytes);
 110 }
 111
 112
 113 CUDAJ2KFrameEncoder::Input::~Input()
 114 {
 115         cudaFree(_pixel_data_d[0]);
 116         cudaFree(_pixel_data_d[1]);
 117         cudaFree(_pixel_data_d[2]);
 118 }
 119
 120
 121 optional<dcp::ArrayData>
 122 CUDAJ2KFrameEncoder::encode(DCPVideo const& vf)
 123 {
 124         auto input = Input(vf, _stream);
 125
 126         auto const size = vf.frame()->out_size();
 127         DCPOMATIC_ASSERT(!_size || size == *_size);
 128         _size = size;
 129
 130         DCPOMATIC_ASSERT(!_resolution || vf.resolution() == *_resolution);
 131         _resolution = vf.resolution();
 132
 133         nvjpeg2kImageComponentInfo_t info[3];
 134         for (int i = 0; i < 3; ++i) {
 135                 info[i].component_width = _size->width;
 136                 info[i].component_height = _size->height;
 137                 info[i].precision = 12;
 138                 info[i].sgn = 0;
 139         }
 140
 141         nvjpeg2kEncodeConfig_t config;
 142         memset(&config, 0, sizeof(config));
 143         config.stream_type = NVJPEG2K_STREAM_J2K;
 144         config.color_space = NVJPEG2K_COLORSPACE_SRGB;
 145         config.image_width = _size->width;
 146         config.image_height = _size->height;
 147         config.num_components = 3;
 148         config.image_comp_info = reinterpret_cast<nvjpeg2kImageComponentInfo_t*>(&info);
 149         config.code_block_w = 32;
 150         config.code_block_h = 32;
 151         config.irreversible = 0;
 152         config.mct_mode = 1;
 153         config.prog_order = NVJPEG2K_CPRL;
 154         config.num_resolutions = *_resolution == Resolution::FOUR_K ? 7 : 6;
 155
 156         auto status = nvjpeg2kEncodeParamsSetEncodeConfig(_encoder_params, &config);
 157         if (status != NVJPEG2K_STATUS_SUCCESS) {
 158                 throw CUDAError("nvjpeg2kEncodeParamsSetEncodeConfig", status);
 159         }
 160
 161         // XXX: quality
 162         status = nvjpeg2kEncodeParamsSetQuality(_encoder_params, 30);
 163         if (status != NVJPEG2K_STATUS_SUCCESS) {
 164                 throw CUDAError("nvjpeg2kEncodeParamsSetQuality", status);
 165         }
 166
 167         status = nvjpeg2kEncode(_encoder_handle, _encoder_state, _encoder_params, input.device_image(), _stream);
 168         if (status != NVJPEG2K_STATUS_SUCCESS) {
 169                 throw CUDAError("nvjpeg2kEncode", status);
 170         }
 171
 172         size_t compressed_size;
 173         status = nvjpeg2kEncodeRetrieveBitstream(_encoder_handle, _encoder_state, nullptr, &compressed_size, _stream);
 174
 175         dcp::ArrayData output(compressed_size);
 176         status = nvjpeg2kEncodeRetrieveBitstream(_encoder_handle, _encoder_state, output.data(), &compressed_size, _stream);
 177         if (status != NVJPEG2K_STATUS_SUCCESS) {
 178                 throw CUDAError("nvjpeg2kEncodeRetrieveBitstream", status);
 179         }
 180
 181         return output;
 182 }
 183
 184
 185 void
 186 CUDAJ2KFrameEncoder::log_thread_start ()
 187 {
 188        LOG_TIMING("start-encoder-thread thread=%1", thread_id());
 189 }
 190
 191
 192 void
 193 CUDAJ2KFrameEncoder::flush()
 194 {
 195
 196 }