Move pk_i4 permute outside, and enable permuteB for B.

example/01_gemm/gemm_xdl_fp16_pk_i4_v3_b_scale.cpp

@@ -22,7 +22,7 @@ using CElementOp = PassThrough;
 
 static constexpr auto GemmDefault = ck::tensor_operation::device::GemmSpecialization::Default;
 
-static constexpr bool PermuteB = false;
+static constexpr bool PermuteB = true;
 
 static constexpr ck::index_t Scale_Block_N = 1;
 static constexpr ck::index_t Scale_Block_K = 128;
@@ -199,7 +199,7 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config)
     }
 
     // vector pk_i4x4 permute
-#if 0
+#if 1
     for(int i = 0; i < N; i++)
     {
         for(int j = 0; j < K; j += 8)

include/ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp

@@ -86,7 +86,7 @@ struct PassThroughPack8
 
     __host__ __device__ constexpr void operator()(ck::half8_t& y, const ck::pk_i4x4_t& x) const
     {
-#if 1
+#if 0
         int x_permute = 0;
         int bits4_0   = (bit_cast<int>(x) >> 0) & 0xF;
         int bits4_1   = (bit_cast<int>(x) >> 4) & 0xF;