Fix hipMemcpy3D (#1798)
Fixes #1790 and #1791. hipMemcpy3D still requires further refactoring for different input and output combinations.
[ROCm/clr commit: ec84c16d75]
このコミットが含まれているのは:
Rahul Garg
GitHub
@@ -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) {
|
||||
|
||||
@@ -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);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@@ -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 <typename T>
|
||||
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<float>(i,i,i, hipChannelFormatKindFloat);
|
||||
runTest<int>(i+1,i,i, hipChannelFormatKindSigned);
|
||||
runTest<char>(i,i+1,i, hipChannelFormatKindSigned);
|
||||
}
|
||||
passed();
|
||||
}
|
||||
@@ -65,7 +65,7 @@ void runTest(int width,int height,int num_layers,texture<T, hipTextureType2DLaye
|
||||
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, num_layers);
|
||||
myparms.extent = make_hipExtent(width , height, num_layers);
|
||||
//myparms.kind = hipMemcpyHostToDevice;
|
||||
HIPCHECK(hipMemcpy3D(&myparms));
|
||||
|
||||
|
||||
@@ -21,15 +21,12 @@ THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
|
||||
* BUILD: %t %s ../test_common.cpp NVCC_OPTIONS -std=c++11
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
#include "test_common.h"
|
||||
|
||||
//typedef char T;
|
||||
const char *sampleName = "simpleTexture3D";
|
||||
|
||||
// Texture reference for 3D texture
|
||||
#if __HIP__
|
||||
__hip_pinned_shadow__
|
||||
@@ -47,7 +44,7 @@ __hip_pinned_shadow__
|
||||
texture<char, hipTextureType3D, hipReadModeElementType> texc;
|
||||
|
||||
template <typename T>
|
||||
__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<T, float>::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<T, int>::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<T, char>::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 <typename T>
|
||||
void runTest(int width,int height,int depth,texture<T, hipTextureType3D, hipReadModeElementType> *tex)
|
||||
void runTest(int width,int height,int depth,texture<T, hipTextureType3D, hipReadModeElementType> *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<T, hipTextureType3D, hipRead
|
||||
}
|
||||
|
||||
// Allocate array and copy image data
|
||||
hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(T)*8, 0, 0, 0, hipChannelFormatKindSigned);
|
||||
hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(T)*8, 0, 0, 0, formatKind);
|
||||
hipArray *arr;
|
||||
|
||||
HIPCHECK(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, height, depth), hipArrayCubemap));
|
||||
HIPCHECK(hipMalloc3DArray(&arr, &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));
|
||||
|
||||
// set texture parameters
|
||||
@@ -119,7 +117,7 @@ void runTest(int width,int height,int depth,texture<T, hipTextureType3D, hipRead
|
||||
|
||||
// copy result from device to host
|
||||
HIPCHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost));
|
||||
HipTest::checkArray(hData,hOutputData,width,height,depth);
|
||||
HipTest::checkArray(hData,hOutputData,width,height,depth);
|
||||
|
||||
hipFree(dData);
|
||||
hipFreeArray(arr);
|
||||
@@ -127,18 +125,13 @@ void runTest(int width,int height,int depth,texture<T, hipTextureType3D, hipRead
|
||||
free(hOutputData);
|
||||
}
|
||||
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
// Program main
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
int main(int argc, char **argv)
|
||||
{
|
||||
printf("%s starting...\n", sampleName);
|
||||
for(int i=1;i<25;i++)
|
||||
{
|
||||
runTest<float>(i,i,i,&texf);
|
||||
runTest<int>(i+1,i,i,&texi);
|
||||
runTest<char>(i,i+1,i,&texc);
|
||||
runTest<float>(i,i,i,&texf, hipChannelFormatKindFloat);
|
||||
runTest<int>(i+1,i,i,&texi, hipChannelFormatKindSigned);
|
||||
runTest<char>(i,i+1,i,&texc, hipChannelFormatKindSigned);
|
||||
}
|
||||
passed();
|
||||
}
|
||||
|
||||
|
||||
新しいイシューから参照
ユーザーをブロックする