[nvCOMP 5.3] ANS FLOAT8_E4M3 single-byte decompression leaves device_statuses unchanged

[KonataSpring]

Summary

With nvCOMP 5.3, nvcompBatchedANSDecompressAsync leaves device_statuses[0] unchanged when decoding a chunk with one uncompressed byte, compressed using NVCOMP_TYPE_FLOAT8_E4M3.

The API launch and CUDA synchronization succeed, the output byte matches the input, and the actual decompressed size is 1. However, a status buffer initialized to 0x55555555 still contains that sentinel afterward, instead of nvcompSuccess.

This is a standalone C++ reproducer using only nvCOMP, CUDA Runtime, and the standard library. No framework, model weights, custom allocator, or workspace pool is involved.

Environment

• GPU: NVIDIA GeForce RTX 3070 Laptop GPU (compute capability 8.6)
• NVIDIA driver reported by nvidia-smi: 617.14
• Linux container: Ubuntu 22.04.5 LTS, running through WSL2
• nvCOMP 5.3; nvcompGetProperties: version=5300, cudart_version=13030
• CUDA toolkit compiler: nvcc V13.0.88 (distinct from the nvCOMP CUDA build version above)
• GCC: 11.4.0

Only this GPU/environment has been tested; I am not claiming the issue affects every architecture.

Expected behavior

After successful decompression and synchronization, each supplied output status should be nvcompSuccess, as described in the ANS C API documentation.

Actual behavior and reproducibility

The program below queries the required alignments and scratch bounds, verifies buffer alignment, checks compression status, synchronizes before reading results, and initializes decode status to a recognizable sentinel.

A counted run of 10 separate process invocations produced:

• 40/40 E4M3 single-byte cases: correct output and actual size, but status remains 0x55555555.
• 320/320 control cases: correct output and status nvcompSuccess.
• A separate compute-sanitizer --tool memcheck run reproduced the issue and reported ERROR SUMMARY: 0 errors.

The four configurations per input combine:

• Explicit decompression dtype vs. CHAR/default bitstream autodetection.
• Default decompression backend vs. explicit NVCOMP_DECOMPRESS_BACKEND_CUDA.

Controls use E4M3 chunks of 2 and 33 bytes; CHAR chunks of 1, 2 and 33 bytes; and FLOAT16 chunks of 2, 4 and 66 bytes.

Representative output (n is the uncompressed byte count, not the compressed size; type 10 is FLOAT8_E4M3, type 0 is CHAR):

nvcomp_version=5300 nvcomp_cudart_version=13030
type=10 n=1 auto=0 cuda=0 comp_bytes=72 temp=0 status=0x55555555 actual=1 equal=1 align=8/1/1
type=10 n=1 auto=0 cuda=1 comp_bytes=72 temp=0 status=0x55555555 actual=1 equal=1 align=8/1/1
type=10 n=1 auto=1 cuda=0 comp_bytes=72 temp=0 status=0x55555555 actual=1 equal=1 align=8/1/1
type=10 n=1 auto=1 cuda=1 comp_bytes=72 temp=0 status=0x55555555 actual=1 equal=1 align=8/1/1
type=10 n=2 auto=0 cuda=0 comp_bytes=560 temp=0 status=0x00000000 actual=2 equal=1 align=8/1/1
type=0 n=1 auto=0 cuda=0 comp_bytes=304 temp=0 status=0x00000000 actual=1 equal=1 align=8/1/1

Standalone reproducer

Save as ans_standalone.cpp. Adjust the library installation paths if needed:

g++ -std=c++17 -O2 ans_standalone.cpp \
  -I/opt/nvcomp/include -I/usr/local/cuda/include \
  -L/opt/nvcomp/lib -L/usr/local/cuda/lib64 \
  -Wl,-rpath,/opt/nvcomp/lib -Wl,-rpath,/usr/local/cuda/lib64 \
  -lnvcomp -lcudart -o ans_standalone

./ans_standalone
for i in {1..10}; do ./ans_standalone || break; done
compute-sanitizer --tool memcheck --error-exitcode 99 ./ans_standalone

The program prints status values rather than treating the unchanged status as an exit failure, so all controls can execute.

Complete C++ source

#include <cuda_runtime_api.h>
#include <nvcomp/ans.h>
#include <nvcomp.h>
#include <cstdio>
#include <cstring>
#include <stdexcept>
#include <vector>
#include <algorithm>
#include <cstdint>

void cu(cudaError_t s) { if (s != cudaSuccess) throw std::runtime_error(cudaGetErrorString(s)); }
void nv(nvcompStatus_t s) { if (s != nvcompSuccess) throw std::runtime_error("nvCOMP API launch/query failed: " + std::to_string(s)); }
struct Buffer {
  void* p = nullptr;
  explicit Buffer(size_t n) { if (n) cu(cudaMalloc(&p, n)); }
  ~Buffer() { if (p) cudaFree(p); }
  Buffer(const Buffer&) = delete;
};
bool aligned(void* p, size_t a) { return !p || reinterpret_cast<std::uintptr_t>(p) % a == 0; }

void run(nvcompType_t type, size_t bytes, bool autodetect, bool explicit_cuda) {
  auto compress = nvcompBatchedANSCompressDefaultOpts;
  compress.data_type = type;
  auto decompress = nvcompBatchedANSDecompressDefaultOpts;
  decompress.data_type = autodetect ? NVCOMP_TYPE_CHAR : type;
  if (explicit_cuda) decompress.backend = NVCOMP_DECOMPRESS_BACKEND_CUDA;
  size_t compression_temp = 0, decompression_temp = 0, bound = 0;
  nvcompAlignmentRequirements_t ca{}, da{};
  nv(nvcompBatchedANSCompressGetRequiredAlignments(compress, &ca));
  nv(nvcompBatchedANSDecompressGetRequiredAlignments(decompress, &da));
  nv(nvcompBatchedANSCompressGetTempSizeAsync(1, bytes, compress, &compression_temp, bytes));
  nv(nvcompBatchedANSCompressGetMaxOutputChunkSize(bytes, compress, &bound));
  nv(nvcompBatchedANSDecompressGetTempSizeAsync(1, bytes, decompress, &decompression_temp, bytes));
  Buffer input(bytes), encoded(bound), decoded(bytes);
  Buffer compress_temp(compression_temp), decode_temp(decompression_temp);
  Buffer input_table(sizeof(void*)), encoded_table(sizeof(void*)), output_table(sizeof(void*));
  Buffer input_size(sizeof(size_t)), encoded_size(sizeof(size_t)), actual_size(sizeof(size_t));
  Buffer compress_status(sizeof(nvcompStatus_t)), decode_status(sizeof(nvcompStatus_t));
  if (!aligned(input.p, ca.input) || !aligned(encoded.p, ca.output) || !aligned(compress_temp.p, ca.temp) ||
      !aligned(encoded.p, da.input) || !aligned(decoded.p, da.output) || !aligned(decode_temp.p, da.temp))
    throw std::runtime_error("invalid buffer alignment");
  std::vector<unsigned char> original(bytes), restored(bytes);
  for (size_t i = 0; i < bytes; ++i) original[i] = static_cast<unsigned char>(i * 37 + 0x38);
  cu(cudaMemcpy(input.p, original.data(), bytes, cudaMemcpyHostToDevice));
  cu(cudaMemcpy(input_table.p, &input.p, sizeof(void*), cudaMemcpyHostToDevice));
  cu(cudaMemcpy(encoded_table.p, &encoded.p, sizeof(void*), cudaMemcpyHostToDevice));
  cu(cudaMemcpy(output_table.p, &decoded.p, sizeof(void*), cudaMemcpyHostToDevice));
  cu(cudaMemcpy(input_size.p, &bytes, sizeof(size_t), cudaMemcpyHostToDevice));
  cu(cudaMemset(compress_status.p, 0x55, sizeof(nvcompStatus_t)));
  nv(nvcompBatchedANSCompressAsync(static_cast<const void* const*>(input_table.p),
      static_cast<const size_t*>(input_size.p), bytes, 1, compress_temp.p, compression_temp,
      static_cast<void* const*>(encoded_table.p), static_cast<size_t*>(encoded_size.p), compress,
      static_cast<nvcompStatus_t*>(compress_status.p), nullptr));
  cu(cudaDeviceSynchronize());
  unsigned int cs = 0, ds = 0;
  size_t encoded_bytes = 0, actual_bytes = 0;
  cu(cudaMemcpy(&cs, compress_status.p, sizeof(cs), cudaMemcpyDeviceToHost));
  cu(cudaMemcpy(&encoded_bytes, encoded_size.p, sizeof(encoded_bytes), cudaMemcpyDeviceToHost));
  if (cs != 0 || encoded_bytes > bound) throw std::runtime_error("compression failed");
  cu(cudaMemset(decode_status.p, 0x55, sizeof(nvcompStatus_t)));
  cu(cudaMemset(actual_size.p, 0, sizeof(size_t)));
  cu(cudaMemset(decoded.p, 0xa5, bytes));
  nv(nvcompBatchedANSDecompressAsync(static_cast<const void* const*>(encoded_table.p),
      static_cast<const size_t*>(encoded_size.p), static_cast<const size_t*>(input_size.p),
      static_cast<size_t*>(actual_size.p), 1, decode_temp.p, decompression_temp,
      static_cast<void* const*>(output_table.p), decompress, static_cast<nvcompStatus_t*>(decode_status.p), nullptr));
  cu(cudaDeviceSynchronize());
  cu(cudaMemcpy(&ds, decode_status.p, sizeof(ds), cudaMemcpyDeviceToHost));
  cu(cudaMemcpy(&actual_bytes, actual_size.p, sizeof(actual_bytes), cudaMemcpyDeviceToHost));
  cu(cudaMemcpy(restored.data(), decoded.p, bytes, cudaMemcpyDeviceToHost));
  std::printf("type=%d n=%zu auto=%d cuda=%d comp_bytes=%zu temp=%zu status=0x%08x actual=%zu equal=%d align=%zu/%zu/%zu\n",
      int(type), bytes, autodetect, explicit_cuda, encoded_bytes, decompression_temp, ds, actual_bytes,
      restored == original, da.input, da.output, da.temp);
}
int main() {
  try {
    nvcompProperties_t properties{};
    nv(nvcompGetProperties(&properties));
    std::printf("nvcomp_version=%d nvcomp_cudart_version=%d\n", properties.version, properties.cudart_version);
    for (auto type : {NVCOMP_TYPE_FLOAT8_E4M3, NVCOMP_TYPE_CHAR, NVCOMP_TYPE_FLOAT16})
      for (size_t n : {size_t(1), size_t(2), size_t(33)})
        for (bool autodetect : {false, true})
          for (bool explicit_cuda : {false, true})
            run(type, type == NVCOMP_TYPE_FLOAT16 ? n * 2 : n, autodetect, explicit_cuda);
  } catch (const std::exception& e) { std::fprintf(stderr, "%s\n", e.what()); return 1; }
}

Additional observation / question

In a separate variant, calling nvcompBatchedANSDecompressGetTempSizeSync with the same status buffer first writes nvcompSuccess; that value remains successful after decoding. This does not establish that DecompressAsync itself writes the success status, so I am not presenting pre-initializing success as a correctness fix.

Could you confirm whether this is a status-write omission in the single-byte E4M3 path, or whether the reproducer violates an undocumented requirement? In particular, is the synchronous query required before this otherwise valid Async-query/DecompressAsync sequence?

Thank you.