Function hipTexRefSetAddress

hipamd/src/hip_texture.cpp:1055–1096

}

hipError_t hipTexRefSetAddress(size_t* ByteOffset, textureReference* texRef, hipDeviceptr_t dptr,
                               size_t bytes) {
  HIP_INIT_API(hipTexRefSetAddress, ByteOffset, texRef, dptr, bytes);

  if (texRef == nullptr) {
    HIP_RETURN(hipErrorInvalidValue);
  }

  hipDeviceptr_t refDevPtr = nullptr;
  size_t refDevSize = 0;
  HIP_RETURN_ONFAIL(PlatformState::instance().getDynTexGlobalVar(texRef, &refDevPtr, &refDevSize));
  assert(refDevSize == sizeof(textureReference));

  // Any previous address or HIP array state associated with the texture reference is superseded by
  // this function. Any memory previously bound to hTexRef is unbound. No need to check for errors.
  hipError_t err = ihipDestroyTextureObject(texRef->textureObject);
  if (err != hipSuccess) {
    HIP_RETURN(err);
  }

  hipResourceDesc resDesc = {};
  resDesc.resType = hipResourceTypeLinear;
  resDesc.res.linear.devPtr = dptr;
  resDesc.res.linear.desc = hip::getChannelFormatDesc(texRef->numChannels, texRef->format);
  resDesc.res.linear.sizeInBytes = bytes;
  err = ihipGetTextureAlignmentOffset(ByteOffset, dptr);
  if (err != hipSuccess) {
    HIP_RETURN(err);
  }

  // Align the user ptr to HW requirments.
  resDesc.res.linear.devPtr = static_cast<char*>(dptr) - (ByteOffset ? *ByteOffset : 0);
  hipTextureDesc texDesc = hip::getTextureDesc(texRef);

  err = ihipCreateTextureObject(&texRef->textureObject, &resDesc, &texDesc, nullptr);
  if (err != hipSuccess) {
    HIP_RETURN(err);
  }
  // Copy to device.
  hip::Stream* stream = hip::getNullStream();
  HIP_RETURN(ihipMemcpy(refDevPtr, texRef, refDevSize, hipMemcpyHostToDevice, *stream));
}

hipError_t hipTexRefSetAddress2D(textureReference* texRef, const HIP_ARRAY_DESCRIPTOR* desc,
                                 hipDeviceptr_t dptr, size_t Pitch) {