[Interpreter][histogram] Fix silent data corruption

[mieshkiwrk]

There's silent data corruption when calling tl.histogram with interpreter.

# test.py
import torch
import ctypes
import triton
import triton.language as tl


@triton.jit
def histogram_kernel(x_ptr, z_ptr):
    offset = tl.arange(0, 1)
    x = tl.load(x_ptr + offset)
    z = tl.histogram(x, 1)
    buf = (ctypes.c_int32 * 2).from_address(int(z_ptr))

    print(f'before store: {list(buf)}')
    tl.store(z_ptr + offset, z) # tl.store treats z values as int64 while they're int32
    print(f'after store: {list(buf)}')


device = 'cpu'
torch.manual_seed(17)
x = torch.ones(1, device=device, dtype=torch.int32)
z = torch.ones(2, dtype=torch.int32, device=device)
histogram_kernel[(1, )](x, z)

# Output:
# TRITON_INTERPRET=1 TRITON_TEST_SUITE=interpreter python test.py 
# before store: [1, 1]
# after store: [1, 0] <- second element shouldn't be cleared

Based on np.histogram docs: https://numpy.org/doc/2.3/reference/generated/numpy.histogram.html
Returned dtype is taken account when optional weights param is passed, int64 othwerwise.
That leads to tl.store thinking it's saving int32 values while there's int64 in my example tensor passed, so it's writing 8 bytes at once instead of 4 bytes, leading to writing 4 bytes exceeding it's data range causing silent data corruption.

import numpy as np

data = np.array([1], dtype=np.int32)
bins = 1

print(f'Data dtype before: {data.dtype}')
histogram = np.histogram(data, bins=bins, range=(0, bins))[0]
print(f'Data dtype after: {histogram.dtype}')

# Data dtype before: int32                                                                                                                                           
# Data dtype after: int64

Applying "dummy_weights" fixes returned data type as expected fixing data corruption.

[Copilot]

Pull Request Overview

This PR fixes silent data corruption in the interpreter's tl.histogram implementation. The issue occurred because np.histogram returns int64 values by default, but when storing int32 tensors, tl.store would write 8 bytes instead of 4, corrupting adjacent memory.

Key changes:

• Added dummy weights parameter to np.histogram to preserve the input data type
• Added detailed comment explaining the fix

---

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

[Copilot]

[nitpick] Add a comment explaining that the dummy weights preserve the input dtype in the histogram output, as this is a subtle behavior of np.histogram that may not be immediately obvious to future maintainers.

Suggested change

	# tl.store to write 8 bytes instead of 4 bytes which lead to silent data corruption
	# tl.store to write 8 bytes instead of 4 bytes which lead to silent data corruption
	# Note: Providing dummy_weights with the same dtype as the input ensures that
	# np.histogram returns a histogram with the same dtype as the weights.
	# This is a subtle behavior of np.histogram; without weights, the output dtype defaults to int64.

[anmyachev]

Great job!