tvm-commits mailing list archives

Site index · List index
Message view « Date » · « Thread »
Top « Date » · « Thread »
From GitBox <...@apache.org>
Subject [GitHub] [incubator-tvm] masahi commented on issue #4983: [CodeGen][CUDA] Enhance CUDA codegen for SelectNode
Date Wed, 11 Mar 2020 02:16:55 GMT
masahi commented on issue #4983: [CodeGen][CUDA] Enhance CUDA codegen for SelectNode
URL: https://github.com/apache/incubator-tvm/pull/4983#issuecomment-597408983
 
 
   LGTM, but it seems indentation is broken in cuda source codegen. Not important, but would
be nice to clean it up.
   
   ```
   extern "C" __global__ void default_function_kernel0( float* __restrict__ B,  float* __restrict__
A) {
       float4 _1;
               int4 _2 = make_int4(37, 37, 37, 37);
               int4 _3 = make_int4(0, 0, 0, 0);
               ushort4 _4;
               _4.x = (_2.x>=_3.x);
               _4.y = (_2.y>=_3.y);
               _4.z = (_2.z>=_3.z);
               _4.w = (_2.w>=_3.w);
                 int4 _5 = (make_int4)(((((int)threadIdx.x) * 4))+(1*0), ((((int)threadIdx.x)
* 4))+(1*1), ((((int)threadIdx.x) * 4))+(1*2), ((((int)threadIdx.x) * 4))+(1*3));
                 int4 _6 = make_int4(37, 37, 37, 37);
                 int4 _7;
                 _7.x = (_5.x%_6.x);
                 _7.y = (_5.y%_6.y);
                 _7.z = (_5.z%_6.z);
                 _7.w = (_5.w%_6.w);
               int4 _8 = make_int4(0, 0, 0, 0);
               ushort4 _9;
               _9.x = (_7.x>=_8.x);
               _9.y = (_7.y>=_8.y);
               _9.z = (_7.z>=_8.z);
               _9.w = (_7.w>=_8.w);
             ushort4 _10;
             _10.x = (_4.x&&_9.x);
             _10.y = (_4.y&&_9.y);
             _10.z = (_4.z&&_9.z);
             _10.w = (_4.w&&_9.w);
               int4 _11 = make_int4(37, 37, 37, 37);
               int4 _12 = make_int4(0, 0, 0, 0);
               ushort4 _13;
               _13.x = (_11.x<_12.x);
               _13.y = (_11.y<_12.y);
               _13.z = (_11.z<_12.z);
               _13.w = (_11.w<_12.w);
                 int4 _14 = (make_int4)(((((int)threadIdx.x) * 4))+(1*0), ((((int)threadIdx.x)
* 4))+(1*1), ((((int)threadIdx.x) * 4))+(1*2), ((((int)threadIdx.x) * 4))+(1*3));
                 int4 _15 = make_int4(37, 37, 37, 37);
                 int4 _16;
                 _16.x = (_14.x%_15.x);
                 _16.y = (_14.y%_15.y);
                 _16.z = (_14.z%_15.z);
                 _16.w = (_14.w%_15.w);
               int4 _17 = make_int4(0, 0, 0, 0);
               ushort4 _18;
               _18.x = (_16.x<=_17.x);
               _18.y = (_16.y<=_17.y);
               _18.z = (_16.z<=_17.z);
               _18.w = (_16.w<=_17.w);
             ushort4 _19;
             _19.x = (_13.x&&_18.x);
             _19.y = (_13.y&&_18.y);
             _19.z = (_13.z&&_18.z);
             _19.w = (_13.w&&_18.w);
           ushort4 _20;
           _20.x = (_10.x||_19.x);
           _20.y = (_10.y||_19.y);
           _20.z = (_10.z||_19.z);
           _20.w = (_10.w||_19.w);
           int4 _21 = (make_int4)(((((int)threadIdx.x) * 4))+(1*0), ((((int)threadIdx.x) *
4))+(1*1), ((((int)threadIdx.x) * 4))+(1*2), ((((int)threadIdx.x) * 4))+(1*3));
           int4 _22 = make_int4(37, 37, 37, 37);
           int4 _23;
           _23.x = (_21.x/_22.x);
           _23.y = (_21.y/_22.y);
           _23.z = (_21.z/_22.z);
           _23.w = (_21.w/_22.w);
             int4 _24 = (make_int4)(((((int)threadIdx.x) * 4))+(1*0), ((((int)threadIdx.x)
* 4))+(1*1), ((((int)threadIdx.x) * 4))+(1*2), ((((int)threadIdx.x) * 4))+(1*3));
             int4 _25 = make_int4(37, 37, 37, 37);
             int4 _26;
             _26.x = (_24.x/_25.x);
             _26.y = (_24.y/_25.y);
             _26.z = (_24.z/_25.z);
             _26.w = (_24.w/_25.w);
           int4 _27 = make_int4(1, 1, 1, 1);
           int4 _28;
           _28.x = (_26.x-_27.x);
           _28.y = (_26.y-_27.y);
           _28.z = (_26.z-_27.z);
           _28.w = (_26.w-_27.w);
         int4 _29;
         _29.x = (bool(_20.x)?_23.x:_28.x);
         _29.y = (bool(_20.y)?_23.y:_28.y);
         _29.z = (bool(_20.z)?_23.z:_28.z);
         _29.w = (bool(_20.w)?_23.w:_28.w);
       _1.x = A[_29.x];
       _1.y = A[_29.y];
       _1.z = A[_29.z];
       _1.w = A[_29.w];
     (( float4*)(B + ((((int)threadIdx.x) * 4))))[0] = _1;
   }
   ```

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


With regards,
Apache Git Services

Mime
View raw message