[GitHub] [tvm] Laurawly commented on a change in pull request #7561: [Codegen][CUDA] Fix: cuda codegen vectorize cast

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

Laurawly commented on a change in pull request #7561:
URL: https://github.com/apache/tvm/pull/7561#discussion_r585287887



##########
File path: src/target/source/codegen_cuda.cc
##########
@@ -238,12 +253,53 @@ void CodeGenCUDA::PrintType(DataType t, std::ostream& os) {  // NOLINT(*)
           break;
         }
       }
-      case 16:
-        os << "short";
-        break;
-      case 32:
-        os << "int";
+      case 16: {
+        if (t.is_scalar()) {
+          os << "short";
+        } else if (t.lanes() <= 4) {
+          os << "short" << lanes;
+        } else if (t.lanes() <= 8) {
+          // Emit CUDA code to access int16 vector elements.
+          //
+          // short4 is stored as int2
+          //
+          // s4.x is emitted as *(short2*)(&(i2.x)).x
+          // s4.y is emitted as *(short2*)(&(i2.x)).y
+          // s4.z is emitted as *(short2*)(&(i2.y)).x
+          // s4.w is emitted as *(short2*)(&(i2.y)).y
+          //
+          ICHECK_EQ(t.lanes() % 2, 0) << "only support even lane for shorT type with lanes > 4";
+          os << "int" << t.lanes() / 2;
+        } else {
+          fail = true;
+        }
+        if (!fail) {
+          return;
+        }

Review comment:
       Missing `break` here?

##########
File path: tests/python/unittest/test_target_codegen_cuda.py
##########
@@ -511,8 +511,8 @@ def check(t0, t1):
 
         # schedule
         s = tvm.te.create_schedule(C.op)
-        ob, ib = s[C].split(s[C].op.axis[0], nparts=32)
-        _, iib = s[C].split(ib, factor=4)
+        ob, ib = s[C].split(s[C].op.axis[0], nparts=n // factor)

Review comment:
       We can also directly say `factor=factor` here.




----------------------------------------------------------------
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