tensorflow/compiler/xla/service/gpu/tests/gpu_ldg_test.cc
// Check that reading a buffer that's modified in-place does not produce
// ld.global.nc.
//
// We do this by creating a reduce that feeds into a sin.  We don't currently
// fuse sin into reduce, and the sin is elementwise, so it reuses its input
// buffer as its output.
//
// It seems like a fair bet that we won't start fusing sin into the output of
// reduce in the foreseeable future.  But if that turns out to be wrong, I give
// you, future reader, permission to delete this test.
//
// On the ROCM platform the "ptx" string is not populated for the compiled
// executable, and hence the call to CompileAdnVerifyPtx does not do the
// "VerifyPtx" part, it merely compiles the executable
//
Check that reading a buffer that's modified in-place does not produce
ld.global.nc.

We do this by creating a reduce that feeds into a sin.  We don't currently
fuse sin into reduce, and the sin is elementwise, so it reuses its input
buffer as its output.

It seems like a fair bet that we won't start fusing sin into the output of
reduce in the foreseeable future.  But if that turns out to be wrong, I give
you, future reader, permission to delete this test.

On the ROCM platform the "ptx" string is not populated for the compiled
executable, and hence the call to CompileAdnVerifyPtx does not do the
"VerifyPtx" part, it merely compiles the executable
Check that reading a buffer that ' s modified in - place does not produce ld . global . nc . We do this by creating a reduce that feeds into a sin . We don ' t currently fuse sin into reduce and the sin is elementwise so it reuses its input buffer as its output . It seems like a fair bet that we won ' t start fusing sin into the output of reduce in the foreseeable future . But if that turns out to be wrong I give you future reader permission to delete this test . On the ROCM platform the " ptx " string is not populated for the compiled executable and hence the call to CompileAdnVerifyPtx does not do the " VerifyPtx " part it merely compiles the executable
TEST_F(GpuLdgTest, NoLdgWhenSharingBuffer) {
  auto hlo_module = CreateNewVerifiedModule();
  HloComputation::Builder builder(TestName());
  HloComputation* reduce_computation;
  {
    auto embedded_builder = HloComputation::Builder("add");
    auto lhs = embedded_builder.AddInstruction(HloInstruction::CreateParameter(
        0, ShapeUtil::MakeShape(F32, {}), "lhs"));
    auto rhs = embedded_builder.AddInstruction(HloInstruction::CreateParameter(
        1, ShapeUtil::MakeShape(F32, {}), "rhs"));
    embedded_builder.AddInstruction(
        HloInstruction::CreateBinary(lhs->shape(), HloOpcode::kAdd, lhs, rhs));
    reduce_computation =
        hlo_module->AddEmbeddedComputation(embedded_builder.Build());
  }
  auto param_shape = ShapeUtil::MakeShape(F32, {32, 32});
  auto reduce_shape = ShapeUtil::MakeShape(F32, {32});
  HloInstruction* param = builder.AddInstruction(
      HloInstruction::CreateParameter(0, param_shape, "x"));
  HloInstruction* reduce = builder.AddInstruction(HloInstruction::CreateReduce(
      reduce_shape,
      builder.AddInstruction(HloInstruction::CreateBinary(
          param_shape, HloOpcode::kAdd, param, param)),
      builder.AddInstruction(
          HloInstruction::CreateConstant(LiteralUtil::CreateR0<float>(0))),
      {0}, reduce_computation));
  builder.AddInstruction(
      HloInstruction::CreateUnary(reduce_shape, HloOpcode::kSin, reduce));
  std::unique_ptr<HloComputation> computation = builder.Build();
  hlo_module->AddEntryComputation(std::move(computation));
  CompileAndOptionallyVerifyPtx(std::move(hlo_module), R"(
    CHECK-LABEL: .entry sin
    CHECK: {
    CHECK-NOT: ld.global.nc.f32
    CHECK: ld.global.f32
    CHECK: }
  )");
}
TEST_F ( GpuLdgTest , NoLdgWhenSharingBuffer ) { auto hlo_module = CreateNewVerifiedModule ( ) ; HloComputation :: Builder builder ( TestName ( ) ) ; HloComputation * reduce_computation ; { auto embedded_builder = HloComputation :: Builder ( " " ) ; auto lhs = embedded_builder . AddInstruction ( HloInstruction :: CreateParameter ( 0 , ShapeUtil :: MakeShape ( F32 , { } ) , " " ) ) ; auto rhs = embedded_builder . AddInstruction ( HloInstruction :: CreateParameter ( 1 , ShapeUtil :: MakeShape ( F32 , { } ) , " " ) ) ; embedded_builder . AddInstruction ( HloInstruction :: CreateBinary ( lhs -> shape ( ) , HloOpcode :: kAdd , lhs , rhs ) ) ; reduce_computation = hlo_module -> AddEmbeddedComputation ( embedded_builder . Build ( ) ) ; } auto param_shape = ShapeUtil :: MakeShape ( F32 , { 32 , 32 } ) ; auto reduce_shape = ShapeUtil :: MakeShape ( F32 , { 32 } ) ; HloInstruction * param = builder . AddInstruction ( HloInstruction :: CreateParameter ( 0 , param_shape , " " ) ) ; HloInstruction * reduce = builder . AddInstruction ( HloInstruction :: CreateReduce ( reduce_shape , builder . AddInstruction ( HloInstruction :: CreateBinary ( param_shape , HloOpcode :: kAdd , param , param ) ) , builder . AddInstruction ( HloInstruction :: CreateConstant ( LiteralUtil :: CreateR0 < float > ( 0 ) ) ) , { 0 } , reduce_computation ) ) ; builder . AddInstruction ( HloInstruction :: CreateUnary ( reduce_shape , HloOpcode :: kSin , reduce ) ) ; std :: unique_ptr < HloComputation > computation = builder . Build ( ) ; hlo_module -> AddEntryComputation ( std :: move ( computation ) ) ; CompileAndOptionallyVerifyPtx ( std :: move ( hlo_module ) , R" ( 
    CHECK-LABEL: .entry sin
    CHECK: {
    CHECK-NOT: ld.global.nc.f32
    CHECK: ld.global.f32
    CHECK: }
   ) " ) ; }
Check that reading a buffer that's modified in-place does not produce
ld.global.nc.
Check that reading a buffer that ' s modified in - place does not produce ld . global . nc .