fix matmul kernel and tests
MODIFY
accel/opencl/matmul.cl
accel/opencl/matmul.cl
__kernel void matmul(
write_only image2d_t output,
__local float *outputScratch,
- read_only image1d_t input,
+ read_only image2d_t input,
read_only image2d_t weights
//ARGS
) {
@@ -18,7 +18,7 @@ __kernel void matmul(
// fast path precompute (32x speedup)
float outputValue = 0.0f;
for (short inputSet = (short)get_global_id(1); inputSet < numPackedInputChannelsForGroup; inputSet += get_global_size(1)) {
- float4 inputValues = read_imagef(input, smp, inputSet);
+ float4 inputValues = read_imagef(input, smp, (int2)(inputSet, 0));
float4 weightValues = read_imagef(weights, smp, (int2)(mad24(inputSet, 4, weightIndex), packedOutputChannel));
outputValue += dot(inputValues, weightValues);
}
@@ -36,24 +36,8 @@ __kernel void matmul(
scratchIndex += 4;
}
- // slow path
- /*for (short packedInputChannel = 0; packedInputChannel < numPackedInputChannelsForGroup; ++packedInputChannel) {
- float4 inputValues = read_imagef(input, smp, packedInputChannel);
-
- float4 weightValues[4];
- for (short outChIdx = 0; outChIdx < 4; ++outChIdx) {
- weightValues[outChIdx] = read_imagef(weights, smp, (int2)(packedInputChannel*4 + outChIdx, packedOutputChannel));
- }
-
- outputValues.x += dot(inputValues, weightValues[0]);
- outputValues.y += dot(inputValues, weightValues[1]);
- outputValues.z += dot(inputValues, weightValues[2]);
- outputValues.w += dot(inputValues, weightValues[3]);
- }*/
-
// insert unary and binary ops here
- int2 outputLocation;
- outputLocation.x = packedOutputChannel;
+ int2 outputLocation = (int2)(packedOutputChannel, 0);
//BINOP
// output to memory

MODIFY
test/test_ops.py
test/test_ops.py
@@ -18,7 +18,7 @@ def helper_test_op(shps, torch_fxn, tinygrad_fxn, atol=1e-6, rtol=1e-3, grad_ato
ret = tinygrad_fxn(*tst)
def compare(s, x,y,atol,rtol):
- assert x.shape == y.shape
+ if y.shape != tuple(): assert x.shape == y.shape, f"shape mismatch {x.shape} != {y.shape}"
try:
np.testing.assert_allclose(x,y, atol=atol, rtol=rtol)
except Exception:

fix matmul kernel and tests