[GitHub] [tvm] echuraev commented on a diff in pull request #14397: [OpenCL][Textures] Always use SSA for texture loading

["echuraev (via GitHub)" <gi...@apache.org>]

echuraev commented on code in PR #14397:
URL: https://github.com/apache/tvm/pull/14397#discussion_r1150134595


##########
src/target/source/codegen_opencl.cc:
##########
@@ -472,20 +445,15 @@ void CodeGenOpenCL::VisitExpr_(const CallNode* op, std::ostream& os) {
     this->PrintExpr(op->args[2], ss);
     ss << ")))";
 
-    // Only use local SSA if texture is not already being stored
-    if (need_texture_ssa_) {
-      std::string rhs = SSAGetID(ss.str(), op->dtype.with_lanes(4));
-      if (op->args.back().as<RampNode>()) {
-        os << rhs;
-      } else {
-        os << "((";
-        this->PrintType(op->dtype.with_lanes(1), os);
-        os << "*)&" << rhs << ")[";
-        this->PrintExpr(op->args.back(), os);
-        os << "]";
-      }

Review Comment:
   Thank you for your review! I have added several new test cases. They cover case when the data was storage and local variable is reused. Also add a new one which check that writing `float [4]` to `image2d_t` will be handled correctly:
   ```c++
   float out_local[4];
   // ...
   write_imagef(out, (int2)((((((int)get_group_id(0)) * 14) + ((int)get_local_id(0))) % 38), (((((int)get_group_id(0)) * 64) + (((int)get_local_id(0)) >> 1)) / 19)), vload4(0, out_local + 0));
   ```



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

To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org

For queries about this service, please contact Infrastructure at:
users@infra.apache.org