in openai_gemm.py [0:0]
def matmul_test(ng, dtype, op, m, n, k, ones=False, out=False):
prefix = "s" if dtype is np.float32 else "h"
if op[0] == 'T':
vec4A = (m & 3) == 0
vec8A = (m & 7) == 0
dimA = (k,m)
cda = m
else:
vec4A = (k & 3) == 0
vec8A = (k & 7) == 0
dimA = (m,k)
cda = k
if op[1] == 'T':
vec4B = (k & 3) == 0
vec8B = (k & 7) == 0
dimB = (n,k)
cdb = k
else:
vec4B = (n & 3) == 0
vec8B = (n & 7) == 0
dimB = (k,n)
cdb = n
vec4C = (n & 3) == 0
dimC = (m,n)
cdc = n
A1 = ng.empty(dimA, dtype=dtype)
B1 = ng.empty(dimB, dtype=dtype)
C1 = ng.empty(dimC, dtype=dtype)
C2 = ng.empty(dimC, dtype=dtype)
if ones:
A1[:] = 1.0
B1[:] = 1.0
else:
# fill with uniform randoms from -1 to 1
A1[:] = 2 * (.5 - ng.rand())
B1[:] = 2 * (.5 - ng.rand())
# for reducing outputs
partial1 = ng.empty((C1.shape[0],1), dtype=np.float32)
partial2 = partial1[0:1,0:1]
cublas_handle = _get_cublas()
for tileM, tileN, tileK, vecA, vecB, vecC, div, base_op, dyn_shared in selections[prefix][op]:
vecA = (vecA == 4 and vec4A) or (vecA == 8 and vec8A)
vecB = (vecB == 4 and vec4B) or (vecB == 8 and vec8B)
vecC = vecC == 1 or vec4C
vec = vecA and vecB and vecC
if base_op:
# The op is part of the base kernel name
base = "%sgemm_%dx%dx%d_%s" % (prefix, tileM, tileN, tileK, op)
opts = ( "vec", ) if vec else ()
else:
# The op is an option passed to a more generic kernel
base = "%sgemm_%dx%dx%d" % (prefix, tileM, tileN, tileK)
opts = ( op, "vec" ) if vec else (op,)
kernel = get_kernel(base, opts)
blk_A = _ceil_div(m, tileM)
blk_B = _ceil_div(n, tileN)
blk_a, blk_A = _closest_divisor(blk_A, div)
blk_b, blk_B = _closest_divisor(blk_B, div)
if blk_a == 1:
blk_a, blk_A = (blk_A, 1)
for alpha, beta in ( (1.0,0.0), (0.5,0.5) ):
try:
if ones:
C1[:] = 1.0
else:
C1[:] = 2 * (.5 - ng.rand())
C2[:] = C1
params = [
(blk_a * blk_b, blk_B, blk_A), (kernel.threads, 1, 1), None,
C1.gpudata, A1.gpudata, B1.gpudata, alpha, beta,
cda, cdb, cdc, m, n, k, blk_a, blk_b ]
kernel.prepared_async_call(*params)
# convert row order to col order
cublasXgemm[prefix](cublas_handle, op[1], op[0], n, m, k, alpha, B1.gpudata, cdb, A1.gpudata, cda, beta, C2.gpudata, cdc)
# Check for NaNs
partial1[:] = ng.min(ng.finite(C1), axis=1)
partial2[:] = ng.min(partial1, axis=0)
if partial2.get()[0,0] == 0.0:
print "Error: NaN kernel: %s mnk: (%d,%d,%d) ab: (%f,%f)" % (kernel.name, m,n,k, alpha,beta)
exit()
# Get Max Diff
partial1[:] = ng.max(abs(C2 - C1), axis=1)
partial2[:] = ng.max(partial1, axis=0)
diff = partial2.get()[0,0]
# Get Mean
partial1[:] = ng.sum(abs(C2), axis=1)
partial2[:] = ng.sum(partial1, axis=0)
mean = partial2.get()[0,0] / C2.size
# Scale diff by the mean
pctErr = 100 * diff / mean
#print "Error: %.3f %s" % (pctErr, kernel.name)
maxerr = .005 if dtype is np.float32 else 0.7
if pctErr > maxerr:
print "Error: %.3f%% diff: %.5f mean %.5f kernel: %s mnk: (%d,%d,%d) ab: (%f,%f)" % (pctErr, diff, mean, kernel.name, m,n,k, alpha,beta)
print params
if out:
C1 = C1.get()
C2 = C2.get()
D = abs(C2 - C1)
np.savetxt("out_diff.txt", D, fmt='%3.1f')
np.savetxt("out_correct.txt", C2, fmt='%5.1f')
np.savetxt("out_error", C1, fmt='%5.1f')
exit()
except drv.Error as e:
print "kernel: %s mnk: (%d,%d,%d) ab: (%f,%f)" % (kernel.name, m,n,k, alpha,beta)
print e
exit()