3D Textures not working with uint8 type

I have edited the [test example](https://github.com/inducer/pycuda/blob/main/test/test_driver.py#L736) and changed the type from `float` to `uint8`, expecting to use it as `unsigned char` in the kernel. The output values are, however not the same, they seem to be bit-shifted or misaligned when being read from the memory. Below is the code, the expected result is that the printed values at the end are the same (simplified to copying of just the first value). Now I am getting the expected 20 and the wrong 161. I apologize if this is my lack of understanding but I haven't found anything regarding this topic in the docs and thus I consider this a bug.

```
import numpy as np
import pycuda.driver as drv
import pycuda.autoinit
from pycuda.compiler import SourceModule

shape = (2, 4, 8)

a = np.asarray(np.full(shape,20), dtype=np.uint8, order="F")
descr = drv.ArrayDescriptor3D()
descr.width = shape[0]
descr.height = shape[1]
descr.depth = shape[2]
descr.format = drv.dtype_to_array_format(a.dtype)
descr.num_channels = 1
descr.flags = 0
ary = drv.Array(descr)
copy = drv.Memcpy3D()
copy.set_src_host(a)
copy.set_dst_array(ary)
copy.width_in_bytes = copy.src_pitch = a.strides[1]
copy.src_height = copy.height = shape[1]
copy.depth = shape[2]
copy()

mod = SourceModule(
"""
texture mtx_tex;

__global__ void copy_texture(unsigned char *dest)
{
if(threadIdx.x + threadIdx.y + threadIdx.z == 0)
dest[0] = tex3D(mtx_tex, 0, 0, 0);
}
"""
)
copy_texture = mod.get_function("copy_texture")
mtx_tex = mod.get_texref("mtx_tex")
mtx_tex.set_array(ary)
dest = np.zeros(shape, dtype=np.uint8, order="F")
copy_texture(drv.Out(dest), block=shape, texrefs=[mtx_tex])
print(a[0][0][0])
print(dest[0][0][0])
```
This seems to be related to the memory alignment rules. When the descriptor is set to `uint32` (`descr.format = drv.dtype_to_array_format(np.uint32)`) instead of the input array type, the right values can be obtained when accessed as `unsigned char` in the kernel. This is a workaround, or a correct solution? In case of 4 8bit channels, it would be necessary to pass the values reinterpreted as an array of integers?