()
| 57 | @pytest.mark.gpu |
| 58 | @pytest.mark.skipif(not env.has_cuda(), reason="need cuda") |
| 59 | def test_cuda_vectorize_add(): |
| 60 | num_thread = 8 |
| 61 | |
| 62 | def check_cuda(dtype, n, lanes): |
| 63 | if dtype == "float16" and not have_fp16(tvm.cuda(0).compute_version): |
| 64 | print("Skip because gpu does not have fp16 support") |
| 65 | return |
| 66 | if dtype == "int8" and not have_int8(tvm.cuda(0).compute_version): |
| 67 | print("skip because gpu does not support int8") |
| 68 | return |
| 69 | vec_dtype = f"{dtype}x{lanes}" |
| 70 | one = tvm.tirx.const(1, vec_dtype) |
| 71 | num_blocks = (n + num_thread - 1) // num_thread |
| 72 | |
| 73 | @I.ir_module(s_tir=True) |
| 74 | class Module: |
| 75 | @T.prim_func(s_tir=True) |
| 76 | def main(A: T.Buffer((n,), vec_dtype), B: T.Buffer((n,), vec_dtype)): |
| 77 | T.func_attr({"tirx.noalias": True}) |
| 78 | for i_0 in T.thread_binding(num_blocks, thread="blockIdx.x"): |
| 79 | for i_1 in T.thread_binding(num_thread, thread="threadIdx.x"): |
| 80 | with T.sblock("B"): |
| 81 | v_i = T.axis.spatial(n, i_0 * num_thread + i_1) |
| 82 | T.where(i_0 * num_thread + i_1 < n) |
| 83 | T.reads(A[v_i]) |
| 84 | T.writes(B[v_i]) |
| 85 | B[v_i] = A[v_i] + one |
| 86 | |
| 87 | fun = tvm.compile(Module, target="cuda") |
| 88 | |
| 89 | dev = tvm.cuda(0) |
| 90 | a = tvm.runtime.empty((n,), vec_dtype, dev).copyfrom(np.random.uniform(size=(n, lanes))) |
| 91 | c = tvm.runtime.empty((n,), vec_dtype, dev) |
| 92 | fun(a, c) |
| 93 | tvm.testing.assert_allclose(c.numpy(), a.numpy() + 1) |
| 94 | |
| 95 | check_cuda("float32", 64, 2) |
| 96 | check_cuda("float32", 64, 3) |
| 97 | check_cuda("float32", 64, 4) |
| 98 | check_cuda("int8", 64, 2) |
| 99 | check_cuda("int8", 64, 3) |
| 100 | check_cuda("int8", 64, 4) |
| 101 | check_cuda("uint8", 64, 2) |
| 102 | check_cuda("uint8", 64, 3) |
| 103 | check_cuda("uint8", 64, 4) |
| 104 | check_cuda("float16", 64, 2) |
| 105 | check_cuda("float16", 64, 4) |
| 106 | check_cuda("float16", 64, 6) |
| 107 | check_cuda("float16", 64, 8) |
| 108 | |
| 109 | |
| 110 | @pytest.mark.gpu |
nothing calls this directly
no test coverage detected
searching dependent graphs…