diff --git a/projects/clr/hipamd/include/hip/hcc_detail/driver_types.h b/projects/clr/hipamd/include/hip/hcc_detail/driver_types.h index 0c29542c7e..1941f44617 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/driver_types.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/driver_types.h @@ -255,14 +255,14 @@ typedef struct hipMemcpy3DParms { hipArray_t srcArray; struct hipPos srcPos; struct hipPitchedPtr srcPtr; - hipArray_t dstArray; struct hipPos dstPos; struct hipPitchedPtr dstPtr; - struct hipExtent extent; enum hipMemcpyKind kind; +} hipMemcpy3DParms; +typedef struct HIP_MEMCPY3D { size_t Depth; size_t Height; size_t WidthInBytes; @@ -283,10 +283,7 @@ typedef struct hipMemcpy3DParms { size_t srcLOD; hipMemoryType srcMemoryType; size_t srcPitch; - size_t srcXInBytes; - size_t srcY; - size_t srcZ; -}hipMemcpy3DParms; +} HIP_MEMCPY3D; static inline struct hipPitchedPtr make_hipPitchedPtr(void* d, size_t p, size_t xsz, size_t ysz) { diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 444e41107a..d965059923 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -1540,111 +1540,144 @@ hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t return ihipLogStatus(e); } +int getByteSizeFromFormat(const hipChannelFormatDesc& desc){ + int byteSize =0; + switch (desc.f) { + case hipChannelFormatKindUnsigned: + switch (desc.x) { + case 32: + byteSize = sizeof(uint32_t); + break; + case 16: + byteSize = sizeof(uint16_t); + break; + case 8: + byteSize = sizeof(uint8_t); + break; + default: + byteSize = sizeof(uint32_t); + } + break; + case hipChannelFormatKindSigned: + switch (desc.x) { + case 32: + byteSize = sizeof(int32_t); + break; + case 16: + byteSize = sizeof(int16_t); + break; + case 8: + byteSize = sizeof(int8_t); + break; + default: + byteSize = sizeof(int32_t); + } + break; + case hipChannelFormatKindFloat: + switch (desc.x) { + case 32: + byteSize = sizeof(float); + break; + case 16: + byteSize = sizeof(_Float16); + break; + default: + byteSize = sizeof(float); + } + break; + case hipChannelFormatKindNone: + default: + break; + } + return byteSize; +} + hipError_t ihipMemcpy3D(const struct hipMemcpy3DParms* p, hipStream_t stream, bool isAsync) { hipError_t e = hipSuccess; if(p) { - size_t byteSize, width, height, depth, widthInBytes, srcPitch, dstPitch, ySize; - hipChannelFormatDesc desc; - void* srcPtr;void* dstPtr; + size_t dstByteSize, srcByteSize, copyWidth, copyHeight, copyDepth, widthInBytes, srcPitch, dstPitch, srcYsize, dstYsize; + size_t srcXoffset, srcYoffset, srcZoffset, dstXoffset, dstYoffset, dstZoffset; + size_t srcWidth, srcHeight, srcDepth, dstWidth, dstHeight, dstDepth; + + void* srcPtr, *dstPtr; + bool copyWidthUpdate= false; + copyDepth = p->extent.depth; + copyHeight = p->extent.height; + copyWidth = p->extent.width; // in bytes ? + dstXoffset = p->dstPos.x; + dstYoffset = p->dstPos.y; + dstZoffset = p->dstPos.z; + srcXoffset = p->srcPos.x; + srcYoffset = p->srcPos.y; + srcZoffset = p->srcPos.z; if (p->dstArray != nullptr) { - if (p->dstArray->isDrv == false) { - 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 = 0; - break; - } - depth = p->extent.depth; - height = p->extent.height; - width = p->extent.width; - widthInBytes = p->extent.width * byteSize; - srcPitch = p->srcPtr.pitch; - srcPtr = p->srcPtr.ptr; - ySize = p->srcPtr.ysize; - desc = p->dstArray->desc; - dstPtr = p->dstArray->data; - hsa_ext_image_data_info_t imageInfo; - if(hipTextureType2DLayered == p->dstArray->textureType) - GetImageInfo(HSA_EXT_IMAGE_GEOMETRY_2DA, width, height, 0, desc, imageInfo, depth); - else - GetImageInfo(HSA_EXT_IMAGE_GEOMETRY_3D, width, height, depth, desc, imageInfo); - dstPitch = imageInfo.size/(height == 0 ? 1 : height)/(depth == 0 ? 1 : depth); - } else { - depth = p->Depth; - height = p->Height; - widthInBytes = p->WidthInBytes; - width = p->dstArray->width; - hsa_ext_image_channel_order_t channelOrder; - switch(p->dstArray->NumChannels) { - case 2: - channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RG; - break; - case 3: - channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGB; - break; - case 4: - channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA; - break; - case 1: - default: - channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_R; - break; - } - hsa_ext_image_channel_type_t channelType; - e = ihipArrayToImageFormat(p->dstArray->Format,channelType); - srcPitch = p->srcPitch; - srcPtr = (void*)p->srcHost; - ySize = p->srcHeight; - dstPtr = p->dstArray->data; - hsa_ext_image_data_info_t imageInfo; - if(hipTextureType2DLayered == p->dstArray->textureType) - GetImageInfo(HSA_EXT_IMAGE_GEOMETRY_2DA, width, height, 0, channelOrder, channelType, imageInfo, depth); - else - GetImageInfo(HSA_EXT_IMAGE_GEOMETRY_3D, width, height, depth, channelOrder, channelType, imageInfo); - dstPitch = imageInfo.size/(height == 0 ? 1 : height)/(depth == 0 ? 1 : depth); + if ((p->dstArray->isDrv == true) ||( p->dstPtr.ptr!= nullptr)){ + return hipErrorInvalidValue; + } + // Array destination + dstByteSize = getByteSizeFromFormat(p->dstArray->desc); + hipChannelFormatDesc desc; + desc = p->dstArray->desc; + dstPtr = p->dstArray->data; + dstWidth = p->dstArray->width; + dstHeight = p->dstArray->height; + dstDepth = p->dstArray->depth; + dstPitch = dstByteSize * alignUp(dstWidth, IMAGE_PITCH_ALIGNMENT); + if(!copyWidthUpdate) { + copyWidth = copyWidth * dstByteSize; + copyWidthUpdate = true; } } else { - // Non array destination - depth = p->extent.depth; - height = p->extent.height; - widthInBytes = p->extent.width; - srcPitch = p->srcPtr.pitch; - srcPtr = p->srcPtr.ptr; + //Non Array destination dstPtr = p->dstPtr.ptr; - ySize = p->srcPtr.ysize; + dstWidth = p->dstPtr.xsize; + dstHeight = p->dstPtr.ysize; dstPitch = p->dstPtr.pitch; } + if (p->srcArray != nullptr) { + if ((p->srcArray->isDrv == true) ||( p->srcPtr.ptr!= nullptr)){ + return hipErrorInvalidValue; + } + // Array source + srcByteSize = getByteSizeFromFormat(p->srcArray->desc); + hipChannelFormatDesc desc; + desc = p->srcArray->desc; + srcPtr = p->srcArray->data; + srcWidth = p->srcArray->width; + srcHeight = p->srcArray->height; + srcDepth = p->srcArray->depth; + srcPitch = srcByteSize * alignUp(srcWidth, IMAGE_PITCH_ALIGNMENT); + if(!copyWidthUpdate) { + copyWidth = copyWidth * srcByteSize; + copyWidthUpdate = true; + } + } else { + //Non Array source + srcPtr = p->srcPtr.ptr; + srcWidth = p->srcPtr.xsize; + srcHeight = p->srcPtr.ysize; + srcPitch = p->srcPtr.pitch; + } + stream = ihipSyncAndResolveStream(stream); try { - if((widthInBytes == dstPitch) && (widthInBytes == srcPitch)) { + if((copyWidth == dstPitch) && (copyWidth == srcPitch)&& (copyHeight == dstHeight) &&(copyHeight == srcHeight)) { if(isAsync) - stream->locked_copyAsync((void*)dstPtr, (void*)srcPtr, widthInBytes*height*depth, p->kind); + stream->locked_copyAsync((void*)dstPtr, (void*)srcPtr, copyWidth*copyHeight*copyDepth, p->kind); else - stream->locked_copySync((void*)dstPtr, (void*)srcPtr, widthInBytes*height*depth, p->kind, false); + stream->locked_copySync((void*)dstPtr, (void*)srcPtr, copyWidth*copyHeight*copyDepth, p->kind, false); } else { - for (int i = 0; i < depth; i++) { - for (int j = 0; j < height; j++) { - // TODO: p->srcPos or p->dstPos are not 0. + for (int i = 0; i < copyDepth; i++) { + for (int j = 0; j < copyHeight; j++) { unsigned char* src = - (unsigned char*)srcPtr + i * ySize * srcPitch + j * srcPitch; + (unsigned char*)srcPtr + (i + srcZoffset) * srcHeight * srcPitch + (j + srcYoffset) * srcPitch + srcXoffset; unsigned char* dst = - (unsigned char*)dstPtr + i * height * dstPitch + j * dstPitch; + (unsigned char*)dstPtr + (i + dstZoffset) * dstHeight * dstPitch + (j + dstYoffset) * dstPitch + dstXoffset; if(isAsync) - stream->locked_copyAsync(dst, src, widthInBytes, p->kind); + stream->locked_copyAsync(dst, src, copyWidth, p->kind); else - stream->locked_copySync(dst, src, widthInBytes, p->kind); + stream->locked_copySync(dst, src, copyWidth, p->kind); } } } diff --git a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemcpy3D.cpp b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemcpy3D.cpp new file mode 100644 index 0000000000..255a3490b6 --- /dev/null +++ b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemcpy3D.cpp @@ -0,0 +1,110 @@ +/* +Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp + * TEST: %t + * HIT_END + */ + +#include "test_common.h" + +template +void runTest(int width,int height,int depth, hipChannelFormatKind formatKind) +{ + unsigned int size = width * height * depth * sizeof(T); + T* hData = (T*) malloc(size); + memset(hData, 0, size); + + for (int i = 0; i < depth; i++) { + for (int j = 0; j < height; j++) { + for (int k = 0; k < width; k++) { + hData[i*width*height + j*width +k] = i*width*height + j*width + k; + } + } + } + printf("test- sizeof(T) =%d\n", sizeof(T)); + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(T)*8, 0, 0, 0, formatKind); + hipArray *arr,*arr1; + + HIPCHECK(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, height, depth), hipArrayDefault)); + HIPCHECK(hipMalloc3DArray(&arr1, &channelDesc, make_hipExtent(width, height, depth), hipArrayDefault)); + hipMemcpy3DParms myparms = {0}; + myparms.srcPos = make_hipPos(0,0,0); + myparms.dstPos = make_hipPos(0,0,0); + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height); + myparms.dstArray = arr; + myparms.extent = make_hipExtent(width , height, depth); +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyHostToDevice; +#else + myparms.kind = hipMemcpyHostToDevice; +#endif + HIPCHECK(hipMemcpy3D(&myparms)); + HIPCHECK(hipDeviceSynchronize()); + //Array to Array + memset(&myparms,0x0, sizeof(hipMemcpy3DParms)); + myparms.srcPos = make_hipPos(0,0,0); + myparms.dstPos = make_hipPos(0,0,0); + myparms.srcArray = arr; + myparms.dstArray = arr1; + myparms.extent = make_hipExtent(width, height, depth); +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyDeviceToDevice; +#else + myparms.kind = hipMemcpyDeviceToDevice; +#endif + HIPCHECK(hipMemcpy3D(&myparms)); + HIPCHECK(hipDeviceSynchronize()); + + T *hOutputData = (T*) malloc(size); + memset(hOutputData, 0, size); + //Device to host + memset(&myparms,0x0, sizeof(hipMemcpy3DParms)); + myparms.srcPos = make_hipPos(0,0,0); + myparms.dstPos = make_hipPos(0,0,0); + myparms.dstPtr = make_hipPitchedPtr(hOutputData, width * sizeof(T), width, height); + myparms.srcArray = arr1; + myparms.extent = make_hipExtent(width, height, depth); +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyDeviceToHost; +#else + myparms.kind = hipMemcpyDeviceToHost; +#endif + HIPCHECK(hipMemcpy3D(&myparms)); + HIPCHECK(hipDeviceSynchronize()); + + // Check result + HipTest::checkArray(hData,hOutputData,width,height,depth); + hipFreeArray(arr); + hipFreeArray(arr1); + free(hData); + free(hOutputData); +} + +int main(int argc, char **argv) +{ + for(int i=1;i<25;i++) + { + runTest(i,i,i, hipChannelFormatKindFloat); + runTest(i+1,i,i, hipChannelFormatKindSigned); + runTest(i,i+1,i, hipChannelFormatKindSigned); + } + passed(); +} diff --git a/projects/clr/hipamd/tests/src/texture/simpleTexture2DLayered.cpp b/projects/clr/hipamd/tests/src/texture/simpleTexture2DLayered.cpp index fa545cb4ca..e5014dae6b 100644 --- a/projects/clr/hipamd/tests/src/texture/simpleTexture2DLayered.cpp +++ b/projects/clr/hipamd/tests/src/texture/simpleTexture2DLayered.cpp @@ -65,7 +65,7 @@ void runTest(int width,int height,int num_layers,texture texc; template -__global__ void simpleKernel3DArray(T* outputData, +__global__ void simpleKernel3DArray(T* outputData, int width, int height,int depth) { @@ -55,21 +52,18 @@ __global__ void simpleKernel3DArray(T* outputData, for (int j = 0; j < height; j++) { for (int k = 0; k < width; k++) { if(std::is_same::value) - outputData[i*width*height + j*width + k] = tex3D(texf, texf.textureObject, k, j, i); + outputData[i*width*height + j*width + k] = tex3D(texf, k, j, i); else if(std::is_same::value) - outputData[i*width*height + j*width + k] = tex3D(texi, texi.textureObject, k, j, i); + outputData[i*width*height + j*width + k] = tex3D(texi, k, j, i); else if(std::is_same::value) - outputData[i*width*height + j*width + k] = tex3D(texc, texc.textureObject, k, j, i); + outputData[i*width*height + j*width + k] = tex3D(texc, k, j, i); } } } } -//////////////////////////////////////////////////////////////////////////////// -//! Run a simple test for tex3D -//////////////////////////////////////////////////////////////////////////////// template -void runTest(int width,int height,int depth,texture *tex) +void runTest(int width,int height,int depth,texture *tex, hipChannelFormatKind formatKind) { unsigned int size = width * height * depth * sizeof(T); T* hData = (T*) malloc(size); @@ -84,17 +78,21 @@ void runTest(int width,int height,int depth,texture(i,i,i,&texf); - runTest(i+1,i,i,&texi); - runTest(i,i+1,i,&texc); + runTest(i,i,i,&texf, hipChannelFormatKindFloat); + runTest(i+1,i,i,&texi, hipChannelFormatKindSigned); + runTest(i,i+1,i,&texc, hipChannelFormatKindSigned); } passed(); } -