AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Allow creating texture from unaligned user ptr

Browse Source 
All we have to do is align the ptr to HW requirments an if it's not zero, then return the offset to the user.

We currently don't have anywhere to store this offset, so hipGetTextureAlignmentOffset() will still always return 0.

Change-Id: If31998127d99a2a3222a026d88249519d6102505
This commit is contained in:
Vladislav Sytchenko
2020-03-19 16:05:06 -04:00
parent 8e8b372726
commit ea701777d8
1 changed files with 43 additions and 21 deletions
Download Patch File Download Diff File Expand all files Collapse all files
vdi/hip_texture.cpp
+43 -21
View File
@@ -382,6 +382,28 @@ hipError_t hipGetTextureObjectTextureDesc(hipTextureDesc* pTexDesc,
HIP_RETURN(hipSuccess);
}
inline bool ihipGetTextureAlignmentOffset(size_t* offset,
const void* devPtr) {
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
const char* alignedDevPtr = amd::alignUp(static_cast<const char*>(devPtr), info.imageBaseAddressAlignment_);
const size_t alignedOffset = alignedDevPtr - static_cast<const char*>(devPtr);
// If the device memory pointer was returned from hipMalloc(),
// the offset is guaranteed to be 0 and NULL may be passed as the offset parameter.
if ((alignedOffset != 0) &&
(offset == nullptr)) {
return false;
}
if (offset != nullptr) {
*offset = alignedOffset;
}
return true;
}
hipError_t ihipBindTexture(size_t* offset,
const textureReference* texref,
const void* devPtr,
@@ -398,19 +420,19 @@ hipError_t ihipBindTexture(size_t* offset,
// No need to check for errors.
ihipDestroyTextureObject(texref->textureObject);
// If the device memory pointer was returned from hipMalloc(),
// the offset is guaranteed to be 0 and NULL may be passed as the offset parameter.
// TODO enforce alignment on devPtr.
if (offset != nullptr) {
*offset = 0;
}
hipResourceDesc resDesc = {};
resDesc.resType = hipResourceTypeLinear;
resDesc.res.linear.devPtr = const_cast<void*>(devPtr);
resDesc.res.linear.desc = *desc;
resDesc.res.linear.sizeInBytes = size;
if (ihipGetTextureAlignmentOffset(offset, devPtr)) {
// Align the user ptr to HW requirments.
resDesc.res.linear.devPtr = static_cast<char*>(const_cast<void*>(devPtr)) - *offset;
} else {
return hipErrorInvalidValue;
}
hipTextureDesc texDesc = hip::getTextureDesc(texref);
return ihipCreateTextureObject(const_cast<hipTextureObject_t*>(&texref->textureObject), &resDesc, &texDesc, nullptr);
@@ -434,13 +456,6 @@ hipError_t ihipBindTexture2D(size_t* offset,
// No need to check for errors.
ihipDestroyTextureObject(texref->textureObject);
// If the device memory pointer was returned from hipMalloc(),
// the offset is guaranteed to be 0 and NULL may be passed as the offset parameter.
// TODO enforce alignment on devPtr.
if (offset != nullptr) {
*offset = 0;
}
hipResourceDesc resDesc = {};
resDesc.resType = hipResourceTypePitch2D;
resDesc.res.pitch2D.devPtr = const_cast<void*>(devPtr);
@@ -449,6 +464,13 @@ hipError_t ihipBindTexture2D(size_t* offset,
resDesc.res.pitch2D.height = height;
resDesc.res.pitch2D.pitchInBytes = pitch;
if (ihipGetTextureAlignmentOffset(offset, devPtr)) {
// Align the user ptr to HW requirments.
resDesc.res.pitch2D.devPtr = static_cast<char*>(const_cast<void*>(devPtr)) - *offset;
} else {
return hipErrorInvalidValue;
}
hipTextureDesc texDesc = hip::getTextureDesc(texref);
return ihipCreateTextureObject(const_cast<hipTextureObject_t*>(&texref->textureObject), &resDesc, &texDesc, nullptr);
@@ -798,19 +820,19 @@ hipError_t hipTexRefSetAddress(size_t* ByteOffset,
// No need to check for errors.
ihipDestroyTextureObject(texRef->textureObject);
// If the device memory pointer was returned from hipMemAlloc(),
// the offset is guaranteed to be 0 and NULL may be passed as the ByteOffset parameter.
// TODO enforce alignment on devPtr.
if (ByteOffset != nullptr) {
*ByteOffset = 0;
}
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;
if (ihipGetTextureAlignmentOffset(ByteOffset, dptr)) {
// Align the user ptr to HW requirments.
resDesc.res.linear.devPtr = static_cast<char*>(dptr) - *ByteOffset;
} else {
return HIP_RETURN(hipErrorInvalidValue);
}
hipTextureDesc texDesc = hip::getTextureDesc(texRef);
HIP_RETURN(ihipCreateTextureObject(&texRef->textureObject, &resDesc, &texDesc, nullptr));