wpan11nv commented on a change in pull request #5226: [CODEGEN][CUDA] Fix
vector load
URL: https://github.com/apache/incubator-tvm/pull/5226#discussion_r404293617
##########
File path: tests/python/unittest/test_target_codegen_cuda.py
##########
@@ -543,6 +543,40 @@ def run_test(dtype):
run_test("uint32")
run_test("uint64")
+def test_cuda_vectorize_load_permute_pad():
+ def check_cuda(dtype, n, l, padding, lanes):
+ if not tvm.gpu(0).exist or not tvm.runtime.enabled("cuda"):
+ print("skip because cuda is not enabled..")
+ return
+
+ ctx = tvm.gpu(0)
+ A = tvm.te.placeholder((n, l), name='A', dtype=dtype)
+ B = tvm.te.compute((n // lanes, l + 2 * padding, lanes),
+ lambda i, j, k: tvm.te.if_then_else(
+ tvm.te.any(j < padding, j >= l + padding),
+ tvm.runtime.convert(0).astype(dtype), A[i * lanes + k, j -
padding]),
+ name='B')
+ s = te.create_schedule(B.op)
+ block, thread, vectorize = s[B].op.axis
+ s[B].bind(block, bx)
+ s[B].bind(thread, tx)
+ s[B].vectorize(vectorize)
+ fun = tvm.build(s, [A, B], "cuda", name="vector_load_permute_pad")
+ np_a = np.random.randint(
+ low=-128, high=127, size=(n, l)).astype(A.dtype)
+ a = tvm.nd.empty((n, l), A.dtype, ctx).copyfrom(np_a)
+ b = tvm.nd.empty((n // lanes, l + padding * 2, lanes), B.dtype, ctx)
+ fun(a, b)
+ np_a_reshape = np_a.reshape(n // lanes, lanes, l).transpose(0, 2, 1)
+ ref = np.pad(np_a_reshape, ((0, 0), (padding, padding),
+ (0, 0)), mode='constant',
constant_values=0)
+ tvm.testing.assert_allclose(b.asnumpy(), ref)
+
+ check_cuda("int8", 64, 16, 3, 4)
Review comment:
uint8 test?
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
[email protected]
With regards,
Apache Git Services
