AI/rocm-systems
2
0
Fork 0
程式碼 問題 合併請求 Actions 套件 專案 版本發布 Wiki 動態

P4 to Git Change 2039536 by jujiang@JJ-HIP on 2019/12/02 14:31:13

瀏覽代碼 
SWDEV-214490 - Update HIP RT for texture3D in HIP/PAL on Windows
	-Update ihipBindTexture

	http://ocltc.amd.com/reviews/r/18333/

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#89 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_texture.cpp#28 edit
此提交包含在：
foreman
2019-12-02 13:37:35 -06:00
父節點 5349bd8036
當前提交 69dcd2b862
共有 2 個檔案被更改，包括 66 行新增 和 38 行删除
下載 Patch 檔 下載差異檔 展開所有檔案 折疊所有檔案
api/hip/hip_memory.cpp
+4 -35
查看文件
@@ -38,6 +38,8 @@ extern void getDrvChannelOrderAndType(const enum hipArray_Format Format,
extern void setDescFromChannelType(cl_channel_type channelType, hipChannelFormatDesc* desc);
extern void getByteSizeFromChannelFormatKind(enum hipChannelFormatKind channelFormatKind, size_t* byteSize);
amd::Memory* getMemoryObject(const void* ptr, size_t& offset) {
amd::Memory *memObj = amd::MemObjMap::FindMemObj(ptr);
if (memObj != nullptr) {
@@ -881,24 +883,7 @@ hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, con
amd::HostQueue* queue = hip::getNullStream();
size_t dpitch = dst->width;
switch (dst[0].desc.f) {
case hipChannelFormatKindSigned:
dpitch *= sizeof(int);
break;
case hipChannelFormatKindUnsigned:
dpitch *= sizeof(unsigned int);
break;
case hipChannelFormatKindFloat:
dpitch *= sizeof(float);
break;
case hipChannelFormatKindNone:
dpitch *= sizeof(size_t);
break;
default:
dpitch *= 1;
break;
}
getByteSizeFromChannelFormatKind(dst[0].desc.f, &dpitch);
if ((wOffset + width > (dpitch)) || width > spitch) {
HIP_RETURN(hipErrorInvalidDevicePointer);
@@ -1133,23 +1118,7 @@ hipError_t hipMemcpy3D(const struct hipMemcpy3DParms* p) {
size_t dstOrigin[3];
size_t region[3];
if (p->dstArray != nullptr) {
switch (p->dstArray->desc.f) {
case hipChannelFormatKindSigned:
byteSize = sizeof(int);
break;
case hipChannelFormatKindUnsigned:
byteSize = sizeof(unsigned int);
break;
case hipChannelFormatKindFloat:
byteSize = sizeof(float);
break;
case hipChannelFormatKindNone:
byteSize = sizeof(size_t);
break;
default:
byteSize = 1;
break;
}
getByteSizeFromChannelFormatKind(p->dstArray->desc.f, &byteSize);
region[2] = p->extent.depth;
region[1] = p->extent.height;
region[0] = p->extent.width;
api/hip/hip_texture.cpp
+62 -3
查看文件
@@ -193,6 +193,27 @@ void getChannelOrderAndType(const hipChannelFormatDesc& desc, enum hipTextureRea
}
}
void getByteSizeFromChannelFormatKind(enum hipChannelFormatKind channelFormatKind, size_t* byteSize) {
switch (channelFormatKind)
{
case hipChannelFormatKindSigned:
*byteSize = sizeof(int);
break;
case hipChannelFormatKindUnsigned:
*byteSize = sizeof(unsigned int);
break;
case hipChannelFormatKindFloat:
*byteSize = sizeof(float);
break;
case hipChannelFormatKindNone:
*byteSize = sizeof(size_t);
break;
default:
*byteSize = 1;
break;
}
}
amd::Sampler* fillSamplerDescriptor(enum hipTextureAddressMode addressMode,
enum hipTextureFilterMode filterMode, int normalizedCoords) {
#ifndef CL_FILTER_NONE
@@ -401,14 +422,33 @@ hipError_t ihipBindTexture(cl_mem_object_type type,
}
if (hip::getCurrentContext()) {
cl_image_format image_format;
size_t byteSize;
size_t rowPitch = 0;
size_t depth = 0;
size_t slicePitch = 0;
getChannelOrderAndType(*desc, hipReadModeElementType,
&image_format.image_channel_order, &image_format.image_channel_data_type);
getByteSizeFromChannelFormatKind(desc->f, &byteSize);
const amd::Image::Format imageFormat(image_format);
amd::Memory* memory = getMemoryObject(devPtr, *offset);
amd::Image* image = new (*hip::getCurrentContext()) amd::Image(*memory->asBuffer(),
type, memory->getMemFlags(), imageFormat, width, height, 1, pitch, 0);
switch (type) {
case CL_MEM_OBJECT_IMAGE3D:
rowPitch = width * byteSize;
depth = pitch;
slicePitch = rowPitch * height;
break;
case CL_MEM_OBJECT_IMAGE2D:
default:
rowPitch = pitch;
depth = 1;
slicePitch = 0;
break;
}
amd::Image* image = new (*hip::getCurrentContext()) amd::Image(*memory->asBuffer(),
type, memory->getMemFlags(), imageFormat, width, height, depth, rowPitch, slicePitch);
if (!image->create()) {
delete image;
return hipErrorMemoryAllocation;
@@ -437,6 +477,19 @@ hipError_t ihipBindTexture(cl_mem_object_type type,
resDesc.res.pitch2D.height = height;
resDesc.res.pitch2D.pitchInBytes = pitch;
break;
case CL_MEM_OBJECT_IMAGE3D:
resDesc.resType = hipResourceTypeArray;
resDesc.res.array.array = (hipArray*)malloc(sizeof(hipArray));
resDesc.res.array.array->desc = *desc;
resDesc.res.array.array->width = width;
resDesc.res.array.array->height = height;
resDesc.res.array.array->depth = depth;
resDesc.res.array.array->Format = tex->format;
resDesc.res.array.array->NumChannels = tex->numChannels;
resDesc.res.array.array->isDrv = false;
resDesc.res.array.array->textureType = hipTextureType3D;
resDesc.res.array.array->data = const_cast<void*>(devPtr);
break;
default:
resDesc.resType = hipResourceTypeArray;
resDesc.res.array.array = nullptr;
@@ -444,7 +497,10 @@ hipError_t ihipBindTexture(cl_mem_object_type type,
}
tex->textureObject = reinterpret_cast<hipTextureObject_t>(ihipCreateTextureObject(resDesc, *image, *sampler));
if(type == CL_MEM_OBJECT_IMAGE3D) {
free(resDesc.res.array.array);
}
memset(&resDesc, 0, sizeof(hipResourceDesc));
return hipSuccess;
}
return hipErrorInvalidValue;
@@ -508,6 +564,9 @@ hipError_t ihipBindTextureToArrayImpl(TlsData* tls, int dim, enum hipTextureRead
case 2:
clType = CL_MEM_OBJECT_IMAGE2D;
break;
case 3:
clType = CL_MEM_OBJECT_IMAGE3D;
break;
default:
HIP_RETURN(hipErrorInvalidValue);
}