Wait for the caller's stream before the nvJPEG hardware decode writes

[karkuspeter]

decode_jpeg(..., device="cuda") can silently corrupt unrelated GPU memory.

nvJPEG's hardware engine does not honour the stream it is given: with work already queued on that
stream, nvjpegDecodeBatched writes its destination buffers before that work has run. Those
destinations come from PyTorch's caching allocator, which hands out a block as soon as it is freed on
the stream that owns it — so the decoder can be handed memory that queued kernels are still reading,
and overwrite it. The caller sees wrong numbers, NaNs, or a device-side assert in code that never
touched the decoder.

The wrapper here is not at fault; it submits on the caller's stream, which is the right thing to do.
The fix belongs in nvJPEG, and until it lands the decoder has to defend itself: this adds one
cudaStreamSynchronize(stream) in decode_batched_hardware, after the outputs are allocated and
before the first nvjpegDecodeBatched. It costs no wall time today, because decode_images already
host-synchronizes that same stream before returning, so the wait moves earlier rather than being
added — median 1920x1080 decode over 200 calls is 1.186 ms with the barrier and 1.183 ms without. The
software path is untouched.

Reproducing it

repro.py below needs only torch and torchcodec, on a GPU with an nvJPEG hardware engine. Each
iteration fills four tensors and synchronizes so they are resident, queues ~100 ms of work followed by
one reduction per tensor, frees the tensors while those reductions are still queued, then decodes into
what the allocator just took back. The sums are exact integers, so they are only wrong if the decode
wrote early.

H100 80GB HBM3, driver 575.57.08, torch 2.13.0+cu129, torchcodec v0.16.0 built against the CUDA 12.8
toolkit (nvJPEG 12.3.5); the released 0.16.0+cu129 wheel (nvJPEG 12.4.0.76) behaves the same.

$ python repro.py
corrupted: 98/100, decode landed on a freed victim block: 98/100

$ python repro.py --barrier    # what this PR does inside the decoder
corrupted: 0/100, decode landed on a freed victim block: 98/100

$ python repro.py --cpu        # the workload on its own is fine
corrupted: 0/100, decode landed on a freed victim block: 0/100

repro.py

"""Show that decode_jpeg(device="cuda") can overwrite memory that queued kernels still read.

Needs only torch and torchcodec, on a GPU with an nvJPEG hardware engine.

    python repro.py            # corrupted: 98/100
    python repro.py --barrier  # corrupted: 0/100  (what this PR does inside the decoder)
    python repro.py --cpu      # corrupted: 0/100  (the workload itself is fine)
"""

import argparse

import torch
from torchcodec.decoders import decode_jpeg
from torchcodec.encoders import JpegEncoder

HEIGHT, WIDTH = 1080, 1920
NUM_BYTES = 3 * HEIGHT * WIDTH  # the decode output's exact byte count
VALUES = (1, 2, 3, 4)


def main() -> None:
    parser = argparse.ArgumentParser()
    parser.add_argument("--iterations", type=int, default=100)
    parser.add_argument("--barrier", action="store_true", help="drain the stream first")
    parser.add_argument("--cpu", action="store_true", help="decode on the CPU instead")
    args = parser.parse_args()

    device = torch.device("cuda", torch.cuda.current_device())
    rows = torch.arange(HEIGHT, dtype=torch.uint8).view(1, HEIGHT, 1)
    cols = torch.arange(WIDTH, dtype=torch.uint8).view(1, 1, WIDTH)
    image = (rows + cols).expand(3, HEIGHT, WIDTH).contiguous()
    encoded = JpegEncoder(image).to_tensor(quality=90)  # baseline JPEG

    expected = float(sum(value * NUM_BYTES for value in VALUES))
    corrupted = landed_on_freed = 0
    for _ in range(args.iterations):
        victims = [
            torch.full((NUM_BYTES,), value, dtype=torch.uint8, device=device)
            for value in VALUES
        ]
        pointers = {victim.data_ptr() for victim in victims}
        torch.cuda.synchronize()  # the victims' contents are resident

        # Queue ~100ms of work, then reductions that read the victims, then free
        # the victims while their reductions are still queued. The sums are exact
        # integers, and are only wrong if something wrote to that memory early.
        torch.cuda._sleep(200_000_000)
        total = torch.zeros((), dtype=torch.float64, device=device)
        for victim in victims:
            total = total + victim.sum(dtype=torch.float64)
        victims.clear()

        if args.barrier:
            torch.cuda.current_stream().synchronize()
        decoded = decode_jpeg(encoded, device="cpu" if args.cpu else device)
        landed_on_freed += decoded.data_ptr() in pointers
        torch.cuda.synchronize()
        corrupted += float(total) != expected

    print(
        f"corrupted: {corrupted}/{args.iterations}, "
        f"decode landed on a freed victim block: {landed_on_freed}/{args.iterations}"
    )


if __name__ == "__main__":
    main()

Why this is nvJPEG and not the wrapper

The program below involves no framework, no allocator and no aliasing: one cudaMalloc'd destination
is reused for the whole run, and its only other user is a checksum kernel on the same stream. Each
iteration zeroes the buffer and synchronizes, queues a ~100 ms spin kernel and then the checksum, then
decodes into that buffer on that stream. The checksum must be zero.

backend	API	between the queued work and the decode	out of order
hardware	nvjpegDecodeBatched	nothing	100/100
hardware	nvjpegDecodeBatched	cudaStreamSynchronize	0/100
hardware	nvjpegDecodeBatched	cudaStreamWaitEvent on a second stream	100/100
hardware	nvjpegDecodeBatched	that same event, host-waited	0/100
hardware	nvjpegDecode	nothing	0/100
default (software)	nvjpegDecodeBatched	nothing	0/100

This contradicts the documented contract, that nvjpegDecodeBatched "is asynchronous with respect to
the host. All GPU tasks for this function will be submitted to the provided stream", together with
in-order execution within a stream. Reproduced on H100 with nvJPEG 12.3.5, 12.4.0.76, 13.0.1 and
13.2.1.

Rows three and four are why the barrier has to be a host wait. A device-side dependency does not gate
the write: submitting the decode on a second stream that holds a pending cudaStreamWaitEvent on the
busy stream is out of order in all 100 iterations, while waiting for that same event on the host first
is clean in all 100. The write happens when the call is made. The host-side call times say the same
thing from the other end — with ~100 ms queued, the hardware call returns in 0.1 ms and writes anyway,
while the software backend's call blocks the host for ~76 ms, so the software path is ordered because
it waits, not because it takes a device-side dependency. NVIDIA's DALI carries the same workaround for
the same class of problem in nvImageCodec (NVIDIA/DALI#5408).

nvjpeg_hw_stream_order.cu, built with nvcc -O2 -std=c++17 nvjpeg_hw_stream_order.cu -o nvjpeg_hw_stream_order -lnvjpeg

// Does nvJPEG's hardware decode honour the stream it is given?
//
// nvjpegDecodeBatched is documented as "asynchronous with respect to the host.
// All GPU tasks for this function will be submitted to the provided stream", and
// work submitted to one stream runs in issue order. So a kernel already queued on
// that stream must finish before the decode overwrites the destination buffers.
//
// This program checks that directly. Nothing here involves a framework, a memory
// pool, or aliasing: one device buffer is allocated once and reused, and its only
// other user is a checksum kernel on the same stream. Each iteration:
//
//   1. zero the buffer, then synchronize, so the zeros are resident;
//   2. queue a ~100ms spin kernel on the stream, then a checksum of the buffer;
//   3. call nvJPEG to decode into that same buffer, on that same stream;
//   4. synchronize and read the checksum back.
//
// The checksum must be 0. Any other value means the decode wrote into the buffer
// before the kernel that was already queued to read it had run.
//
// The third argument picks what stands between the queued work and the decode,
// which is what decides whether a caller can order the decode at all:
//
//   none           submit on the same stream, nothing else               (the bug)
//   presync        drain that stream on the host, then submit on it (the workaround)
//   event          record an event on the stream after the checksum, have a second
//                  stream wait on it, and submit there: device-side ordering only
//   event-presync  the same, plus draining that second stream on the host before
//                  submitting, which separates "the event was ignored" from "that
//                  stream was never ordered in the first place"
//
// Build and run (any baseline JPEG will do; one is encoded here with nvJPEG itself):
//   nvcc -O2 -std=c++17 nvjpeg_hw_stream_order.cu -o nvjpeg_hw_stream_order -lnvjpeg
//   ./nvjpeg_hw_stream_order                          # backend=hardware api=batched
//   ./nvjpeg_hw_stream_order hardware batched presync  # workaround: host barrier
//   ./nvjpeg_hw_stream_order hardware batched event    # cudaStreamWaitEvent only
//   ./nvjpeg_hw_stream_order default batched           # control: software backend

#include <cuda_runtime.h>
#include <nvjpeg.h>

#include <algorithm>
#include <chrono>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <string>
#include <vector>

#define CHECK_CUDA(call)                                                    \
  do {                                                                      \
    cudaError_t status = (call);                                            \
    if (status != cudaSuccess) {                                            \
      fprintf(stderr, "%s:%d %s failed: %s\n", __FILE__, __LINE__, #call,   \
              cudaGetErrorString(status));                                  \
      exit(EXIT_FAILURE);                                                   \
    }                                                                       \
  } while (0)

#define CHECK_NVJPEG(call)                                                  \
  do {                                                                      \
    nvjpegStatus_t status = (call);                                         \
    if (status != NVJPEG_STATUS_SUCCESS) {                                  \
      fprintf(stderr, "%s:%d %s failed: %d\n", __FILE__, __LINE__, #call,   \
              static_cast<int>(status));                                    \
      exit(EXIT_FAILURE);                                                   \
    }                                                                       \
  } while (0)

namespace {

constexpr int kHeight = 1080;
constexpr int kWidth = 1920;
constexpr size_t kPlane = static_cast<size_t>(kHeight) * kWidth;
constexpr size_t kBytes = 3 * kPlane;

// Keeps the stream busy long enough that the decode cannot accidentally be late.
__global__ void spin(long long cycles) {
  long long start = clock64();
  while (clock64() - start < cycles) {
  }
}

__global__ void checksum(const unsigned char* data, size_t count,
                         unsigned long long* total) {
  size_t index = blockIdx.x * blockDim.x + threadIdx.x;
  size_t stride = static_cast<size_t>(gridDim.x) * blockDim.x;
  unsigned long long local = 0;
  for (size_t i = index; i < count; i += stride) {
    local += data[i];
  }
  atomicAdd(total, local);
}

__global__ void gradient(unsigned char* planes, int height, int width) {
  size_t index = blockIdx.x * blockDim.x + threadIdx.x;
  size_t stride = static_cast<size_t>(gridDim.x) * blockDim.x;
  size_t count = static_cast<size_t>(height) * width;
  for (size_t i = index; i < count; i += stride) {
    int row = static_cast<int>(i / width);
    int column = static_cast<int>(i % width);
    unsigned char value = static_cast<unsigned char>((row * 7 + column * 13) & 0xFF);
    planes[i] = value;                    // R
    planes[count + i] = 255 - value;      // G
    planes[2 * count + i] = value / 2;    // B
  }
}

// A baseline JPEG of the gradient above, encoded with nvJPEG so that this file has
// no external dependency. Baseline single-scan 4:2:0 is what the hardware engine
// accepts, which is also what a camera or a dataset gives you.
std::vector<unsigned char> encode_baseline_jpeg() {
  nvjpegHandle_t handle;
  CHECK_NVJPEG(nvjpegCreateEx(NVJPEG_BACKEND_DEFAULT, nullptr, nullptr,
                              NVJPEG_FLAGS_DEFAULT, &handle));
  nvjpegEncoderState_t state;
  CHECK_NVJPEG(nvjpegEncoderStateCreate(handle, &state, nullptr));
  nvjpegEncoderParams_t params;
  CHECK_NVJPEG(nvjpegEncoderParamsCreate(handle, &params, nullptr));
  CHECK_NVJPEG(nvjpegEncoderParamsSetQuality(params, 90, nullptr));
  CHECK_NVJPEG(nvjpegEncoderParamsSetSamplingFactors(params, NVJPEG_CSS_420, nullptr));

  unsigned char* planes = nullptr;
  CHECK_CUDA(cudaMalloc(&planes, kBytes));
  gradient<<<256, 256>>>(planes, kHeight, kWidth);
  CHECK_CUDA(cudaDeviceSynchronize());

  nvjpegImage_t source{};
  for (int channel = 0; channel < 3; ++channel) {
    source.channel[channel] = planes + channel * kPlane;
    source.pitch[channel] = kWidth;
  }
  CHECK_NVJPEG(nvjpegEncodeImage(handle, state, params, &source, NVJPEG_INPUT_RGB,
                                 kWidth, kHeight, nullptr));
  size_t length = 0;
  CHECK_NVJPEG(nvjpegEncodeRetrieveBitstream(handle, state, nullptr, &length, nullptr));
  std::vector<unsigned char> jpeg(length);
  CHECK_NVJPEG(nvjpegEncodeRetrieveBitstream(handle, state, jpeg.data(), &length, nullptr));
  CHECK_CUDA(cudaDeviceSynchronize());
  CHECK_CUDA(cudaFree(planes));
  nvjpegEncoderParamsDestroy(params);
  nvjpegEncoderStateDestroy(state);
  nvjpegDestroy(handle);
  return jpeg;
}

} // namespace

int main(int argc, char** argv) {
  std::string backend_name = argc > 1 ? argv[1] : "hardware";
  std::string api = argc > 2 ? argv[2] : "batched";
  // What stands between the queued work and the decode; see the header.
  std::string barrier = (argc > 3 && *argv[3]) ? argv[3] : "none";
  int iterations = argc > 4 ? atoi(argv[4]) : 100;
  // The batch parameters never change here, so one initialize call is enough; pass
  // "init-each" to instead re-initialize per call, the way the PyTorch wrappers do.
  bool init_each = argc > 5 && std::string(argv[5]) == "init-each";

  bool submit_on_side = barrier == "event" || barrier == "event-presync";
  if (barrier != "none" && barrier != "presync" && !submit_on_side) {
    fprintf(stderr, "unknown barrier mode: %s\n", barrier.c_str());
    return EXIT_FAILURE;
  }

  int major = 0, minor = 0, patch = 0;
  nvjpegGetProperty(MAJOR_VERSION, &major);
  nvjpegGetProperty(MINOR_VERSION, &minor);
  nvjpegGetProperty(PATCH_LEVEL, &patch);
  cudaDeviceProp properties{};
  CHECK_CUDA(cudaGetDeviceProperties(&properties, 0));
  printf("nvjpeg %d.%d.%d  %s (sm_%d%d)  backend=%s api=%s barrier=%s iterations=%d "
         "init=%s\n",
         major, minor, patch, properties.name, properties.major, properties.minor,
         backend_name.c_str(), api.c_str(), barrier.c_str(), iterations,
         init_each ? "each-call" : "once");

  std::vector<unsigned char> jpeg = encode_baseline_jpeg();
  printf("encoded a %dx%d baseline JPEG, %zu bytes\n", kWidth, kHeight, jpeg.size());

  nvjpegBackend_t backend =
      backend_name == "hardware" ? NVJPEG_BACKEND_HARDWARE : NVJPEG_BACKEND_DEFAULT;
  nvjpegHandle_t handle;
  nvjpegStatus_t created = nvjpegCreateEx(backend, nullptr, nullptr,
                                          NVJPEG_FLAGS_DEFAULT, &handle);
  if (created == NVJPEG_STATUS_ARCH_MISMATCH) {
    printf("no hardware JPEG engine on this GPU, nothing to test\n");
    return 0;
  }
  CHECK_NVJPEG(created);
  nvjpegJpegState_t state;
  CHECK_NVJPEG(nvjpegJpegStateCreate(handle, &state));

  cudaStream_t stream;
  CHECK_CUDA(cudaStreamCreate(&stream));
  // Used by the event modes only: the decode goes here, waiting on `stream`.
  cudaStream_t side;
  CHECK_CUDA(cudaStreamCreate(&side));
  cudaEvent_t queued;
  CHECK_CUDA(cudaEventCreate(&queued));

  unsigned char* buffer = nullptr;  // the decode destination, ours for the whole run
  CHECK_CUDA(cudaMalloc(&buffer, kBytes));
  unsigned long long* total = nullptr;
  CHECK_CUDA(cudaMalloc(&total, sizeof(unsigned long long)));

  nvjpegImage_t destination{};
  for (int channel = 0; channel < 3; ++channel) {
    destination.channel[channel] = buffer + channel * kPlane;
    destination.pitch[channel] = kWidth;
  }
  const unsigned char* data = jpeg.data();
  size_t length = jpeg.size();

  if (api == "batched") {
    CHECK_NVJPEG(nvjpegDecodeBatchedInitialize(handle, state, 1, 1, NVJPEG_OUTPUT_RGB));
  }
  auto decode = [&](cudaStream_t on) {
    if (api == "batched") {
      if (init_each) {
        CHECK_NVJPEG(nvjpegDecodeBatchedInitialize(handle, state, 1, 1, NVJPEG_OUTPUT_RGB));
      }
      CHECK_NVJPEG(nvjpegDecodeBatched(handle, state, &data, &length, &destination, on));
    } else {
      CHECK_NVJPEG(nvjpegDecode(handle, state, data, length, NVJPEG_OUTPUT_RGB,
                                &destination, on));
    }
  };

  decode(stream);  // warm-up: first-call setup stays out of the measurement
  CHECK_CUDA(cudaStreamSynchronize(stream));

  // Report whether nvJPEG considers this image hardware-decodable, and how long a
  // decode takes on an idle stream: the engine path is several times faster than
  // the software one, so together these say which path the numbers below describe.
  {
    nvjpegJpegStream_t parsed;
    CHECK_NVJPEG(nvjpegJpegStreamCreate(handle, &parsed));
    CHECK_NVJPEG(nvjpegJpegStreamParseHeader(handle, data, length, parsed));
    int supported = -1;
    nvjpegDecodeBatchedSupported(handle, parsed, &supported);
    nvjpegJpegStreamDestroy(parsed);

    cudaEvent_t begin, end;
    CHECK_CUDA(cudaEventCreate(&begin));
    CHECK_CUDA(cudaEventCreate(&end));
    float best = 1e9f;
    for (int trial = 0; trial < 10; ++trial) {
      CHECK_CUDA(cudaEventRecord(begin, stream));
      decode(stream);
      CHECK_CUDA(cudaEventRecord(end, stream));
      CHECK_CUDA(cudaStreamSynchronize(stream));
      float elapsed = 0.0f;
      CHECK_CUDA(cudaEventElapsedTime(&elapsed, begin, end));
      best = elapsed < best ? elapsed : best;
    }
    CHECK_CUDA(cudaEventDestroy(begin));
    CHECK_CUDA(cudaEventDestroy(end));
    printf("batched_hw_supported=%d (0 means yes)  fastest decode %.2f ms\n", supported, best);
  }

  // The decoded image is non-zero, so a checksum of 0 below really does mean the
  // queued kernel ran before the decode wrote.
  unsigned long long decoded_sum = 0;
  CHECK_CUDA(cudaMemsetAsync(total, 0, sizeof(unsigned long long), stream));
  checksum<<<256, 256, 0, stream>>>(buffer, kBytes, total);
  CHECK_CUDA(cudaMemcpyAsync(&decoded_sum, total, sizeof(unsigned long long),
                             cudaMemcpyDeviceToHost, stream));
  CHECK_CUDA(cudaStreamSynchronize(stream));
  printf("checksum of a decoded frame: %llu\n", decoded_sum);

  int violations = 0;
  std::vector<double> call_ms;
  call_ms.reserve(iterations);
  for (int iteration = 0; iteration < iterations; ++iteration) {
    CHECK_CUDA(cudaMemsetAsync(buffer, 0, kBytes, stream));
    CHECK_CUDA(cudaMemsetAsync(total, 0, sizeof(unsigned long long), stream));
    CHECK_CUDA(cudaStreamSynchronize(stream));  // zeros are resident

    spin<<<1, 1, 0, stream>>>(150000000LL);          // ~100ms of stream time
    checksum<<<256, 256, 0, stream>>>(buffer, kBytes, total);  // queued read

    cudaStream_t submit = stream;
    if (barrier == "presync") {
      CHECK_CUDA(cudaStreamSynchronize(stream));  // the workaround: drain first
    } else if (submit_on_side) {
      // Recorded after the checksum, so the event is reached only once that read
      // has run, and `side` cannot legally run anything before then.
      CHECK_CUDA(cudaEventRecord(queued, stream));
      CHECK_CUDA(cudaStreamWaitEvent(side, queued, 0));
      if (barrier == "event-presync") {
        CHECK_CUDA(cudaStreamSynchronize(side));  // wait for the event on the host
      }
      submit = side;
    }

    auto begin = std::chrono::steady_clock::now();
    decode(submit);  // writes `buffer`
    call_ms.push_back(std::chrono::duration<double, std::milli>(
                          std::chrono::steady_clock::now() - begin)
                          .count());
    if (submit_on_side) {
      CHECK_CUDA(cudaStreamSynchronize(side));
    }

    unsigned long long seen = 0;
    CHECK_CUDA(cudaMemcpyAsync(&seen, total, sizeof(unsigned long long),
                               cudaMemcpyDeviceToHost, stream));
    CHECK_CUDA(cudaStreamSynchronize(stream));
    if (seen != 0) {
      if (violations == 0) {
        printf("first violation on iteration %d: queued read saw %llu, expected 0\n",
               iteration, seen);
      }
      ++violations;
    }
  }
  // The call time says where the host waited, if it waited at all: roughly the
  // spin's 100ms means the decode call itself blocked until the queued work was
  // done, and a fraction of that means it returned while that work was pending.
  std::sort(call_ms.begin(), call_ms.end());
  printf("out_of_order=%d/%d  decode call median %.1f ms\n", violations, iterations,
         call_ms.empty() ? 0.0 : call_ms[call_ms.size() / 2]);

  CHECK_CUDA(cudaFree(buffer));
  CHECK_CUDA(cudaFree(total));
  nvjpegJpegStateDestroy(state);
  nvjpegDestroy(handle);
  CHECK_CUDA(cudaEventDestroy(queued));
  CHECK_CUDA(cudaStreamDestroy(side));
  CHECK_CUDA(cudaStreamDestroy(stream));
  return violations == 0 ? 0 : 1;
}

Test

TestImageDecoder::test_cuda_jpeg_waits_for_callers_stream, the same shape as the reproducer. It
skips itself if the allocator never reused a freed block, so it cannot quietly become a test of
nothing.

Built from v0.16.0 with and without the barrier, same tree and same build configuration: the test
passes 5 runs out of 5 with it and fails 3 out of 3 without it. -k "jpeg or Jpeg" is otherwise
unchanged at 121 passed, with one pre-existing failure either way (test_cuda_jpeg_errors expects
nvjpegDecode failed: for corrupt input, which this GPU's hardware engine does not produce).

Scope

The hardware batched path only; the software path is unaffected, which the nvjpegDecode and
software-backend rows above show. An A100 on driver 535.129.03 was clean over 300 iterations while two
H100 machines reproduce, so architecture and driver version move together across the hardware I have
and I cannot say which one matters. The test passes vacuously on unaffected GPUs.

Related

• Fix CUDA decode_jpeg writing into memory the caller's stream is still using pytorch/vision#9606, the same symptom through torchvision.io.decode_jpeg. That decoder has an
  additional defect of its own — it decodes on a private stream that is never ordered against the
  caller's, so it corrupts on nvJPEG's software path too. This wrapper does not have that problem.
• Encode jpeg cuda sync pytorch/vision#8929, torchvision's encode-side stream synchronization fix. Out of scope here;
  EncodeJpegCuda.cpp follows the same caller's-stream-plus-host-sync design as decode, and whether
  nvJPEG's encode path has the same defect is untested.

[pytorch-bot]

🔗 Helpful Links

🧪 See artifacts and rendered test results at hud.pytorch.org/pr/meta-pytorch/torchcodec/1634

• 📄 Preview Python docs built from this PR

Note: Links to docs will display an error until the docs builds have been completed.

This comment was automatically generated by Dr. CI and updates every 15 minutes.

[meta-cla]

Hi @karkuspeter!

Thank you for your pull request and welcome to our community.

Action Required

In order to merge any pull request (code, docs, etc.), we require contributors to sign our Contributor License Agreement, and we don't seem to have one on file for you.

Process

In order for us to review and merge your suggested changes, please sign at https://code.facebook.com/cla. If you are contributing on behalf of someone else (eg your employer), the individual CLA may not be sufficient and your employer may need to sign the corporate CLA.

Once the CLA is signed, our tooling will perform checks and validations. Afterwards, the pull request will be tagged with CLA signed. The tagging process may take up to 1 hour after signing. Please give it that time before contacting us about it.

If you have received this in error or have any questions, please contact us at cla@meta.com. Thanks!

[NicolasHug]

Thanks @karkuspeter I can reproduce on an H100 (not on an A100, similar to yours).

I'll get to this eventually. In the mean time, out of curiosity is that a bug that was reported to NVJPEG?

[karkuspeter]

Thanks @NicolasHug great to hear you were able to reproduce.

I am preparing a bugreport / fix for nvjpeg as well