[GitHub] [tvm] wejoncy opened a new issue #7247: [BUG] Generate float type of index for vectorize schedule

[GitBox <gi...@apache.org>]

wejoncy opened a new issue #7247:
URL: https://github.com/apache/tvm/issues/7247


   ## Description
   
   TVM will generate float type of  index expr, like `B_local[((((float)get_local_id(1)) + ((float)(((int)get_group_id(2)) * 8))))];` , which is illegal.
   
   ## Error information
   ```
   <kernel>:41:227: error: array subscript is not an integer
       B[((((float)((((((int)get_group_id(0)) * 131072) + (yy_inner * 16384)) + (((int)get_group_id(1)) * 4096)) + (((int)get_local_id(0)) * 512))) + (((float)get_local_id(1)) + ((float)(((int)get_group_id(2)) * 8)))))] = B_local[((((float)get_local_id(1)) + ((float)(((int)get_group_id(2)) * 8))))];
   ```
   
   ## Env
   TVM::main
   
   ## Minimal code to reproduce.
   ```
   
   import numpy as np
   import tvm
   from tvm import te
   
   # The sizes of inputs and filters
   batch = 1
   in_channel = 256
   out_channel = 512
   in_size = 32
   kernel = 1
   pad = 0
   stride = 1
   
   # Algorithm
   A = te.placeholder((in_size, in_size, in_channel, batch), name="A")
   W = te.placeholder((kernel, kernel, in_channel, out_channel), name="W")
   out_size = (in_size - kernel + 2 * pad) // stride + 1
   # Pad input
   Apad = te.compute(
       (in_size + 2 * pad, in_size + 2 * pad, in_channel, batch),
       lambda yy, xx, cc, nn: tvm.tir.if_then_else(
           tvm.tir.all(yy >= pad, yy - pad < in_size, xx >= pad, xx - pad < in_size),
           A[yy - pad, xx - pad, cc, nn],
           tvm.tir.const(0.0, "float32"),
       ),
       name="Apad",
   )
   # Create reduction variables
   rc = te.reduce_axis((0, in_channel), name="rc")
   ry = te.reduce_axis((0, kernel), name="ry")
   rx = te.reduce_axis((0, kernel), name="rx")
   # Compute the convolution
   B = te.compute(
       (out_size, out_size, out_channel, batch),
       lambda yy, xx, ff, nn: te.sum(
           Apad[yy * stride + ry, xx * stride + rx, rc, nn] * W[ry, rx, rc, ff], axis=[ry, rx, rc]
       ),
       name="B",
   )
   
   # Designate the memory hierarchy
   s = te.create_schedule(B.op)
   s[Apad].compute_inline()  # compute Apad inline
   AA = s.cache_read(Apad, "shared", [B])
   WW = s.cache_read(W, "shared", [B])
   AL = s.cache_read(AA, "local", [B])
   WL = s.cache_read(WW, "local", [B])
   BL = s.cache_write(B, "local")
   
   # tile consts
   tile = 1
   num_thread = 8
   block_factor = tile * num_thread
   step = 8
   vthread = 2
   
   # Get the GPU thread indices
   block_x = te.thread_axis("blockIdx.x")
   block_y = te.thread_axis("blockIdx.y")
   block_z = te.thread_axis("blockIdx.z")
   thread_x = te.thread_axis((0, num_thread), "threadIdx.x")
   thread_y = te.thread_axis((0, num_thread/4), "threadIdx.y")
   thread_xz = te.thread_axis((0, vthread), "vthread", name="vx")
   thread_yz = te.thread_axis((0, vthread), "vthread", name="vy")
   
   # Split the workloads
   hi, wi, fi, ni = s[B].op.axis
   bz, fi = s[B].split(fi, factor=block_factor)
   bx, hi = s[B].split(hi, factor=block_factor)
   by, wi = s[B].split(wi, factor=block_factor)
   # Bind the iteration variables to GPU thread indices
   s[B].bind(bz, block_z)
   s[B].bind(by, block_y)
   s[B].bind(bx, block_x)
   
   
   ty, fi = s[B].split(fi, nparts=num_thread)
   tx, wi = s[B].split(wi, nparts=num_thread)
   
   s[B].reorder(bz,by,bx,ty,wi, fi)
   
   
   s[B].bind(ty, thread_y)
   s[B].bind(tx, thread_x)
   
   
   # Schedule BL local write
   s[BL].compute_at(s[B],ni)
   fi, noo = s[B].split(fi, factor=4)
   s[B].vectorize(noo)  # vectorize memory load
   yi, xi, fi, ni = s[BL].op.axis
   ry, rx, rc = s[BL].op.reduce_axis
   rco, rci = s[BL].split(rc, factor=step)
   s[BL].reorder(rco, ry, rx, rci, fi, ni)
   
   # Attach computation to iteration variables
   s[AA].compute_at(s[BL], rx)
   s[WW].compute_at(s[BL], rx)
   s[AL].compute_at(s[BL], rci)
   s[WL].compute_at(s[BL], rci)
   
   yi, xi, ci, ni = s[AA].op.axis
   tx, ni = s[AA].split(ni, nparts=num_thread)
   _, ci = s[AA].split(ci, factor=4)
   
   s[AA].bind(tx, thread_x)
   s[AA].vectorize(ci)  # vectorize memory load
   
   # Schedule for W's shared memory load
   yi,xi,ci,fi=s[WW].op.axis
   ty, ci = s[WW].split(ci, nparts=num_thread)
   fi, fv = s[WW].split(fi, factor=4)
   s[WW].bind(ty, thread_y)
   s[WW].vectorize(fv)  # vectorize memory load
   
   target="opencl"
   func = tvm.build(s, [A, W, B], target)
   print("------opencl code------")
   print(func.imported_modules[0].get_source()) if len(func.imported_modules) > 0 else print("source not imported")
   ctx = tvm.context(target, 0)
   np.random.seed(5)
   a_np = np.random.uniform(size=(in_size, in_size, in_channel, batch)).astype("float32")
   w_np = np.random.uniform(size=(kernel, kernel, in_channel, out_channel)).astype("float32")
   
   a = tvm.nd.array(a_np, ctx)
   w = tvm.nd.array(w_np, ctx)
   b = tvm.nd.array(np.zeros((out_size, out_size, out_channel, batch), dtype="float32"), ctx)
   func(a, w, b)
   np.savetxt("filename.txt",b.asnumpy()[:,:,0,0])
   evaluator = func.time_evaluator(func.entry_name, ctx, number=5)
   print("Convolution: %f ms" % (evaluator(a, w, b).mean * 1e3))
   ```
   


----------------------------------------------------------------
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:
users@infra.apache.org

[GitHub] [tvm] wejoncy closed issue #7247: [BUG]( OPENCL Code-generate) Generate float type of index for local memory

[GitBox <gi...@apache.org>]

wejoncy closed issue #7247:
URL: https://github.com/apache/tvm/issues/7247


   


----------------------------------------------------------------
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:
users@infra.apache.org

[GitHub] [tvm] wejoncy commented on issue #7247: [BUG]( OPENCL Code-generate) Generate float type of index for local memory

[GitBox <gi...@apache.org>]

wejoncy commented on issue #7247:
URL: https://github.com/apache/tvm/issues/7247#issuecomment-758630626


   
   Have root cause the issue, `te.thread_axis` do not check if dom type is int. So add a cast operation in `te.thread_axis`  could resolve it.


----------------------------------------------------------------
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:
users@infra.apache.org