diff --git a/include/hip/hcc_detail/host_defines.h b/include/hip/hcc_detail/host_defines.h index b21946e99f..ad28cc7626 100644 --- a/include/hip/hcc_detail/host_defines.h +++ b/include/hip/hcc_detail/host_defines.h @@ -72,7 +72,6 @@ THE SOFTWARE. #define __noinline__ __attribute__((noinline)) #define __forceinline__ inline __attribute__((always_inline)) -#define __hip_pinned_shadow__ __attribute__((hip_pinned_shadow)) #else diff --git a/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp b/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp index 6fd49fdb0f..5831da0e9d 100644 --- a/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp +++ b/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp @@ -21,11 +21,7 @@ THE SOFTWARE. */ #include "hip/hip_runtime.h" -#if __HIP__ -__hip_pinned_shadow__ -#else -extern -#endif + texture tex; extern "C" __global__ void tex2dKernel(float* outputData, int width, int height) { diff --git a/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp b/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp index e7c254e9fd..3903acd125 100644 --- a/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp +++ b/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp @@ -33,9 +33,6 @@ THE SOFTWARE. #define fileName "tex2d_kernel.code" -#if __HIP__ -__hip_pinned_shadow__ -#endif texture tex; bool testResult = false; diff --git a/tests/src/runtimeApi/module/tex2d_kernel.cpp b/tests/src/runtimeApi/module/tex2d_kernel.cpp index e744d88776..e52843441b 100644 --- a/tests/src/runtimeApi/module/tex2d_kernel.cpp +++ b/tests/src/runtimeApi/module/tex2d_kernel.cpp @@ -27,9 +27,6 @@ THE SOFTWARE. #include "hip/hip_runtime.h" -#if __HIP__ -__hip_pinned_shadow__ -#endif extern texture tex; extern "C" __global__ void tex2dKernel(float* outputData, int width, int height) { diff --git a/tests/src/texture/hipBindTex2DPitch.cpp b/tests/src/texture/hipBindTex2DPitch.cpp index 8c57520c00..6cee22a45d 100644 --- a/tests/src/texture/hipBindTex2DPitch.cpp +++ b/tests/src/texture/hipBindTex2DPitch.cpp @@ -28,9 +28,6 @@ THE SOFTWARE. #define SIZE_W 12 #define TYPE_t float -#if __HIP__ -__hip_pinned_shadow__ -#endif texture tex; // texture object is a kernel argument diff --git a/tests/src/texture/hipBindTexRef1DFetch.cpp b/tests/src/texture/hipBindTexRef1DFetch.cpp index 2e962fb05d..af79153fe0 100644 --- a/tests/src/texture/hipBindTexRef1DFetch.cpp +++ b/tests/src/texture/hipBindTexRef1DFetch.cpp @@ -32,9 +32,6 @@ THE SOFTWARE. #define N 512 -#if __HIP__ -__hip_pinned_shadow__ -#endif texture tex; __global__ void kernel(float *out) { diff --git a/tests/src/texture/hipNormalizedFloatValueTex.cpp b/tests/src/texture/hipNormalizedFloatValueTex.cpp index b4aa3e9c05..af33a29d3c 100644 --- a/tests/src/texture/hipNormalizedFloatValueTex.cpp +++ b/tests/src/texture/hipNormalizedFloatValueTex.cpp @@ -42,24 +42,12 @@ static float getNormalizedValue(const float value, return value; } -#if __HIP__ -__hip_pinned_shadow__ -#endif texture texc; -#if __HIP__ -__hip_pinned_shadow__ -#endif texture texuc; -#if __HIP__ -__hip_pinned_shadow__ -#endif texture texs; -#if __HIP__ -__hip_pinned_shadow__ -#endif texture texus; diff --git a/tests/src/texture/hipTextureRef2D.cpp b/tests/src/texture/hipTextureRef2D.cpp index 5573cf6884..5247f81fe0 100644 --- a/tests/src/texture/hipTextureRef2D.cpp +++ b/tests/src/texture/hipTextureRef2D.cpp @@ -9,9 +9,7 @@ #include #include "test_common.h" -#if __HIP__ -__hip_pinned_shadow__ -#endif + texture tex; __global__ void tex2DKernel(float* outputData, diff --git a/tests/src/texture/simpleTexture2DLayered.cpp b/tests/src/texture/simpleTexture2DLayered.cpp index f4d3aac1e5..8b1bbb64a3 100644 --- a/tests/src/texture/simpleTexture2DLayered.cpp +++ b/tests/src/texture/simpleTexture2DLayered.cpp @@ -30,9 +30,6 @@ THE SOFTWARE. typedef float T; // Texture reference for 2D Layered texture -#if __HIP__ -__hip_pinned_shadow__ -#endif texture tex2DL; __global__ void simpleKernelLayeredArray(T* outputData,int width,int height,int layer) diff --git a/tests/src/texture/simpleTexture3D.cpp b/tests/src/texture/simpleTexture3D.cpp index a494a1a6c0..82f6cf5e99 100644 --- a/tests/src/texture/simpleTexture3D.cpp +++ b/tests/src/texture/simpleTexture3D.cpp @@ -31,19 +31,10 @@ THE SOFTWARE. const char *sampleName = "simpleTexture3D"; // Texture reference for 3D texture -#if __HIP__ -__hip_pinned_shadow__ -#endif texture texf; -#if __HIP__ -__hip_pinned_shadow__ -#endif texture texi; -#if __HIP__ -__hip_pinned_shadow__ -#endif texture texc; template diff --git a/vdi/hip_hcc.def.in b/vdi/hip_hcc.def.in index 5eaedf6851..d8101a1cb8 100755 --- a/vdi/hip_hcc.def.in +++ b/vdi/hip_hcc.def.in @@ -161,6 +161,8 @@ __hipPushCallConfiguration __hipRegisterFatBinary __hipRegisterFunction __hipRegisterVar +__hipRegisterSurface +__hipRegisterTexture __hipUnregisterFatBinary __gnu_h2f_ieee __gnu_f2h_ieee diff --git a/vdi/hip_hcc.map.in b/vdi/hip_hcc.map.in index 98a3479f40..2cd55b5581 100755 --- a/vdi/hip_hcc.map.in +++ b/vdi/hip_hcc.map.in @@ -161,6 +161,8 @@ global: __hipRegisterFatBinary; __hipRegisterFunction; __hipRegisterVar; + __hipRegisterSurface; + __hipRegisterTexture; __hipUnregisterFatBinary; __gnu_h2f_ieee; __gnu_f2h_ieee; diff --git a/vdi/hip_internal.hpp b/vdi/hip_internal.hpp index 0d0caada14..3e09df03d4 100755 --- a/vdi/hip_internal.hpp +++ b/vdi/hip_internal.hpp @@ -222,13 +222,22 @@ public: std::vector< std::pair< hipModule_t, bool > >* modules; std::vector functions; }; + enum DeviceVarKind { + DVK_Variable, + DVK_Surface, + DVK_Texture + }; struct DeviceVar { + DeviceVarKind kind; void* shadowVptr; std::string hostVar; size_t size; std::vector< std::pair< hipModule_t, bool > >* modules; std::vector rvars; bool dyn_undef; + int type; // surface/texture type + int norm; // texture has normalized output + bool shadowAllocated = false; // shadow ptr is allocated on-demand and needs freeing. }; private: class Module { @@ -278,6 +287,9 @@ public: hipDeviceptr_t* dev_ptr, size_t* size_ptr); bool getTexRef(const char* hostVar, hipModule_t hmod, textureReference** texRef); + bool getGlobalVarFromSymbol(const void* hostVar, int deviceId, + hipDeviceptr_t* dev_ptr, size_t* size_ptr); + bool getShadowVarInfo(std::string var_name, hipModule_t hmod, void** var_addr, size_t* var_size); void setupArgument(const void *arg, size_t size, size_t offset); diff --git a/vdi/hip_module.cpp b/vdi/hip_module.cpp index f523a86a0f..5d09f88293 100755 --- a/vdi/hip_module.cpp +++ b/vdi/hip_module.cpp @@ -150,8 +150,15 @@ inline bool ihipModuleRegisterUndefined(amd::Program* program, hipModule_t* modu = new texture(); memset(tex_hptr, 0x00, sizeof(texture)); - PlatformState::DeviceVar dvar{ reinterpret_cast(tex_hptr), it->c_str(), sizeof(*tex_hptr), modules, - std::vector{ g_devices.size()}, true }; + PlatformState::DeviceVar dvar{PlatformState::DVK_Variable, + reinterpret_cast(tex_hptr), + it->c_str(), + sizeof(*tex_hptr), + modules, + std::vector{g_devices.size()}, + true, + /*type*/ 0, + /*norm*/ 0}; PlatformState::instance().registerVar(it->c_str(), dvar); } @@ -194,8 +201,15 @@ inline bool ihipModuleRegisterGlobal(amd::Program* program, hipModule_t* module) modules->at(dev) = std::make_pair(*module, true); } - PlatformState::DeviceVar dvar{nullptr, it->c_str(), 0, modules, - std::vector{ g_devices.size()}, false }; + PlatformState::DeviceVar dvar{PlatformState::DVK_Variable, + nullptr, + it->c_str(), + 0, + modules, + std::vector{g_devices.size()}, + false, + /*type*/ 0, + /*norm*/ 0}; PlatformState::instance().registerVar(it->c_str(), dvar); } @@ -673,4 +687,3 @@ hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const HIP_RETURN(hipSuccess); } - diff --git a/vdi/hip_platform.cpp b/vdi/hip_platform.cpp index 7b63d5225b..d00974c6bc 100755 --- a/vdi/hip_platform.cpp +++ b/vdi/hip_platform.cpp @@ -19,7 +19,7 @@ THE SOFTWARE. */ #include - +#include #include "hip_internal.hpp" #include "platform/program.hpp" #include "platform/runtime.hpp" @@ -220,7 +220,7 @@ std::vector< std::pair >* PlatformState::unregisterVar(hipMod DeviceVar& dvar = it->second; if ((*dvar.modules)[0].first == hmod) { rmodules = dvar.modules; - if (dvar.dyn_undef) { + if (dvar.shadowAllocated) { texture* tex_hptr = reinterpret_cast *>(dvar.shadowVptr); delete tex_hptr; @@ -474,12 +474,27 @@ bool PlatformState::getTexRef(const char* hostVar, hipModule_t hmod, textureRefe return false; } - if (!dvar->dyn_undef) { - DevLogPrintfError("HostVar: %s is not created through hipModuleLoad \n", hostVar); + switch (dvar->kind) { + case PlatformState::DVK_Variable: + // TODO: Need to define a target-specific symbol info to indicate the device + // variable kind, i.e. regular variable, texture or surface. + // Before that, have to assume the specified variable is a texture or + // surface reference variable. + dvar->kind = DVK_Texture; + // FALL THROUGH + case PlatformState::DVK_Texture: + break; + default: + // If it's already used as non-texture variable, bail out. return false; } - *texRef = new (dvar->shadowVptr) texture{}; + if (!dvar->shadowVptr) { + dvar->shadowVptr = new texture{}; + dvar->shadowAllocated = true; + } + *texRef = reinterpret_cast(dvar->shadowVptr); + registerVarSym(dvar->shadowVptr, hostVar); return true; } @@ -523,6 +538,18 @@ bool PlatformState::getGlobalVar(const char* hostVar, int deviceId, hipModule_t } } +bool PlatformState::getGlobalVarFromSymbol(const void* hostVar, int deviceId, + hipDeviceptr_t* dev_ptr, + size_t* size_ptr) { + std::string symbolName; + if (!PlatformState::instance().findSymbol(hostVar, symbolName)) { + return false; + } + return PlatformState::instance().getGlobalVar(symbolName.c_str(), + ihipGetDevice(), nullptr, + dev_ptr, size_ptr); +} + void PlatformState::setupArgument(const void *arg, size_t size, size_t offset) { auto& arguments = execStack_.top().arguments_; @@ -577,11 +604,56 @@ extern "C" void __hipRegisterVar( int constant, // Whether this variable is constant int global) // Unknown, always 0 { - PlatformState::DeviceVar dvar{var, std::string{ hostVar }, size, modules, - std::vector{g_devices.size()}, false }; + PlatformState::DeviceVar dvar{PlatformState::DVK_Variable, + var, + std::string{hostVar}, + size, + modules, + std::vector{g_devices.size()}, + false, + /*type*/ 0, + /*norm*/ 0}; - PlatformState::instance().registerVar(hostVar, dvar); - PlatformState::instance().registerVarSym(var, deviceVar); + PlatformState::instance().registerVar(hostVar, dvar); + PlatformState::instance().registerVarSym(var, deviceVar); +} + +extern "C" void __hipRegisterSurface(std::vector>* + modules, // The device modules containing code object + void* var, // The shadow variable in host code + char* hostVar, // Variable name in host code + char* deviceVar, // Variable name in device code + int type, int ext) { + PlatformState::DeviceVar dvar{PlatformState::DVK_Surface, + var, + std::string{hostVar}, + sizeof(surfaceReference), // Copy whole surfaceReference + modules, + std::vector{g_devices.size()}, + false, + type, + /*norm*/ 0}; + PlatformState::instance().registerVar(hostVar, dvar); + PlatformState::instance().registerVarSym(var, deviceVar); +} + +extern "C" void __hipRegisterTexture(std::vector>* + modules, // The device modules containing code object + void* var, // The shadow variable in host code + char* hostVar, // Variable name in host code + char* deviceVar, // Variable name in device code + int type, int norm, int ext) { + PlatformState::DeviceVar dvar{PlatformState::DVK_Texture, + var, + std::string{hostVar}, + sizeof(textureReference), // Copy whole textureReference so far. + modules, + std::vector{g_devices.size()}, + false, + type, + norm}; + PlatformState::instance().registerVar(hostVar, dvar); + PlatformState::instance().registerVarSym(var, deviceVar); } extern "C" void __hipUnregisterFatBinary(std::vector< std::pair >* modules) diff --git a/vdi/hip_texture.cpp b/vdi/hip_texture.cpp index da24d663d1..94026c8e33 100755 --- a/vdi/hip_texture.cpp +++ b/vdi/hip_texture.cpp @@ -24,6 +24,9 @@ #include "hip_conversions.hpp" #include "platform/sampler.hpp" +hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, + amd::HostQueue& queue, bool isAsync = false); + struct __hip_texture { uint32_t imageSRD[HIP_IMAGE_OBJECT_SIZE_DWORD]; uint32_t samplerSRD[HIP_SAMPLER_OBJECT_SIZE_DWORD]; @@ -473,7 +476,20 @@ hipError_t hipBindTexture2D(size_t* offset, size_t pitch) { HIP_INIT_API(hipBindTexture2D, offset, texref, devPtr, desc, width, height, pitch); - HIP_RETURN(ihipBindTexture2D(offset, texref, devPtr, desc, width, height, pitch)); + hipDeviceptr_t refDevPtr = nullptr; + size_t refDevSize = 0; + if (!PlatformState::instance().getGlobalVarFromSymbol(texref, ihipGetDevice(), &refDevPtr, + &refDevSize)) { + HIP_RETURN(hipErrorInvalidSymbol); + } + assert(refDevSize == sizeof(textureReference)); + hipError_t err = ihipBindTexture2D(offset, texref, devPtr, desc, width, height, pitch); + if (err != hipSuccess) { + HIP_RETURN(err); + } + // Copy to device. + amd::HostQueue* queue = hip::getNullStream(); + HIP_RETURN(ihipMemcpy(refDevPtr, texref, refDevSize, hipMemcpyHostToDevice, *queue)); } hipError_t ihipBindTextureToArray(const textureReference* texref, @@ -507,7 +523,20 @@ hipError_t hipBindTextureToArray(const textureReference* texref, const hipChannelFormatDesc* desc) { HIP_INIT_API(hipBindTextureToArray, texref, array, desc); - HIP_RETURN(ihipBindTextureToArray(texref, array, desc)); + hipDeviceptr_t refDevPtr = nullptr; + size_t refDevSize = 0; + if (!PlatformState::instance().getGlobalVarFromSymbol(texref, ihipGetDevice(), &refDevPtr, + &refDevSize)) { + HIP_RETURN(hipErrorInvalidSymbol); + } + assert(refDevSize == sizeof(textureReference)); + hipError_t err = ihipBindTextureToArray(texref, array, desc); + if (err != hipSuccess) { + HIP_RETURN(err); + } + // Copy to device. + amd::HostQueue* queue = hip::getNullStream(); + HIP_RETURN(ihipMemcpy(refDevPtr, texref, refDevSize, hipMemcpyHostToDevice, *queue)); } hipError_t ihipBindTextureToMipmappedArray(const textureReference* texref, @@ -541,7 +570,20 @@ hipError_t hipBindTextureToMipmappedArray(const textureReference* texref, const hipChannelFormatDesc* desc) { HIP_INIT_API(hipBindTextureToMipmappedArray, texref, mipmappedArray, desc); - HIP_RETURN(ihipBindTextureToMipmappedArray(texref, mipmappedArray, desc)); + hipDeviceptr_t refDevPtr = nullptr; + size_t refDevSize = 0; + if (!PlatformState::instance().getGlobalVarFromSymbol(texref, ihipGetDevice(), &refDevPtr, + &refDevSize)) { + HIP_RETURN(hipErrorInvalidSymbol); + } + assert(refDevSize == sizeof(textureReference)); + hipError_t err = ihipBindTextureToMipmappedArray(texref, mipmappedArray, desc); + if (err != hipSuccess) { + HIP_RETURN(err); + } + // Copy to device. + amd::HostQueue* queue = hip::getNullStream(); + HIP_RETURN(ihipMemcpy(refDevPtr, texref, refDevSize, hipMemcpyHostToDevice, *queue)); } hipError_t hipUnbindTexture(const textureReference* texref) { @@ -564,7 +606,20 @@ hipError_t hipBindTexture(size_t* offset, size_t size) { HIP_INIT_API(hipBindTexture, offset, texref, devPtr, desc, size); - HIP_RETURN(ihipBindTexture(offset, texref, devPtr, desc, size)); + hipDeviceptr_t refDevPtr = nullptr; + size_t refDevSize = 0; + if (!PlatformState::instance().getGlobalVarFromSymbol(texref, ihipGetDevice(), &refDevPtr, + &refDevSize)) { + HIP_RETURN(hipErrorInvalidSymbol); + } + assert(refDevSize == sizeof(textureReference)); + hipError_t err = ihipBindTexture(offset, texref, devPtr, desc, size); + if (err != hipSuccess) { + HIP_RETURN(err); + } + // Copy to device. + amd::HostQueue* queue = hip::getNullStream(); + HIP_RETURN(ihipMemcpy(refDevPtr, texref, refDevSize, hipMemcpyHostToDevice, *queue)); } hipError_t hipGetChannelDesc(hipChannelFormatDesc* desc, @@ -599,9 +654,12 @@ hipError_t hipGetTextureAlignmentOffset(size_t* offset, hipError_t hipGetTextureReference(const textureReference** texref, const void* symbol) { HIP_INIT_API(hipGetTextureReference, texref, symbol); - assert(0 && "Unimplemented"); + if (texref == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + *texref = reinterpret_cast(symbol); - HIP_RETURN(hipErrorNotSupported); + HIP_RETURN(hipSuccess); } hipError_t hipTexRefSetFormat(textureReference* texRef, @@ -744,6 +802,14 @@ hipError_t hipTexRefSetArray(textureReference* texRef, HIP_RETURN(hipErrorInvalidValue); } + hipDeviceptr_t refDevPtr = nullptr; + size_t refDevSize = 0; + if (!PlatformState::instance().getGlobalVarFromSymbol(texRef, ihipGetDevice(), &refDevPtr, + &refDevSize)) { + HIP_RETURN(hipErrorInvalidSymbol); + } + 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. @@ -758,7 +824,13 @@ hipError_t hipTexRefSetArray(textureReference* texRef, hipResourceViewFormat format = hip::getResourceViewFormat(hip::getChannelFormatDesc(texRef->numChannels, texRef->format)); hipResourceViewDesc resViewDesc = hip::getResourceViewDesc(array, format); - HIP_RETURN(ihipCreateTextureObject(&texRef->textureObject, &resDesc, &texDesc, &resViewDesc)); + hipError_t err = ihipCreateTextureObject(&texRef->textureObject, &resDesc, &texDesc, &resViewDesc); + if (err != hipSuccess) { + HIP_RETURN(err); + } + // Copy to device. + amd::HostQueue* queue = hip::getNullStream(); + HIP_RETURN(ihipMemcpy(refDevPtr, texRef, refDevSize, hipMemcpyHostToDevice, *queue)); } hipError_t hipTexRefGetAddress(hipDeviceptr_t* dptr, @@ -808,6 +880,14 @@ hipError_t hipTexRefSetAddress(size_t* ByteOffset, HIP_RETURN(hipErrorInvalidValue); } + hipDeviceptr_t refDevPtr = nullptr; + size_t refDevSize = 0; + if (!PlatformState::instance().getGlobalVarFromSymbol(texRef, ihipGetDevice(), &refDevPtr, + &refDevSize)) { + HIP_RETURN(hipErrorInvalidSymbol); + } + 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. @@ -828,7 +908,13 @@ hipError_t hipTexRefSetAddress(size_t* ByteOffset, hipTextureDesc texDesc = hip::getTextureDesc(texRef); - HIP_RETURN(ihipCreateTextureObject(&texRef->textureObject, &resDesc, &texDesc, nullptr)); + hipError_t err = ihipCreateTextureObject(&texRef->textureObject, &resDesc, &texDesc, nullptr); + if (err != hipSuccess) { + HIP_RETURN(err); + } + // Copy to device. + amd::HostQueue* queue = hip::getNullStream(); + HIP_RETURN(ihipMemcpy(refDevPtr, texRef, refDevSize, hipMemcpyHostToDevice, *queue)); } hipError_t hipTexRefSetAddress2D(textureReference* texRef, @@ -841,6 +927,14 @@ hipError_t hipTexRefSetAddress2D(textureReference* texRef, HIP_RETURN(hipErrorInvalidValue); } + hipDeviceptr_t refDevPtr = nullptr; + size_t refDevSize = 0; + if (!PlatformState::instance().getGlobalVarFromSymbol(texRef, ihipGetDevice(), &refDevPtr, + &refDevSize)) { + HIP_RETURN(hipErrorInvalidSymbol); + } + 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. @@ -856,7 +950,13 @@ hipError_t hipTexRefSetAddress2D(textureReference* texRef, hipTextureDesc texDesc = hip::getTextureDesc(texRef); - HIP_RETURN(ihipCreateTextureObject(&texRef->textureObject, &resDesc, &texDesc, nullptr)); + hipError_t err = ihipCreateTextureObject(&texRef->textureObject, &resDesc, &texDesc, nullptr); + if (err != hipSuccess) { + HIP_RETURN(err); + } + // Copy to device. + amd::HostQueue* queue = hip::getNullStream(); + HIP_RETURN(ihipMemcpy(refDevPtr, texRef, refDevSize, hipMemcpyHostToDevice, *queue)); } hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f) { @@ -1107,6 +1207,14 @@ hipError_t hipTexRefSetMipmappedArray(textureReference* texRef, HIP_RETURN(hipErrorInvalidValue); } + hipDeviceptr_t refDevPtr = nullptr; + size_t refDevSize = 0; + if (!PlatformState::instance().getGlobalVarFromSymbol(texRef, ihipGetDevice(), &refDevPtr, + &refDevSize)) { + HIP_RETURN(hipErrorInvalidSymbol); + } + 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. @@ -1121,7 +1229,13 @@ hipError_t hipTexRefSetMipmappedArray(textureReference* texRef, hipResourceViewFormat format = hip::getResourceViewFormat(hip::getChannelFormatDesc(texRef->numChannels, texRef->format)); hipResourceViewDesc resViewDesc = hip::getResourceViewDesc(mipmappedArray, format); - HIP_RETURN(ihipCreateTextureObject(&texRef->textureObject, &resDesc, &texDesc, &resViewDesc)); + hipError_t err = ihipCreateTextureObject(&texRef->textureObject, &resDesc, &texDesc, &resViewDesc); + if (err != hipSuccess) { + HIP_RETURN(err); + } + // Copy to device. + amd::HostQueue* queue = hip::getNullStream(); + HIP_RETURN(ihipMemcpy(refDevPtr, texRef, refDevSize, hipMemcpyHostToDevice, *queue)); } hipError_t hipTexObjectCreate(hipTextureObject_t* pTexObject,