tensorflow/compiler/xla/service/gpu/tests/gpu_ldg_test.cc
// Parameters are never overwritten, so parameter reads should get ld.global.nc
// reads.
//
// 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
//
Parameters are never overwritten, so parameter reads should get ld.global.nc
reads.

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
Parameters are never overwritten so parameter reads should get ld . global . nc reads . 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, LdgForParamRead) {
  HloComputation::Builder builder(TestName());
  auto shape = ShapeUtil::MakeShape(F32, {2, 2});
  HloInstruction* param =
      builder.AddInstruction(HloInstruction::CreateParameter(0, shape, "x"));
  builder.AddInstruction(
      HloInstruction::CreateBinary(shape, HloOpcode::kAdd, param, param));
  std::unique_ptr<HloComputation> computation = builder.Build();
  auto hlo_module = CreateNewVerifiedModule();
  hlo_module->AddEntryComputation(std::move(computation));
  CompileAndOptionallyVerifyPtx(std::move(hlo_module), R"(
    CHECK-NOT: ld.global.f32
    CHECK: ld.global.nc.f32
  )");
}
TEST_F ( GpuLdgTest , LdgForParamRead ) { HloComputation :: Builder builder ( TestName ( ) ) ; auto shape = ShapeUtil :: MakeShape ( F32 , { 2 , 2 } ) ; HloInstruction * param = builder . AddInstruction ( HloInstruction :: CreateParameter ( 0 , shape , " " ) ) ; builder . AddInstruction ( HloInstruction :: CreateBinary ( shape , HloOpcode :: kAdd , param , param ) ) ; std :: unique_ptr < HloComputation > computation = builder . Build ( ) ; auto hlo_module = CreateNewVerifiedModule ( ) ; hlo_module -> AddEntryComputation ( std :: move ( computation ) ) ; CompileAndOptionallyVerifyPtx ( std :: move ( hlo_module ) , R" ( 
    CHECK-NOT: ld.global.f32
    CHECK: ld.global.nc.f32
   ) " ) ; }
Parameters are never overwritten, so parameter reads should get ld.global.nc
reads.
Parameters are never overwritten so parameter reads should get ld . global . nc reads .