tensorflow/compiler/xla/service/gpu/tests/gpu_ldg_test.cc
// Check that reading a buffer produced by a non-parameter HLO also results in
// ld.global.nc, if that buffer isn't modified within the instruction that reads
// it.
//
// 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 produced by a non-parameter HLO also results in
ld.global.nc, if that buffer isn't modified within the instruction that reads
it.

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 produced by a non - parameter HLO also results in ld . global . nc if that buffer isn ' t modified within the instruction that reads it . 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, LdgForNonParamRead) {
  HloComputation::Builder builder(TestName());
  auto shape = ShapeUtil::MakeShape(F32, {2, 2});
  HloInstruction* param =
      builder.AddInstruction(HloInstruction::CreateParameter(0, shape, "x"));
  HloInstruction* add = builder.AddInstruction(
      HloInstruction::CreateBinary(shape, HloOpcode::kAdd, param, param));
  HloInstruction* square = builder.AddInstruction(
      HloInstruction::CreateBinary(shape, HloOpcode::kMultiply, add, add));
  builder.AddInstruction(HloInstruction::CreateTuple({add, square}));
  std::unique_ptr<HloComputation> computation = builder.Build();
  auto hlo_module = CreateNewVerifiedModule();
  hlo_module->AddEntryComputation(std::move(computation));
  CompileAndOptionallyVerifyPtx(std::move(hlo_module), R"(
    CHECK: {
    CHECK-NOT: ld.global.f32
    CHECK: ld.global.nc.f32
    CHECK: }
  )");
}
TEST_F ( GpuLdgTest , LdgForNonParamRead ) { HloComputation :: Builder builder ( TestName ( ) ) ; auto shape = ShapeUtil :: MakeShape ( F32 , { 2 , 2 } ) ; HloInstruction * param = builder . AddInstruction ( HloInstruction :: CreateParameter ( 0 , shape , " " ) ) ; HloInstruction * add = builder . AddInstruction ( HloInstruction :: CreateBinary ( shape , HloOpcode :: kAdd , param , param ) ) ; HloInstruction * square = builder . AddInstruction ( HloInstruction :: CreateBinary ( shape , HloOpcode :: kMultiply , add , add ) ) ; builder . AddInstruction ( HloInstruction :: CreateTuple ( { add , square } ) ) ; std :: unique_ptr < HloComputation > computation = builder . Build ( ) ; auto hlo_module = CreateNewVerifiedModule ( ) ; hlo_module -> AddEntryComputation ( std :: move ( computation ) ) ; CompileAndOptionallyVerifyPtx ( std :: move ( hlo_module ) , R" ( 
    CHECK: {
    CHECK-NOT: ld.global.f32
    CHECK: ld.global.nc.f32
    CHECK: }
   ) " ) ; }
Check that reading a buffer produced by a non-parameter HLO also results in
ld.global.nc, if that buffer isn't modified within the instruction that reads
it.
Check that reading a buffer produced by a non - parameter HLO also results in ld . global . nc if that buffer isn ' t modified within the instruction that reads it .