Shawn-Inspur commented on a change in pull request #5600:
URL: https://github.com/apache/incubator-tvm/pull/5600#discussion_r427832441
##########
File path: topi/python/topi/cuda/softmax.py
##########
@@ -53,13 +54,62 @@ def schedule_softmax(outs):
raise ValueError('Tag is expected to be softmax_output or
log_softmax_output. \
Got {0}'.format(op_tag))
+ # The nvptx backend only supports 32-bits warp shuffle instructions.
+ #
+ # TODO(tvm-team) Fix nvptx codegen or deprecate nvptx backend.
+ def sched_warp_softmax():
+ if tgt.target_name == "nvptx":
+ return softmax.dtype == "float32" or softmax.dtype == "int32"
+ return True
+
if len(softmax.shape) > 2:
ops = [max_elem.op, expsum.op, softmax.op]
if exp is not None:
ops.append(exp.op)
for op in ops:
s = schedule_injective_from_existing(s, op.output(0))
+
+ elif sched_warp_softmax():
+ # A warp of 32 threads performs a row reduction.
+ num_thread = tgt.thread_warp_size
+ block_x = te.thread_axis("blockIdx.x")
+ thread_x = te.thread_axis((0, num_thread), "threadIdx.x")
+
+ # (4) softmax
+ xo, xi = s[softmax].split(softmax.op.axis[1], nparts=num_thread)
+ if tgt.target_name != "nvptx":
+ _, xii = s[softmax].split(xi, factor=4)
+ s[softmax].vectorize(xii)
+ s[softmax].bind(xo, thread_x)
+ s[softmax].bind(softmax.op.axis[0], block_x)
+
+ # (3) expsum
+ k = expsum.op.reduce_axis[0]
+ ko, _ = s[expsum].split(k, nparts=num_thread)
+ s[expsum].bind(ko, thread_x)
+ s[expsum].compute_at(s[softmax], xo)
+
+ # (2) exp
+ if exp is not None:
+ xo, xi = s[exp].split(exp.op.axis[1], nparts=num_thread)
+ _, xii = s[exp].split(xi, factor=4)
+ s[exp].vectorize(xii)
Review comment:
I am wondering whether the vectorizing of “xii” should be included in an
if block, which checked "tgt.target_name != "nvptx", for the logic here looks
the same as presented in line 81.
----------------------------------------------------------------
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]
