test/cuda/test_tc_mapper_bugs.cc
// This test exercises the unrolling and reduction matching which used to fail
This test exercises the unrolling and reduction matching which used to fail
This test exercises the unrolling and reduction matching which used to fail
TEST_F(TensorDot_32_512_8_2_28_28, ReductionUnroll) {
  Init();
  auto options =
      tc::CudaMappingOptions::makeNaiveCudaMappingOptions()
          .outerScheduleFusionStrategy(tc::FusionStrategy::Preserve3Coincident)
          .outerScheduleAllowSkewing(false)
          .outerSchedulePositiveOrthant(true)
          .intraTileScheduleFusionStrategy(tc::FusionStrategy::Max)
          .intraTileScheduleAllowSkewing(false)
          .intraTileSchedulePositiveOrthant(true)
          .tile(28, 14, 7, 32, 28, 256)
          .mapToBlocks(1)
          .mapToThreads(7, 28)
          .unroll(64)
          .tileImperfectlyNested(false)
          .useSharedMemory(true)
          .usePrivateMemory(false)
          .unrollCopyShared(true)
          .matchLibraryCalls(true);
  Check(options
            .tile(28, 14, 7, 32, 28, 256)
            .mapToThreads(7, 28));
}
TEST_F ( TensorDot_32_512_8_2_28_28 , ReductionUnroll ) { Init ( ) ; auto options = tc :: CudaMappingOptions :: makeNaiveCudaMappingOptions ( ) . outerScheduleFusionStrategy ( tc :: FusionStrategy :: Preserve3Coincident ) . outerScheduleAllowSkewing ( false ) . outerSchedulePositiveOrthant ( true ) . intraTileScheduleFusionStrategy ( tc :: FusionStrategy :: Max ) . intraTileScheduleAllowSkewing ( false ) . intraTileSchedulePositiveOrthant ( true ) . tile ( 28 , 14 , 7 , 32 , 28 , 256 ) . mapToBlocks ( 1 ) . mapToThreads ( 7 , 28 ) . unroll ( 64 ) . tileImperfectlyNested ( false ) . useSharedMemory ( true ) . usePrivateMemory ( false ) . unrollCopyShared ( true ) . matchLibraryCalls ( true ) ; Check ( options . tile ( 28 , 14 , 7 , 32 , 28 , 256 ) . mapToThreads ( 7 , 28 ) ) ; }
This test exercises the unrolling and reduction matching which used to fail
This test exercises the unrolling and reduction matching which used to fail
// Tiling by 28 (problem size after scheduling) and setting // threadIdx.y to 28 seems to work fine. //                tile for y //                   ~~ //              map y //               ~~