Manually perform 8-bit dot product (__dp4a)

Hi,
The candle-kernel crate throws an error on the Maxwell architecture when migrating code that works fine in PyTorch.
It seems the only problem is '__dp4a' which used for performing four 8-bit integer dot product operations for quantization.
It would be great if the kernel crate has implementation to compute the dot product for '__CUDA_ARCH__ < 610'.

```cuda
#if __CUDA_ARCH__ < 610
// Manually perform 8-bit dot product
__device__ int manual_dp4a(int a, int b, int c) {
int result = c;
for (int i = 0; i < 4; ++i) {
int8_t a_byte = (a >> (i * 8)) & 0xFF;
int8_t b_byte = (b >> (i * 8)) & 0xFF;
result += a_byte * b_byte;
}
return result;
}
#endif

#if __CUDA_ARCH__ >= 610
...
sumi = __dp4a(vi0, u[2*i+0], sumi);
...
#else
...
sumi = manual_dp4a(vi0, u[2*i+0], sumi);
...
#endif
```

Start in the candle-kernel crate and locate the __dp4a call sites and their __CUDA_ARCH__ guards. Compare the existing path for newer architectures with the manual_dp4a example in this issue, then verify that Maxwell builds and computes the same quantized dot products without __dp4a.

Written by the indexing model from the issue text.