From 5e8b10d9a0ea518dcb6c3c88b0a9cfd143363bcd Mon Sep 17 00:00:00 2001 From: Amruth Date: Tue, 27 Mar 2018 14:12:05 -0700 Subject: support for pinned memories - temporary fix --- src/abstract_hardware_model.h | 3 +++ 1 file changed, 3 insertions(+) (limited to 'src/abstract_hardware_model.h') diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index aaa4b00..67b36c7 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -182,6 +182,9 @@ void increment_x_then_y_then_z( dim3 &i, const dim3 &bound); class stream_manager; struct CUstream_st; extern stream_manager * g_stream_manager; +//support for pinned memories added +extern std::map pinned_memory; +extern std::map pinned_memory_size; class kernel_info_t { public: -- cgit v1.3 From 4fb2d23f350bd8921417f5c09bde73594e5c8a0b Mon Sep 17 00:00:00 2001 From: Jonathan Date: Thu, 17 May 2018 09:27:07 -0700 Subject: allows gpgpusim to select a set of texture array,attr,info but maybe not the right one --- src/abstract_hardware_model.h | 73 +++++++++++++++++++++++++++++++++++++------ src/cuda-sim/cuda-sim.cc | 30 ++++++++++++++++++ 2 files changed, 94 insertions(+), 9 deletions(-) (limited to 'src/abstract_hardware_model.h') diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index f5708bc..ca41e68 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -522,22 +522,77 @@ public: } const struct cudaArray* get_texarray( const struct textureReference *texref ) const { - std::map::const_iterator t=m_TextureRefToCudaArray.find(texref); - assert(t != m_TextureRefToCudaArray.end()); - return t->second; + + for (std::map::const_iterator kv = m_TextureRefToCudaArray.begin(); kv!= m_TextureRefToCudaArray.end(); kv ++){ + const struct textureReference* tr = kv->first; + if (tr->normalized==texref->normalized&& + tr->filterMode==texref->filterMode&& + tr->addressMode[0]==texref->addressMode[0]&& + tr->addressMode[1]==texref->addressMode[1]&& + tr->addressMode[2]==texref->addressMode[2]&& + tr->channelDesc.x==texref->channelDesc.x&& + tr->channelDesc.y==texref->channelDesc.y&& + tr->channelDesc.z==texref->channelDesc.z&& + tr->channelDesc.w==texref->channelDesc.w&& + tr->channelDesc.f==texref->channelDesc.f){ + + return kv->second; + } + } + + assert(false); +// std::map::const_iterator t=m_TextureRefToCudaArray.find(texref); +// assert(t != m_TextureRefToCudaArray.end()); +// return t->second; } const struct textureInfo* get_texinfo( const struct textureReference *texref ) const { - std::map::const_iterator t=m_TextureRefToTexureInfo.find(texref); - assert(t != m_TextureRefToTexureInfo.end()); - return t->second; + for (std::map::const_iterator kv = m_TextureRefToTexureInfo.begin(); kv!= m_TextureRefToTexureInfo.end(); kv ++){ + const struct textureReference* tr = kv->first; + if (tr->normalized==texref->normalized&& + tr->filterMode==texref->filterMode&& + tr->addressMode[0]==texref->addressMode[0]&& + tr->addressMode[1]==texref->addressMode[1]&& + tr->addressMode[2]==texref->addressMode[2]&& + tr->channelDesc.x==texref->channelDesc.x&& + tr->channelDesc.y==texref->channelDesc.y&& + tr->channelDesc.z==texref->channelDesc.z&& + tr->channelDesc.w==texref->channelDesc.w&& + tr->channelDesc.f==texref->channelDesc.f){ + + return kv->second; + } + } + + assert(false); +// std::map::const_iterator t=m_TextureRefToTexureInfo.find(texref); +// assert(t != m_TextureRefToTexureInfo.end()); +// return t->second; } const struct textureReferenceAttr* get_texattr( const struct textureReference *texref ) const { - std::map::const_iterator t=m_TextureRefToAttribute.find(texref); - assert(t != m_TextureRefToAttribute.end()); - return t->second; + for (std::map::const_iterator kv = m_TextureRefToAttribute.begin(); kv!= m_TextureRefToAttribute.end(); kv ++){ + const struct textureReference* tr = kv->first; + if (tr->normalized==texref->normalized&& + tr->filterMode==texref->filterMode&& + tr->addressMode[0]==texref->addressMode[0]&& + tr->addressMode[1]==texref->addressMode[1]&& + tr->addressMode[2]==texref->addressMode[2]&& + tr->channelDesc.x==texref->channelDesc.x&& + tr->channelDesc.y==texref->channelDesc.y&& + tr->channelDesc.z==texref->channelDesc.z&& + tr->channelDesc.w==texref->channelDesc.w&& + tr->channelDesc.f==texref->channelDesc.f){ + + return kv->second; + } + } + + assert(false); +// std::map::const_iterator t=m_TextureRefToAttribute.find(texref); +// assert(t != m_TextureRefToAttribute.end()); +// return t->second; } const gpgpu_functional_sim_config &get_config() const { return m_function_model_config; } diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 946043a..6e04ca8 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -140,6 +140,36 @@ unsigned int intLOGB2( unsigned int v ) { void gpgpu_t::gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* texref, const struct cudaArray* array) { + // counts number of matches +// int normalized; +// enum cudaTextureFilterMode filterMode; +// enum cudaTextureAddressMode addressMode[3]; +// struct cudaChannelFormatDesc channelDesc; +// int x; +// int y; +// int z; +// int w; +// enum cudaChannelFormatKind f; +// int trMatches = 0; +// for (auto& kv : m_NameToTextureRef){ +// const struct textureReference* tr = kv.second; +// if (tr->normalized==texref->normalized&& +// tr->filterMode==texref->filterMode&& +// tr->addressMode[0]==texref->addressMode[0]&& +// tr->addressMode[1]==texref->addressMode[1]&& +// tr->addressMode[2]==texref->addressMode[2]&& +// tr->channelDesc.x==texref->channelDesc.x&& +// tr->channelDesc.y==texref->channelDesc.y&& +// tr->channelDesc.z==texref->channelDesc.z&& +// tr->channelDesc.w==texref->channelDesc.w&& +// tr->channelDesc.f==texref->channelDesc.f){ +// +// m_TextureRefToCudaArray[tr] = array; +// trMatches++; +// } +// } +// printf("GPGPU-Sim PTX: matches to texref = %d\n", trMatches); +// assert(trMatches==1); m_TextureRefToCudaArray[texref] = array; unsigned int texel_size_bits = array->desc.w + array->desc.x + array->desc.y + array->desc.z; unsigned int texel_size = texel_size_bits/8; -- cgit v1.3 From d36216db7a9d1d2360acf00afb5303b1f76fb7c7 Mon Sep 17 00:00:00 2001 From: Jonathan Date: Thu, 17 May 2018 16:24:41 -0700 Subject: counting matches with same texref contents, added cuda8 fields to texref struct --- src/abstract_hardware_model.h | 143 ++++++++++++++++++++++++++++++++++++++---- 1 file changed, 130 insertions(+), 13 deletions(-) (limited to 'src/abstract_hardware_model.h') diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index ca41e68..c21ff6d 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -25,6 +25,11 @@ // OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + +#ifndef VERSION_EIGHT +#define VERSION_EIGHT +#endif + #ifndef ABSTRACT_HARDWARE_MODEL_INCLUDED #define ABSTRACT_HARDWARE_MODEL_INCLUDED @@ -445,6 +450,35 @@ struct textureReference { enum cudaTextureFilterMode filterMode; enum cudaTextureAddressMode addressMode[3]; struct cudaChannelFormatDesc channelDesc; + +#ifdef VERSION_EIGHT + /** + * Perform sRGB->linear conversion during texture read + */ + int sRGB; + /** + * Limit to the anisotropy ratio + */ + unsigned int maxAnisotropy; + /** + * Mipmap filter mode + */ + enum cudaTextureFilterMode mipmapFilterMode; + /** + * Offset applied to the supplied mipmap level + */ + float mipmapLevelBias; + /** + * Lower end of the mipmap level range to clamp access to + */ + float minMipmapLevelClamp; + /** + * Upper end of the mipmap level range to clamp access to + */ + float maxMipmapLevelClamp; + int __cudaReserved[15]; +#endif + }; #endif @@ -522,7 +556,8 @@ public: } const struct cudaArray* get_texarray( const struct textureReference *texref ) const { - + int matches = 0; + const struct cudaArray* t = NULL; for (std::map::const_iterator kv = m_TextureRefToCudaArray.begin(); kv!= m_TextureRefToCudaArray.end(); kv ++){ const struct textureReference* tr = kv->first; if (tr->normalized==texref->normalized&& @@ -534,19 +569,49 @@ public: tr->channelDesc.y==texref->channelDesc.y&& tr->channelDesc.z==texref->channelDesc.z&& tr->channelDesc.w==texref->channelDesc.w&& - tr->channelDesc.f==texref->channelDesc.f){ - - return kv->second; + tr->channelDesc.f==texref->channelDesc.f && + tr->sRGB==texref->sRGB&& + tr->maxAnisotropy==texref->maxAnisotropy&& + tr->mipmapFilterMode==texref->mipmapFilterMode&& + tr->mipmapLevelBias==texref->mipmapLevelBias&& + tr->minMipmapLevelClamp==texref->minMipmapLevelClamp&& + tr->maxMipmapLevelClamp==texref->maxMipmapLevelClamp&& + tr->__cudaReserved[0] ==texref->__cudaReserved[0]&& + tr->__cudaReserved[1] ==texref->__cudaReserved[1]&& + tr->__cudaReserved[2] ==texref->__cudaReserved[2]&& + tr->__cudaReserved[3] ==texref->__cudaReserved[3]&& + tr->__cudaReserved[4] ==texref->__cudaReserved[4]&& + tr->__cudaReserved[5] ==texref->__cudaReserved[5]&& + tr->__cudaReserved[6] ==texref->__cudaReserved[6]&& + tr->__cudaReserved[7] ==texref->__cudaReserved[7]&& + tr->__cudaReserved[8] ==texref->__cudaReserved[8]&& + tr->__cudaReserved[9] ==texref->__cudaReserved[9]&& + tr->__cudaReserved[10]==texref->__cudaReserved[10]&& + tr->__cudaReserved[11]==texref->__cudaReserved[11]&& + tr->__cudaReserved[12]==texref->__cudaReserved[12]&& + tr->__cudaReserved[13]==texref->__cudaReserved[13]&& + tr->__cudaReserved[14]==texref->__cudaReserved[14] + ){ + matches++; + t = kv->second; + //return kv->second; } } - assert(false); + printf("matches (texarray) = %d\n", matches); + //assert(matches==1); + return t; + + //assert(false); + // std::map::const_iterator t=m_TextureRefToCudaArray.find(texref); // assert(t != m_TextureRefToCudaArray.end()); // return t->second; } const struct textureInfo* get_texinfo( const struct textureReference *texref ) const { + int matches = 0; + const struct textureInfo* t = NULL; for (std::map::const_iterator kv = m_TextureRefToTexureInfo.begin(); kv!= m_TextureRefToTexureInfo.end(); kv ++){ const struct textureReference* tr = kv->first; if (tr->normalized==texref->normalized&& @@ -558,13 +623,38 @@ public: tr->channelDesc.y==texref->channelDesc.y&& tr->channelDesc.z==texref->channelDesc.z&& tr->channelDesc.w==texref->channelDesc.w&& - tr->channelDesc.f==texref->channelDesc.f){ - - return kv->second; + tr->channelDesc.f==texref->channelDesc.f&& + tr->sRGB==texref->sRGB&& + tr->maxAnisotropy==texref->maxAnisotropy&& + tr->mipmapFilterMode==texref->mipmapFilterMode&& + tr->mipmapLevelBias==texref->mipmapLevelBias&& + tr->minMipmapLevelClamp==texref->minMipmapLevelClamp&& + tr->maxMipmapLevelClamp==texref->maxMipmapLevelClamp&& + tr->__cudaReserved[0] ==texref->__cudaReserved[0]&& + tr->__cudaReserved[1] ==texref->__cudaReserved[1]&& + tr->__cudaReserved[2] ==texref->__cudaReserved[2]&& + tr->__cudaReserved[3] ==texref->__cudaReserved[3]&& + tr->__cudaReserved[4] ==texref->__cudaReserved[4]&& + tr->__cudaReserved[5] ==texref->__cudaReserved[5]&& + tr->__cudaReserved[6] ==texref->__cudaReserved[6]&& + tr->__cudaReserved[7] ==texref->__cudaReserved[7]&& + tr->__cudaReserved[8] ==texref->__cudaReserved[8]&& + tr->__cudaReserved[9] ==texref->__cudaReserved[9]&& + tr->__cudaReserved[10]==texref->__cudaReserved[10]&& + tr->__cudaReserved[11]==texref->__cudaReserved[11]&& + tr->__cudaReserved[12]==texref->__cudaReserved[12]&& + tr->__cudaReserved[13]==texref->__cudaReserved[13]&& + tr->__cudaReserved[14]==texref->__cudaReserved[14]){ + matches++; + t = kv->second; + //return kv->second; } } - assert(false); + printf("matches (texinfo) = %d\n", matches); + //assert(matches==1); + return t; + // std::map::const_iterator t=m_TextureRefToTexureInfo.find(texref); // assert(t != m_TextureRefToTexureInfo.end()); // return t->second; @@ -572,6 +662,8 @@ public: const struct textureReferenceAttr* get_texattr( const struct textureReference *texref ) const { + int matches = 0; + const struct textureReferenceAttr* t = NULL; for (std::map::const_iterator kv = m_TextureRefToAttribute.begin(); kv!= m_TextureRefToAttribute.end(); kv ++){ const struct textureReference* tr = kv->first; if (tr->normalized==texref->normalized&& @@ -583,13 +675,38 @@ public: tr->channelDesc.y==texref->channelDesc.y&& tr->channelDesc.z==texref->channelDesc.z&& tr->channelDesc.w==texref->channelDesc.w&& - tr->channelDesc.f==texref->channelDesc.f){ - - return kv->second; + tr->channelDesc.f==texref->channelDesc.f&& + tr->sRGB==texref->sRGB&& + tr->maxAnisotropy==texref->maxAnisotropy&& + tr->mipmapFilterMode==texref->mipmapFilterMode&& + tr->mipmapLevelBias==texref->mipmapLevelBias&& + tr->minMipmapLevelClamp==texref->minMipmapLevelClamp&& + tr->maxMipmapLevelClamp==texref->maxMipmapLevelClamp&& + tr->__cudaReserved[0] ==texref->__cudaReserved[0]&& + tr->__cudaReserved[1] ==texref->__cudaReserved[1]&& + tr->__cudaReserved[2] ==texref->__cudaReserved[2]&& + tr->__cudaReserved[3] ==texref->__cudaReserved[3]&& + tr->__cudaReserved[4] ==texref->__cudaReserved[4]&& + tr->__cudaReserved[5] ==texref->__cudaReserved[5]&& + tr->__cudaReserved[6] ==texref->__cudaReserved[6]&& + tr->__cudaReserved[7] ==texref->__cudaReserved[7]&& + tr->__cudaReserved[8] ==texref->__cudaReserved[8]&& + tr->__cudaReserved[9] ==texref->__cudaReserved[9]&& + tr->__cudaReserved[10]==texref->__cudaReserved[10]&& + tr->__cudaReserved[11]==texref->__cudaReserved[11]&& + tr->__cudaReserved[12]==texref->__cudaReserved[12]&& + tr->__cudaReserved[13]==texref->__cudaReserved[13]&& + tr->__cudaReserved[14]==texref->__cudaReserved[14]){ + matches++; + t = kv->second; + //return kv->second; } } - assert(false); + printf("matches (texattr) = %d\n", matches); + //assert(matches==1); + return t; + // std::map::const_iterator t=m_TextureRefToAttribute.find(texref); // assert(t != m_TextureRefToAttribute.end()); // return t->second; -- cgit v1.3 From d75898c2bf867ac6ea45594f3da9f18525f2ad6f Mon Sep 17 00:00:00 2001 From: Jonathan Date: Thu, 17 May 2018 16:26:55 -0700 Subject: abstract_hardware_model.h:texrefAttr is different and can be searched via pointer cuda-sim.cc: counts matches between cudaBinTextureToArray texref param (pointer) to texref pointers in m_NameToTextureRef --- src/abstract_hardware_model.h | 94 +++++++++++++++++++++---------------------- src/cuda-sim/cuda-sim.cc | 47 +++++++++++++++++----- 2 files changed, 84 insertions(+), 57 deletions(-) (limited to 'src/abstract_hardware_model.h') diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index c21ff6d..608a7e2 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -662,54 +662,54 @@ public: const struct textureReferenceAttr* get_texattr( const struct textureReference *texref ) const { - int matches = 0; - const struct textureReferenceAttr* t = NULL; - for (std::map::const_iterator kv = m_TextureRefToAttribute.begin(); kv!= m_TextureRefToAttribute.end(); kv ++){ - const struct textureReference* tr = kv->first; - if (tr->normalized==texref->normalized&& - tr->filterMode==texref->filterMode&& - tr->addressMode[0]==texref->addressMode[0]&& - tr->addressMode[1]==texref->addressMode[1]&& - tr->addressMode[2]==texref->addressMode[2]&& - tr->channelDesc.x==texref->channelDesc.x&& - tr->channelDesc.y==texref->channelDesc.y&& - tr->channelDesc.z==texref->channelDesc.z&& - tr->channelDesc.w==texref->channelDesc.w&& - tr->channelDesc.f==texref->channelDesc.f&& - tr->sRGB==texref->sRGB&& - tr->maxAnisotropy==texref->maxAnisotropy&& - tr->mipmapFilterMode==texref->mipmapFilterMode&& - tr->mipmapLevelBias==texref->mipmapLevelBias&& - tr->minMipmapLevelClamp==texref->minMipmapLevelClamp&& - tr->maxMipmapLevelClamp==texref->maxMipmapLevelClamp&& - tr->__cudaReserved[0] ==texref->__cudaReserved[0]&& - tr->__cudaReserved[1] ==texref->__cudaReserved[1]&& - tr->__cudaReserved[2] ==texref->__cudaReserved[2]&& - tr->__cudaReserved[3] ==texref->__cudaReserved[3]&& - tr->__cudaReserved[4] ==texref->__cudaReserved[4]&& - tr->__cudaReserved[5] ==texref->__cudaReserved[5]&& - tr->__cudaReserved[6] ==texref->__cudaReserved[6]&& - tr->__cudaReserved[7] ==texref->__cudaReserved[7]&& - tr->__cudaReserved[8] ==texref->__cudaReserved[8]&& - tr->__cudaReserved[9] ==texref->__cudaReserved[9]&& - tr->__cudaReserved[10]==texref->__cudaReserved[10]&& - tr->__cudaReserved[11]==texref->__cudaReserved[11]&& - tr->__cudaReserved[12]==texref->__cudaReserved[12]&& - tr->__cudaReserved[13]==texref->__cudaReserved[13]&& - tr->__cudaReserved[14]==texref->__cudaReserved[14]){ - matches++; - t = kv->second; - //return kv->second; - } - } - - printf("matches (texattr) = %d\n", matches); - //assert(matches==1); - return t; +// int matches = 0; +// const struct textureReferenceAttr* t = NULL; +// for (std::map::const_iterator kv = m_TextureRefToAttribute.begin(); kv!= m_TextureRefToAttribute.end(); kv ++){ +// const struct textureReference* tr = kv->first; +// if (tr->normalized==texref->normalized&& +// tr->filterMode==texref->filterMode&& +// tr->addressMode[0]==texref->addressMode[0]&& +// tr->addressMode[1]==texref->addressMode[1]&& +// tr->addressMode[2]==texref->addressMode[2]&& +// tr->channelDesc.x==texref->channelDesc.x&& +// tr->channelDesc.y==texref->channelDesc.y&& +// tr->channelDesc.z==texref->channelDesc.z&& +// tr->channelDesc.w==texref->channelDesc.w&& +// tr->channelDesc.f==texref->channelDesc.f&& +// tr->sRGB==texref->sRGB&& +// tr->maxAnisotropy==texref->maxAnisotropy&& +// tr->mipmapFilterMode==texref->mipmapFilterMode&& +// tr->mipmapLevelBias==texref->mipmapLevelBias&& +// tr->minMipmapLevelClamp==texref->minMipmapLevelClamp&& +// tr->maxMipmapLevelClamp==texref->maxMipmapLevelClamp&& +// tr->__cudaReserved[0] ==texref->__cudaReserved[0]&& +// tr->__cudaReserved[1] ==texref->__cudaReserved[1]&& +// tr->__cudaReserved[2] ==texref->__cudaReserved[2]&& +// tr->__cudaReserved[3] ==texref->__cudaReserved[3]&& +// tr->__cudaReserved[4] ==texref->__cudaReserved[4]&& +// tr->__cudaReserved[5] ==texref->__cudaReserved[5]&& +// tr->__cudaReserved[6] ==texref->__cudaReserved[6]&& +// tr->__cudaReserved[7] ==texref->__cudaReserved[7]&& +// tr->__cudaReserved[8] ==texref->__cudaReserved[8]&& +// tr->__cudaReserved[9] ==texref->__cudaReserved[9]&& +// tr->__cudaReserved[10]==texref->__cudaReserved[10]&& +// tr->__cudaReserved[11]==texref->__cudaReserved[11]&& +// tr->__cudaReserved[12]==texref->__cudaReserved[12]&& +// tr->__cudaReserved[13]==texref->__cudaReserved[13]&& +// tr->__cudaReserved[14]==texref->__cudaReserved[14]){ +// matches++; +// t = kv->second; +// //return kv->second; +// } +// } +// +// printf("matches (texattr) = %d\n", matches); +// //assert(matches==1); +// return t; -// std::map::const_iterator t=m_TextureRefToAttribute.find(texref); -// assert(t != m_TextureRefToAttribute.end()); -// return t->second; + std::map::const_iterator t=m_TextureRefToAttribute.find(texref); + assert(t != m_TextureRefToAttribute.end()); + return t->second; } const gpgpu_functional_sim_config &get_config() const { return m_function_model_config; } diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 6e04ca8..39ffa63 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -26,6 +26,10 @@ // OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +#ifndef VERSION_EIGHT +#define VERSION_EIGHT +#endif + #include "cuda-sim.h" #include "instructions.h" @@ -104,6 +108,16 @@ static address_type get_converge_point(address_type pc); void gpgpu_t::gpgpu_ptx_sim_bindNameToTexture(const char* name, const struct textureReference* texref, int dim, int readmode, int ext) { + +//#ifdef VERSION_EIGHT +// int i; +// printf("%s ", name); +// printf("__cudaReserved:"); +// for (i = 0; i<15; i++){ +// printf(" %i", texref->__cudaReserved[i]); +// } +// printf("\n"); +//#endif std::string texname(name); m_NameToTextureRef[texname] = texref; const textureReferenceAttr *texAttr = new textureReferenceAttr(texref, dim, (enum cudaTextureReadMode)readmode, ext); @@ -140,16 +154,15 @@ unsigned int intLOGB2( unsigned int v ) { void gpgpu_t::gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* texref, const struct cudaArray* array) { +//#ifdef VERSION_EIGHT +// int i; +// printf("__cudaReserved:"); +// for (i = 0; i<15; i++){ +// printf(" %i", texref->__cudaReserved[i]); +// } +// printf("\n"); +//#endif // counts number of matches -// int normalized; -// enum cudaTextureFilterMode filterMode; -// enum cudaTextureAddressMode addressMode[3]; -// struct cudaChannelFormatDesc channelDesc; -// int x; -// int y; -// int z; -// int w; -// enum cudaChannelFormatKind f; // int trMatches = 0; // for (auto& kv : m_NameToTextureRef){ // const struct textureReference* tr = kv.second; @@ -170,7 +183,21 @@ void gpgpu_t::gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* te // } // printf("GPGPU-Sim PTX: matches to texref = %d\n", trMatches); // assert(trMatches==1); - m_TextureRefToCudaArray[texref] = array; + + //tests if texref pointer matches any pointer in m_NameToTextureRef map + int trMatches = 0; + for (auto& kv : m_NameToTextureRef){ + const struct textureReference* tr = kv.second; + if (tr==texref){ + m_TextureRefToCudaArray[tr] = array; + //printf("%s\n", kv.first); + trMatches++; + } + } + printf("GPGPU-Sim PTX: matches to texref = %d\n", trMatches); + //assert(trMatches==1); + + //m_TextureRefToCudaArray[texref] = array; unsigned int texel_size_bits = array->desc.w + array->desc.x + array->desc.y + array->desc.z; unsigned int texel_size = texel_size_bits/8; unsigned int Tx, Ty; -- cgit v1.3 From 3b4c3898771ac5e774bca9445a5a4a81670b7b17 Mon Sep 17 00:00:00 2001 From: Jonathan Date: Fri, 18 May 2018 10:50:19 -0700 Subject: reverted cuda8 changes (extra fields) and assert texture bug fix's assumption --- src/abstract_hardware_model.h | 134 ++++++++++++++---------------------------- src/cuda-sim/cuda-sim.cc | 59 +------------------ 2 files changed, 45 insertions(+), 148 deletions(-) (limited to 'src/abstract_hardware_model.h') diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 608a7e2..412c0a8 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -451,33 +451,32 @@ struct textureReference { enum cudaTextureAddressMode addressMode[3]; struct cudaChannelFormatDesc channelDesc; -#ifdef VERSION_EIGHT - /** - * Perform sRGB->linear conversion during texture read - */ - int sRGB; - /** - * Limit to the anisotropy ratio - */ - unsigned int maxAnisotropy; - /** - * Mipmap filter mode - */ - enum cudaTextureFilterMode mipmapFilterMode; - /** - * Offset applied to the supplied mipmap level - */ - float mipmapLevelBias; - /** - * Lower end of the mipmap level range to clamp access to - */ - float minMipmapLevelClamp; - /** - * Upper end of the mipmap level range to clamp access to - */ - float maxMipmapLevelClamp; - int __cudaReserved[15]; -#endif +//following commented section applies only to CUDA_VERSION 8+ +// /** +// * Perform sRGB->linear conversion during texture read +// */ +// int sRGB; +// /** +// * Limit to the anisotropy ratio +// */ +// unsigned int maxAnisotropy; +// /** +// * Mipmap filter mode +// */ +// enum cudaTextureFilterMode mipmapFilterMode; +// /** +// * Offset applied to the supplied mipmap level +// */ +// float mipmapLevelBias; +// /** +// * Lower end of the mipmap level range to clamp access to +// */ +// float minMipmapLevelClamp; +// /** +// * Upper end of the mipmap level range to clamp access to +// */ +// float maxMipmapLevelClamp; +// int __cudaReserved[15]; }; @@ -569,7 +568,10 @@ public: tr->channelDesc.y==texref->channelDesc.y&& tr->channelDesc.z==texref->channelDesc.z&& tr->channelDesc.w==texref->channelDesc.w&& - tr->channelDesc.f==texref->channelDesc.f && + tr->channelDesc.f==texref->channelDesc.f + //following commented section applies only to CUDA_VERSION 8+ + /* + && tr->sRGB==texref->sRGB&& tr->maxAnisotropy==texref->maxAnisotropy&& tr->mipmapFilterMode==texref->mipmapFilterMode&& @@ -591,6 +593,7 @@ public: tr->__cudaReserved[12]==texref->__cudaReserved[12]&& tr->__cudaReserved[13]==texref->__cudaReserved[13]&& tr->__cudaReserved[14]==texref->__cudaReserved[14] + */ ){ matches++; t = kv->second; @@ -598,15 +601,9 @@ public: } } - printf("matches (texarray) = %d\n", matches); - //assert(matches==1); + //printf("matches (texarray) = %d\n", matches); + assert(matches==1); return t; - - //assert(false); - -// std::map::const_iterator t=m_TextureRefToCudaArray.find(texref); -// assert(t != m_TextureRefToCudaArray.end()); -// return t->second; } const struct textureInfo* get_texinfo( const struct textureReference *texref ) const { @@ -623,7 +620,10 @@ public: tr->channelDesc.y==texref->channelDesc.y&& tr->channelDesc.z==texref->channelDesc.z&& tr->channelDesc.w==texref->channelDesc.w&& - tr->channelDesc.f==texref->channelDesc.f&& + tr->channelDesc.f==texref->channelDesc.f + //following commented section applies only to CUDA_VERSION 8+ + /* + && tr->sRGB==texref->sRGB&& tr->maxAnisotropy==texref->maxAnisotropy&& tr->mipmapFilterMode==texref->mipmapFilterMode&& @@ -644,69 +644,23 @@ public: tr->__cudaReserved[11]==texref->__cudaReserved[11]&& tr->__cudaReserved[12]==texref->__cudaReserved[12]&& tr->__cudaReserved[13]==texref->__cudaReserved[13]&& - tr->__cudaReserved[14]==texref->__cudaReserved[14]){ + tr->__cudaReserved[14]==texref->__cudaReserved[14] + */ + ){ matches++; t = kv->second; - //return kv->second; } } - printf("matches (texinfo) = %d\n", matches); - //assert(matches==1); + //printf("matches (texinfo) = %d\n", matches); + assert(matches==1); return t; - -// std::map::const_iterator t=m_TextureRefToTexureInfo.find(texref); -// assert(t != m_TextureRefToTexureInfo.end()); -// return t->second; } const struct textureReferenceAttr* get_texattr( const struct textureReference *texref ) const { -// int matches = 0; -// const struct textureReferenceAttr* t = NULL; -// for (std::map::const_iterator kv = m_TextureRefToAttribute.begin(); kv!= m_TextureRefToAttribute.end(); kv ++){ -// const struct textureReference* tr = kv->first; -// if (tr->normalized==texref->normalized&& -// tr->filterMode==texref->filterMode&& -// tr->addressMode[0]==texref->addressMode[0]&& -// tr->addressMode[1]==texref->addressMode[1]&& -// tr->addressMode[2]==texref->addressMode[2]&& -// tr->channelDesc.x==texref->channelDesc.x&& -// tr->channelDesc.y==texref->channelDesc.y&& -// tr->channelDesc.z==texref->channelDesc.z&& -// tr->channelDesc.w==texref->channelDesc.w&& -// tr->channelDesc.f==texref->channelDesc.f&& -// tr->sRGB==texref->sRGB&& -// tr->maxAnisotropy==texref->maxAnisotropy&& -// tr->mipmapFilterMode==texref->mipmapFilterMode&& -// tr->mipmapLevelBias==texref->mipmapLevelBias&& -// tr->minMipmapLevelClamp==texref->minMipmapLevelClamp&& -// tr->maxMipmapLevelClamp==texref->maxMipmapLevelClamp&& -// tr->__cudaReserved[0] ==texref->__cudaReserved[0]&& -// tr->__cudaReserved[1] ==texref->__cudaReserved[1]&& -// tr->__cudaReserved[2] ==texref->__cudaReserved[2]&& -// tr->__cudaReserved[3] ==texref->__cudaReserved[3]&& -// tr->__cudaReserved[4] ==texref->__cudaReserved[4]&& -// tr->__cudaReserved[5] ==texref->__cudaReserved[5]&& -// tr->__cudaReserved[6] ==texref->__cudaReserved[6]&& -// tr->__cudaReserved[7] ==texref->__cudaReserved[7]&& -// tr->__cudaReserved[8] ==texref->__cudaReserved[8]&& -// tr->__cudaReserved[9] ==texref->__cudaReserved[9]&& -// tr->__cudaReserved[10]==texref->__cudaReserved[10]&& -// tr->__cudaReserved[11]==texref->__cudaReserved[11]&& -// tr->__cudaReserved[12]==texref->__cudaReserved[12]&& -// tr->__cudaReserved[13]==texref->__cudaReserved[13]&& -// tr->__cudaReserved[14]==texref->__cudaReserved[14]){ -// matches++; -// t = kv->second; -// //return kv->second; -// } -// } -// -// printf("matches (texattr) = %d\n", matches); -// //assert(matches==1); -// return t; - + //note textureReferenceAttr map behaves differently from cudaArray and + //textureInfo maps std::map::const_iterator t=m_TextureRefToAttribute.find(texref); assert(t != m_TextureRefToAttribute.end()); return t->second; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 656091c..946043a 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -26,10 +26,6 @@ // OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. -#ifndef VERSION_EIGHT -#define VERSION_EIGHT -#endif - #include "cuda-sim.h" #include "instructions.h" @@ -108,16 +104,6 @@ static address_type get_converge_point(address_type pc); void gpgpu_t::gpgpu_ptx_sim_bindNameToTexture(const char* name, const struct textureReference* texref, int dim, int readmode, int ext) { - -//#ifdef VERSION_EIGHT -// int i; -// printf("%s ", name); -// printf("__cudaReserved:"); -// for (i = 0; i<15; i++){ -// printf(" %i", texref->__cudaReserved[i]); -// } -// printf("\n"); -//#endif std::string texname(name); m_NameToTextureRef[texname] = texref; const textureReferenceAttr *texAttr = new textureReferenceAttr(texref, dim, (enum cudaTextureReadMode)readmode, ext); @@ -154,50 +140,7 @@ unsigned int intLOGB2( unsigned int v ) { void gpgpu_t::gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* texref, const struct cudaArray* array) { -//#ifdef VERSION_EIGHT -// int i; -// printf("__cudaReserved:"); -// for (i = 0; i<15; i++){ -// printf(" %i", texref->__cudaReserved[i]); -// } -// printf("\n"); -//#endif - // counts number of matches -// int trMatches = 0; -// for (auto& kv : m_NameToTextureRef){ -// const struct textureReference* tr = kv.second; -// if (tr->normalized==texref->normalized&& -// tr->filterMode==texref->filterMode&& -// tr->addressMode[0]==texref->addressMode[0]&& -// tr->addressMode[1]==texref->addressMode[1]&& -// tr->addressMode[2]==texref->addressMode[2]&& -// tr->channelDesc.x==texref->channelDesc.x&& -// tr->channelDesc.y==texref->channelDesc.y&& -// tr->channelDesc.z==texref->channelDesc.z&& -// tr->channelDesc.w==texref->channelDesc.w&& -// tr->channelDesc.f==texref->channelDesc.f){ -// -// m_TextureRefToCudaArray[tr] = array; -// trMatches++; -// } -// } -// printf("GPGPU-Sim PTX: matches to texref = %d\n", trMatches); -// assert(trMatches==1); - - //tests if texref pointer matches any pointer in m_NameToTextureRef map - int trMatches = 0; - for (std::map::const_iterator kv = m_NameToTextureRef.begin(); kv!= m_NameToTextureRef.end(); kv ++){ - const struct textureReference* tr = kv->second; - if (tr==texref){ - m_TextureRefToCudaArray[tr] = array; - printf("%s\n", kv->first.c_str()); - trMatches++; - } - } - printf("GPGPU-Sim PTX: matches to texref = %d\n", trMatches); - //assert(trMatches==1); - - //m_TextureRefToCudaArray[texref] = array; + m_TextureRefToCudaArray[texref] = array; unsigned int texel_size_bits = array->desc.w + array->desc.x + array->desc.y + array->desc.z; unsigned int texel_size = texel_size_bits/8; unsigned int Tx, Ty; -- cgit v1.3 From 6a39e5c5964f23a97dafaa6a66f2a9d9c37bbfdd Mon Sep 17 00:00:00 2001 From: Jonathan Date: Fri, 18 May 2018 11:52:31 -0700 Subject: implemented unbind, currently only affects cudaArray map --- 0517_13-24cudaReserved.log | 0 517-1418cudaReserved.log | 0 libcuda/cuda_runtime_api.cc | 11 +++++++++++ src/abstract_hardware_model.h | 1 + src/cuda-sim/cuda-sim.cc | 5 +++++ 5 files changed, 17 insertions(+) delete mode 100644 0517_13-24cudaReserved.log delete mode 100644 517-1418cudaReserved.log (limited to 'src/abstract_hardware_model.h') diff --git a/0517_13-24cudaReserved.log b/0517_13-24cudaReserved.log deleted file mode 100644 index e69de29..0000000 diff --git a/517-1418cudaReserved.log b/517-1418cudaReserved.log deleted file mode 100644 index e69de29..0000000 diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index e3c2542..71926f8 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -991,6 +991,14 @@ __host__ cudaError_t CUDARTAPI cudaBindTextureToArray(const struct textureRefere __host__ cudaError_t CUDARTAPI cudaUnbindTexture(const struct textureReference *texref) { + CUctx_st *context = GPGPUSim_Context(); + gpgpu_t *gpu = context->get_device()->get_gpgpu(); + printf("GPGPU-Sim PTX: in cudaUnbindTexture: sizeof(struct textureReference) = %zu\n", sizeof(struct textureReference)); + printf("GPGPU-Sim PTX: Name corresponding to textureReference: %s\n", gpu->gpgpu_ptx_sim_findNamefromTexture(texref)); + const struct cudaArray* array = gpu->get_texarray(texref); + printf("GPGPU-Sim PTX: texref = %p, array = %p\n", texref, array); + + gpu->gpgpu_ptx_sim_unbindTexture(texref); return g_last_cudaError = cudaSuccess; } @@ -2073,10 +2081,12 @@ void __cudaUnregisterFatBinary(void **fatCubinHandle) cudaError_t cudaDeviceReset ( void ) { // Should reset the simulated GPU + // TODO: Implement return g_last_cudaError = cudaSuccess; } cudaError_t CUDARTAPI cudaDeviceSynchronize(void){ // I don't know what this should do + // TODO: Implement return g_last_cudaError = cudaSuccess; } @@ -2178,6 +2188,7 @@ typedef unsigned long GLuint; cudaError_t cudaGLRegisterBufferObject(GLuint bufferObj) { printf("GPGPU-Sim PTX: Execution warning: ignoring call to \"%s\"\n", __my_func__ ); + // TODO: Implement return g_last_cudaError = cudaSuccess; } diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 412c0a8..3ef450e 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -544,6 +544,7 @@ public: class memory_space *get_surf_memory() { return m_surf_mem; } void gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* texref, const struct cudaArray* array); + void gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref); void gpgpu_ptx_sim_bindNameToTexture(const char* name, const struct textureReference* texref, int dim, int readmode, int ext); const char* gpgpu_ptx_sim_findNamefromTexture(const struct textureReference* texref); diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 946043a..ef16f43 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -183,6 +183,11 @@ void gpgpu_t::gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* te m_TextureRefToTexureInfo[texref] = texInfo; } +void gpgpu_t::gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref) +{ + m_TextureRefToCudaArray.erase(texref); +} + unsigned g_assemble_code_next_pc=0; std::map g_pc_to_finfo; std::vector function_info::s_g_pc_to_insn; -- cgit v1.3 From a93aa28a94140ee912c0cba0d9414d9da1588d54 Mon Sep 17 00:00:00 2001 From: Jonathan Date: Fri, 18 May 2018 14:41:23 -0700 Subject: erase texinfo in unbind and disable assert --- src/abstract_hardware_model.h | 195 ++++++++++++++++++++++-------------------- src/cuda-sim/cuda-sim.cc | 2 + 2 files changed, 104 insertions(+), 93 deletions(-) (limited to 'src/abstract_hardware_model.h') diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 3ef450e..6dd5436 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -556,106 +556,115 @@ public: } const struct cudaArray* get_texarray( const struct textureReference *texref ) const { - int matches = 0; - const struct cudaArray* t = NULL; - for (std::map::const_iterator kv = m_TextureRefToCudaArray.begin(); kv!= m_TextureRefToCudaArray.end(); kv ++){ - const struct textureReference* tr = kv->first; - if (tr->normalized==texref->normalized&& - tr->filterMode==texref->filterMode&& - tr->addressMode[0]==texref->addressMode[0]&& - tr->addressMode[1]==texref->addressMode[1]&& - tr->addressMode[2]==texref->addressMode[2]&& - tr->channelDesc.x==texref->channelDesc.x&& - tr->channelDesc.y==texref->channelDesc.y&& - tr->channelDesc.z==texref->channelDesc.z&& - tr->channelDesc.w==texref->channelDesc.w&& - tr->channelDesc.f==texref->channelDesc.f - //following commented section applies only to CUDA_VERSION 8+ - /* - && - tr->sRGB==texref->sRGB&& - tr->maxAnisotropy==texref->maxAnisotropy&& - tr->mipmapFilterMode==texref->mipmapFilterMode&& - tr->mipmapLevelBias==texref->mipmapLevelBias&& - tr->minMipmapLevelClamp==texref->minMipmapLevelClamp&& - tr->maxMipmapLevelClamp==texref->maxMipmapLevelClamp&& - tr->__cudaReserved[0] ==texref->__cudaReserved[0]&& - tr->__cudaReserved[1] ==texref->__cudaReserved[1]&& - tr->__cudaReserved[2] ==texref->__cudaReserved[2]&& - tr->__cudaReserved[3] ==texref->__cudaReserved[3]&& - tr->__cudaReserved[4] ==texref->__cudaReserved[4]&& - tr->__cudaReserved[5] ==texref->__cudaReserved[5]&& - tr->__cudaReserved[6] ==texref->__cudaReserved[6]&& - tr->__cudaReserved[7] ==texref->__cudaReserved[7]&& - tr->__cudaReserved[8] ==texref->__cudaReserved[8]&& - tr->__cudaReserved[9] ==texref->__cudaReserved[9]&& - tr->__cudaReserved[10]==texref->__cudaReserved[10]&& - tr->__cudaReserved[11]==texref->__cudaReserved[11]&& - tr->__cudaReserved[12]==texref->__cudaReserved[12]&& - tr->__cudaReserved[13]==texref->__cudaReserved[13]&& - tr->__cudaReserved[14]==texref->__cudaReserved[14] - */ - ){ - matches++; - t = kv->second; - //return kv->second; + + std::map::const_iterator t=m_TextureRefToCudaArray.find(texref); + if(t != m_TextureRefToCudaArray.end()){ + return t->second; + } else{ + int matches = 0; + const struct cudaArray* t = NULL; + for (std::map::const_iterator kv = m_TextureRefToCudaArray.begin(); + kv!= m_TextureRefToCudaArray.end(); kv ++){ + const struct textureReference* tr = kv->first; + if (tr->normalized==texref->normalized&& + tr->filterMode==texref->filterMode&& + tr->addressMode[0]==texref->addressMode[0]&& + tr->addressMode[1]==texref->addressMode[1]&& + tr->addressMode[2]==texref->addressMode[2]&& + tr->channelDesc.x==texref->channelDesc.x&& + tr->channelDesc.y==texref->channelDesc.y&& + tr->channelDesc.z==texref->channelDesc.z&& + tr->channelDesc.w==texref->channelDesc.w&& + tr->channelDesc.f==texref->channelDesc.f + //following commented section applies only to CUDA_VERSION 8+ + /* + && + tr->sRGB==texref->sRGB&& + tr->maxAnisotropy==texref->maxAnisotropy&& + tr->mipmapFilterMode==texref->mipmapFilterMode&& + tr->mipmapLevelBias==texref->mipmapLevelBias&& + tr->minMipmapLevelClamp==texref->minMipmapLevelClamp&& + tr->maxMipmapLevelClamp==texref->maxMipmapLevelClamp&& + tr->__cudaReserved[0] ==texref->__cudaReserved[0]&& + tr->__cudaReserved[1] ==texref->__cudaReserved[1]&& + tr->__cudaReserved[2] ==texref->__cudaReserved[2]&& + tr->__cudaReserved[3] ==texref->__cudaReserved[3]&& + tr->__cudaReserved[4] ==texref->__cudaReserved[4]&& + tr->__cudaReserved[5] ==texref->__cudaReserved[5]&& + tr->__cudaReserved[6] ==texref->__cudaReserved[6]&& + tr->__cudaReserved[7] ==texref->__cudaReserved[7]&& + tr->__cudaReserved[8] ==texref->__cudaReserved[8]&& + tr->__cudaReserved[9] ==texref->__cudaReserved[9]&& + tr->__cudaReserved[10]==texref->__cudaReserved[10]&& + tr->__cudaReserved[11]==texref->__cudaReserved[11]&& + tr->__cudaReserved[12]==texref->__cudaReserved[12]&& + tr->__cudaReserved[13]==texref->__cudaReserved[13]&& + tr->__cudaReserved[14]==texref->__cudaReserved[14] + */ + ){ + matches++; + t = kv->second; + } } + + //assert(matches==1); + return t; } - - //printf("matches (texarray) = %d\n", matches); - assert(matches==1); - return t; } const struct textureInfo* get_texinfo( const struct textureReference *texref ) const { - int matches = 0; - const struct textureInfo* t = NULL; - for (std::map::const_iterator kv = m_TextureRefToTexureInfo.begin(); kv!= m_TextureRefToTexureInfo.end(); kv ++){ - const struct textureReference* tr = kv->first; - if (tr->normalized==texref->normalized&& - tr->filterMode==texref->filterMode&& - tr->addressMode[0]==texref->addressMode[0]&& - tr->addressMode[1]==texref->addressMode[1]&& - tr->addressMode[2]==texref->addressMode[2]&& - tr->channelDesc.x==texref->channelDesc.x&& - tr->channelDesc.y==texref->channelDesc.y&& - tr->channelDesc.z==texref->channelDesc.z&& - tr->channelDesc.w==texref->channelDesc.w&& - tr->channelDesc.f==texref->channelDesc.f - //following commented section applies only to CUDA_VERSION 8+ - /* - && - tr->sRGB==texref->sRGB&& - tr->maxAnisotropy==texref->maxAnisotropy&& - tr->mipmapFilterMode==texref->mipmapFilterMode&& - tr->mipmapLevelBias==texref->mipmapLevelBias&& - tr->minMipmapLevelClamp==texref->minMipmapLevelClamp&& - tr->maxMipmapLevelClamp==texref->maxMipmapLevelClamp&& - tr->__cudaReserved[0] ==texref->__cudaReserved[0]&& - tr->__cudaReserved[1] ==texref->__cudaReserved[1]&& - tr->__cudaReserved[2] ==texref->__cudaReserved[2]&& - tr->__cudaReserved[3] ==texref->__cudaReserved[3]&& - tr->__cudaReserved[4] ==texref->__cudaReserved[4]&& - tr->__cudaReserved[5] ==texref->__cudaReserved[5]&& - tr->__cudaReserved[6] ==texref->__cudaReserved[6]&& - tr->__cudaReserved[7] ==texref->__cudaReserved[7]&& - tr->__cudaReserved[8] ==texref->__cudaReserved[8]&& - tr->__cudaReserved[9] ==texref->__cudaReserved[9]&& - tr->__cudaReserved[10]==texref->__cudaReserved[10]&& - tr->__cudaReserved[11]==texref->__cudaReserved[11]&& - tr->__cudaReserved[12]==texref->__cudaReserved[12]&& - tr->__cudaReserved[13]==texref->__cudaReserved[13]&& - tr->__cudaReserved[14]==texref->__cudaReserved[14] - */ - ){ - matches++; - t = kv->second; + std::map::const_iterator t=m_TextureRefToTexureInfo.find(texref); + if(t != m_TextureRefToTexureInfo.end()){ + return t->second; + }else{ + int matches = 0; + const struct textureInfo* t = NULL; + for (std::map::const_iterator kv = m_TextureRefToTexureInfo.begin(); + kv!= m_TextureRefToTexureInfo.end(); kv ++){ + const struct textureReference* tr = kv->first; + if (tr->normalized==texref->normalized&& + tr->filterMode==texref->filterMode&& + tr->addressMode[0]==texref->addressMode[0]&& + tr->addressMode[1]==texref->addressMode[1]&& + tr->addressMode[2]==texref->addressMode[2]&& + tr->channelDesc.x==texref->channelDesc.x&& + tr->channelDesc.y==texref->channelDesc.y&& + tr->channelDesc.z==texref->channelDesc.z&& + tr->channelDesc.w==texref->channelDesc.w&& + tr->channelDesc.f==texref->channelDesc.f + //following commented section applies only to CUDA_VERSION 8+ + /* + && + tr->sRGB==texref->sRGB&& + tr->maxAnisotropy==texref->maxAnisotropy&& + tr->mipmapFilterMode==texref->mipmapFilterMode&& + tr->mipmapLevelBias==texref->mipmapLevelBias&& + tr->minMipmapLevelClamp==texref->minMipmapLevelClamp&& + tr->maxMipmapLevelClamp==texref->maxMipmapLevelClamp&& + tr->__cudaReserved[0] ==texref->__cudaReserved[0]&& + tr->__cudaReserved[1] ==texref->__cudaReserved[1]&& + tr->__cudaReserved[2] ==texref->__cudaReserved[2]&& + tr->__cudaReserved[3] ==texref->__cudaReserved[3]&& + tr->__cudaReserved[4] ==texref->__cudaReserved[4]&& + tr->__cudaReserved[5] ==texref->__cudaReserved[5]&& + tr->__cudaReserved[6] ==texref->__cudaReserved[6]&& + tr->__cudaReserved[7] ==texref->__cudaReserved[7]&& + tr->__cudaReserved[8] ==texref->__cudaReserved[8]&& + tr->__cudaReserved[9] ==texref->__cudaReserved[9]&& + tr->__cudaReserved[10]==texref->__cudaReserved[10]&& + tr->__cudaReserved[11]==texref->__cudaReserved[11]&& + tr->__cudaReserved[12]==texref->__cudaReserved[12]&& + tr->__cudaReserved[13]==texref->__cudaReserved[13]&& + tr->__cudaReserved[14]==texref->__cudaReserved[14] + */ + ){ + matches++; + t = kv->second; + } } + //assert(matches==1); + return t; } - - //printf("matches (texinfo) = %d\n", matches); - assert(matches==1); - return t; } const struct textureReferenceAttr* get_texattr( const struct textureReference *texref ) const diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index ef16f43..6bdf75f 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -140,6 +140,7 @@ unsigned int intLOGB2( unsigned int v ) { void gpgpu_t::gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* texref, const struct cudaArray* array) { + printf("GPGPU-Simm PTX: name from texture = %s\n", gpgpu_ptx_sim_findNamefromTexture(texref)); m_TextureRefToCudaArray[texref] = array; unsigned int texel_size_bits = array->desc.w + array->desc.x + array->desc.y + array->desc.z; unsigned int texel_size = texel_size_bits/8; @@ -186,6 +187,7 @@ void gpgpu_t::gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* te void gpgpu_t::gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref) { m_TextureRefToCudaArray.erase(texref); + m_TextureRefToTexureInfo.erase(texref); } unsigned g_assemble_code_next_pc=0; -- cgit v1.3 From 958e430266cb3de73033a7e0aab6e7c697fdc6bc Mon Sep 17 00:00:00 2001 From: Jonathan Date: Tue, 22 May 2018 09:49:26 -0700 Subject: revert to before texture bug --- libcuda/cuda_runtime_api.cc | 11 ---- src/abstract_hardware_model.h | 144 ++---------------------------------------- src/cuda-sim/cuda-sim.cc | 7 -- 3 files changed, 4 insertions(+), 158 deletions(-) (limited to 'src/abstract_hardware_model.h') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 71926f8..e3c2542 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -991,14 +991,6 @@ __host__ cudaError_t CUDARTAPI cudaBindTextureToArray(const struct textureRefere __host__ cudaError_t CUDARTAPI cudaUnbindTexture(const struct textureReference *texref) { - CUctx_st *context = GPGPUSim_Context(); - gpgpu_t *gpu = context->get_device()->get_gpgpu(); - printf("GPGPU-Sim PTX: in cudaUnbindTexture: sizeof(struct textureReference) = %zu\n", sizeof(struct textureReference)); - printf("GPGPU-Sim PTX: Name corresponding to textureReference: %s\n", gpu->gpgpu_ptx_sim_findNamefromTexture(texref)); - const struct cudaArray* array = gpu->get_texarray(texref); - printf("GPGPU-Sim PTX: texref = %p, array = %p\n", texref, array); - - gpu->gpgpu_ptx_sim_unbindTexture(texref); return g_last_cudaError = cudaSuccess; } @@ -2081,12 +2073,10 @@ void __cudaUnregisterFatBinary(void **fatCubinHandle) cudaError_t cudaDeviceReset ( void ) { // Should reset the simulated GPU - // TODO: Implement return g_last_cudaError = cudaSuccess; } cudaError_t CUDARTAPI cudaDeviceSynchronize(void){ // I don't know what this should do - // TODO: Implement return g_last_cudaError = cudaSuccess; } @@ -2188,7 +2178,6 @@ typedef unsigned long GLuint; cudaError_t cudaGLRegisterBufferObject(GLuint bufferObj) { printf("GPGPU-Sim PTX: Execution warning: ignoring call to \"%s\"\n", __my_func__ ); - // TODO: Implement return g_last_cudaError = cudaSuccess; } diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 6dd5436..f5708bc 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -25,11 +25,6 @@ // OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. - -#ifndef VERSION_EIGHT -#define VERSION_EIGHT -#endif - #ifndef ABSTRACT_HARDWARE_MODEL_INCLUDED #define ABSTRACT_HARDWARE_MODEL_INCLUDED @@ -450,34 +445,6 @@ struct textureReference { enum cudaTextureFilterMode filterMode; enum cudaTextureAddressMode addressMode[3]; struct cudaChannelFormatDesc channelDesc; - -//following commented section applies only to CUDA_VERSION 8+ -// /** -// * Perform sRGB->linear conversion during texture read -// */ -// int sRGB; -// /** -// * Limit to the anisotropy ratio -// */ -// unsigned int maxAnisotropy; -// /** -// * Mipmap filter mode -// */ -// enum cudaTextureFilterMode mipmapFilterMode; -// /** -// * Offset applied to the supplied mipmap level -// */ -// float mipmapLevelBias; -// /** -// * Lower end of the mipmap level range to clamp access to -// */ -// float minMipmapLevelClamp; -// /** -// * Upper end of the mipmap level range to clamp access to -// */ -// float maxMipmapLevelClamp; -// int __cudaReserved[15]; - }; #endif @@ -544,7 +511,6 @@ public: class memory_space *get_surf_memory() { return m_surf_mem; } void gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* texref, const struct cudaArray* array); - void gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref); void gpgpu_ptx_sim_bindNameToTexture(const char* name, const struct textureReference* texref, int dim, int readmode, int ext); const char* gpgpu_ptx_sim_findNamefromTexture(const struct textureReference* texref); @@ -556,121 +522,19 @@ public: } const struct cudaArray* get_texarray( const struct textureReference *texref ) const { - std::map::const_iterator t=m_TextureRefToCudaArray.find(texref); - if(t != m_TextureRefToCudaArray.end()){ - return t->second; - } else{ - int matches = 0; - const struct cudaArray* t = NULL; - for (std::map::const_iterator kv = m_TextureRefToCudaArray.begin(); - kv!= m_TextureRefToCudaArray.end(); kv ++){ - const struct textureReference* tr = kv->first; - if (tr->normalized==texref->normalized&& - tr->filterMode==texref->filterMode&& - tr->addressMode[0]==texref->addressMode[0]&& - tr->addressMode[1]==texref->addressMode[1]&& - tr->addressMode[2]==texref->addressMode[2]&& - tr->channelDesc.x==texref->channelDesc.x&& - tr->channelDesc.y==texref->channelDesc.y&& - tr->channelDesc.z==texref->channelDesc.z&& - tr->channelDesc.w==texref->channelDesc.w&& - tr->channelDesc.f==texref->channelDesc.f - //following commented section applies only to CUDA_VERSION 8+ - /* - && - tr->sRGB==texref->sRGB&& - tr->maxAnisotropy==texref->maxAnisotropy&& - tr->mipmapFilterMode==texref->mipmapFilterMode&& - tr->mipmapLevelBias==texref->mipmapLevelBias&& - tr->minMipmapLevelClamp==texref->minMipmapLevelClamp&& - tr->maxMipmapLevelClamp==texref->maxMipmapLevelClamp&& - tr->__cudaReserved[0] ==texref->__cudaReserved[0]&& - tr->__cudaReserved[1] ==texref->__cudaReserved[1]&& - tr->__cudaReserved[2] ==texref->__cudaReserved[2]&& - tr->__cudaReserved[3] ==texref->__cudaReserved[3]&& - tr->__cudaReserved[4] ==texref->__cudaReserved[4]&& - tr->__cudaReserved[5] ==texref->__cudaReserved[5]&& - tr->__cudaReserved[6] ==texref->__cudaReserved[6]&& - tr->__cudaReserved[7] ==texref->__cudaReserved[7]&& - tr->__cudaReserved[8] ==texref->__cudaReserved[8]&& - tr->__cudaReserved[9] ==texref->__cudaReserved[9]&& - tr->__cudaReserved[10]==texref->__cudaReserved[10]&& - tr->__cudaReserved[11]==texref->__cudaReserved[11]&& - tr->__cudaReserved[12]==texref->__cudaReserved[12]&& - tr->__cudaReserved[13]==texref->__cudaReserved[13]&& - tr->__cudaReserved[14]==texref->__cudaReserved[14] - */ - ){ - matches++; - t = kv->second; - } - } - - //assert(matches==1); - return t; - } + assert(t != m_TextureRefToCudaArray.end()); + return t->second; } const struct textureInfo* get_texinfo( const struct textureReference *texref ) const { std::map::const_iterator t=m_TextureRefToTexureInfo.find(texref); - if(t != m_TextureRefToTexureInfo.end()){ - return t->second; - }else{ - int matches = 0; - const struct textureInfo* t = NULL; - for (std::map::const_iterator kv = m_TextureRefToTexureInfo.begin(); - kv!= m_TextureRefToTexureInfo.end(); kv ++){ - const struct textureReference* tr = kv->first; - if (tr->normalized==texref->normalized&& - tr->filterMode==texref->filterMode&& - tr->addressMode[0]==texref->addressMode[0]&& - tr->addressMode[1]==texref->addressMode[1]&& - tr->addressMode[2]==texref->addressMode[2]&& - tr->channelDesc.x==texref->channelDesc.x&& - tr->channelDesc.y==texref->channelDesc.y&& - tr->channelDesc.z==texref->channelDesc.z&& - tr->channelDesc.w==texref->channelDesc.w&& - tr->channelDesc.f==texref->channelDesc.f - //following commented section applies only to CUDA_VERSION 8+ - /* - && - tr->sRGB==texref->sRGB&& - tr->maxAnisotropy==texref->maxAnisotropy&& - tr->mipmapFilterMode==texref->mipmapFilterMode&& - tr->mipmapLevelBias==texref->mipmapLevelBias&& - tr->minMipmapLevelClamp==texref->minMipmapLevelClamp&& - tr->maxMipmapLevelClamp==texref->maxMipmapLevelClamp&& - tr->__cudaReserved[0] ==texref->__cudaReserved[0]&& - tr->__cudaReserved[1] ==texref->__cudaReserved[1]&& - tr->__cudaReserved[2] ==texref->__cudaReserved[2]&& - tr->__cudaReserved[3] ==texref->__cudaReserved[3]&& - tr->__cudaReserved[4] ==texref->__cudaReserved[4]&& - tr->__cudaReserved[5] ==texref->__cudaReserved[5]&& - tr->__cudaReserved[6] ==texref->__cudaReserved[6]&& - tr->__cudaReserved[7] ==texref->__cudaReserved[7]&& - tr->__cudaReserved[8] ==texref->__cudaReserved[8]&& - tr->__cudaReserved[9] ==texref->__cudaReserved[9]&& - tr->__cudaReserved[10]==texref->__cudaReserved[10]&& - tr->__cudaReserved[11]==texref->__cudaReserved[11]&& - tr->__cudaReserved[12]==texref->__cudaReserved[12]&& - tr->__cudaReserved[13]==texref->__cudaReserved[13]&& - tr->__cudaReserved[14]==texref->__cudaReserved[14] - */ - ){ - matches++; - t = kv->second; - } - } - //assert(matches==1); - return t; - } + assert(t != m_TextureRefToTexureInfo.end()); + return t->second; } const struct textureReferenceAttr* get_texattr( const struct textureReference *texref ) const { - //note textureReferenceAttr map behaves differently from cudaArray and - //textureInfo maps std::map::const_iterator t=m_TextureRefToAttribute.find(texref); assert(t != m_TextureRefToAttribute.end()); return t->second; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 6bdf75f..946043a 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -140,7 +140,6 @@ unsigned int intLOGB2( unsigned int v ) { void gpgpu_t::gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* texref, const struct cudaArray* array) { - printf("GPGPU-Simm PTX: name from texture = %s\n", gpgpu_ptx_sim_findNamefromTexture(texref)); m_TextureRefToCudaArray[texref] = array; unsigned int texel_size_bits = array->desc.w + array->desc.x + array->desc.y + array->desc.z; unsigned int texel_size = texel_size_bits/8; @@ -184,12 +183,6 @@ void gpgpu_t::gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* te m_TextureRefToTexureInfo[texref] = texInfo; } -void gpgpu_t::gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref) -{ - m_TextureRefToCudaArray.erase(texref); - m_TextureRefToTexureInfo.erase(texref); -} - unsigned g_assemble_code_next_pc=0; std::map g_pc_to_finfo; std::vector function_info::s_g_pc_to_insn; -- cgit v1.3 From 4a567a94aca58fdd2298b50ae9e9da3a889e2173 Mon Sep 17 00:00:00 2001 From: Jonathan Date: Tue, 22 May 2018 09:56:07 -0700 Subject: unbind implementation --- libcuda/cuda_runtime_api.cc | 11 ++++++++--- src/abstract_hardware_model.h | 1 + src/cuda-sim/cuda-sim.cc | 6 ++++++ 3 files changed, 15 insertions(+), 3 deletions(-) (limited to 'src/abstract_hardware_model.h') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index e3c2542..d971222 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -989,9 +989,14 @@ __host__ cudaError_t CUDARTAPI cudaBindTextureToArray(const struct textureRefere return g_last_cudaError = cudaSuccess; } -__host__ cudaError_t CUDARTAPI cudaUnbindTexture(const struct textureReference *texref) -{ - return g_last_cudaError = cudaSuccess; +__host__ cudaError_t CUDARTAPI cudaUnbindTexture(const struct textureReference *texref) { + CUctx_st *context = GPGPUSim_Context(); + gpgpu_t *gpu = context->get_device()->get_gpgpu(); + printf("GPGPU-Sim PTX: in cudaUnbindTexture: sizeof(struct textureReference) = %zu\n", sizeof(struct textureReference)); + printf("GPGPU-Sim PTX: Name corresponding to textureReference: %s\n", gpu->gpgpu_ptx_sim_findNamefromTexture(texref)); + + gpu->gpgpu_ptx_sim_unbindTexture(texref); + return g_last_cudaError = cudaSuccess; } __host__ cudaError_t CUDARTAPI cudaGetTextureAlignmentOffset(size_t *offset, const struct textureReference *texref) diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index f5708bc..ab94ded 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -512,6 +512,7 @@ public: void gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* texref, const struct cudaArray* array); void gpgpu_ptx_sim_bindNameToTexture(const char* name, const struct textureReference* texref, int dim, int readmode, int ext); + void gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref); const char* gpgpu_ptx_sim_findNamefromTexture(const struct textureReference* texref); const struct textureReference* get_texref(const std::string &texname) const diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 946043a..6125422 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -183,6 +183,12 @@ void gpgpu_t::gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* te m_TextureRefToTexureInfo[texref] = texInfo; } +void gpgpu_t::gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref) +{ + m_TextureRefToCudaArray.erase(texref); + m_TextureRefToTexureInfo.erase(texref); +} + unsigned g_assemble_code_next_pc=0; std::map g_pc_to_finfo; std::vector function_info::s_g_pc_to_insn; -- cgit v1.3 From ff958f36689c9c217eb099326ceb8f70ed3ac447 Mon Sep 17 00:00:00 2001 From: Jonathan Date: Tue, 22 May 2018 14:00:59 -0700 Subject: restructured texref maps to fix texture bug --- src/abstract_hardware_model.h | 36 +++++++++++++++++-------------- src/cuda-sim/cuda-sim.cc | 49 ++++++++++++++++++++++++++++--------------- src/cuda-sim/instructions.cc | 6 +++--- 3 files changed, 55 insertions(+), 36 deletions(-) (limited to 'src/abstract_hardware_model.h') diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index ab94ded..d0af1ea 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -65,6 +65,7 @@ enum FuncCache #include #include +#include typedef unsigned long long new_addr_type; typedef unsigned address_type; @@ -515,29 +516,31 @@ public: void gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref); const char* gpgpu_ptx_sim_findNamefromTexture(const struct textureReference* texref); - const struct textureReference* get_texref(const std::string &texname) const + const struct textureReference* get_texref( const std::string &texname ) const { - std::map::const_iterator t=m_NameToTextureRef.find(texname); + std::map >::const_iterator t=m_NameToTextureRef.find(texname); assert( t != m_NameToTextureRef.end() ); - return t->second; + return *(t->second.begin()); } - const struct cudaArray* get_texarray( const struct textureReference *texref ) const + + const struct cudaArray* get_texarray( const std::string &texname ) const { - std::map::const_iterator t=m_TextureRefToCudaArray.find(texref); - assert(t != m_TextureRefToCudaArray.end()); + std::map::const_iterator t=m_NameToCudaArray.find(texname); + assert(t != m_NameToCudaArray.end()); return t->second; } - const struct textureInfo* get_texinfo( const struct textureReference *texref ) const + + const struct textureInfo* get_texinfo( const std::string &texname ) const { - std::map::const_iterator t=m_TextureRefToTexureInfo.find(texref); - assert(t != m_TextureRefToTexureInfo.end()); + std::map::const_iterator t=m_NameToTexureInfo.find(texname); + assert(t != m_NameToTexureInfo.end()); return t->second; } - const struct textureReferenceAttr* get_texattr( const struct textureReference *texref ) const + const struct textureReferenceAttr* get_texattr( const std::string &texname ) const { - std::map::const_iterator t=m_TextureRefToAttribute.find(texref); - assert(t != m_TextureRefToAttribute.end()); + std::map::const_iterator t=m_NameToAttribute.find(texname); + assert(t != m_NameToAttribute.end()); return t->second; } @@ -554,10 +557,11 @@ protected: unsigned long long m_dev_malloc; - std::map m_NameToTextureRef; - std::map m_TextureRefToCudaArray; - std::map m_TextureRefToTexureInfo; - std::map m_TextureRefToAttribute; + std::map > m_NameToTextureRef; + std::map m_TextureRefToName; + std::map m_NameToCudaArray; + std::map m_NameToTexureInfo; + std::map m_NameToAttribute; }; struct gpgpu_ptx_sim_info diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 6125422..05b6201 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -105,22 +105,36 @@ static address_type get_converge_point(address_type pc); void gpgpu_t::gpgpu_ptx_sim_bindNameToTexture(const char* name, const struct textureReference* texref, int dim, int readmode, int ext) { std::string texname(name); - m_NameToTextureRef[texname] = texref; + if (m_NameToTextureRef.find(texname)==m_NameToTextureRef.end()){ + m_NameToTextureRef[texname] = std::set(); + m_NameToTextureRef[texname].insert(texref); + }else{ + const struct textureReference* tr = *m_NameToTextureRef[texname].begin(); + assert(tr!=NULL); + //asserts that all texrefs in set have same fields + assert(tr->normalized==texref->normalized&& + tr->filterMode==texref->filterMode&& + tr->addressMode[0]==texref->addressMode[0]&& + tr->addressMode[1]==texref->addressMode[1]&& + tr->addressMode[2]==texref->addressMode[2]&& + tr->channelDesc.x==texref->channelDesc.x&& + tr->channelDesc.y==texref->channelDesc.y&& + tr->channelDesc.z==texref->channelDesc.z&& + tr->channelDesc.w==texref->channelDesc.w&& + tr->channelDesc.f==texref->channelDesc.f + ); + m_NameToTextureRef[texname].insert(texref); + } + m_TextureRefToName[texref] = texname; const textureReferenceAttr *texAttr = new textureReferenceAttr(texref, dim, (enum cudaTextureReadMode)readmode, ext); - m_TextureRefToAttribute[texref] = texAttr; + m_NameToAttribute[texname] = texAttr; } const char* gpgpu_t::gpgpu_ptx_sim_findNamefromTexture(const struct textureReference* texref) { - std::map::iterator itr = m_NameToTextureRef.begin(); - while (itr != m_NameToTextureRef.end()) { - if ((*itr).second == texref) { - const char *p = ((*itr).first).c_str(); - return p; - } - itr++; - } - return NULL; + std::map::const_iterator t=m_TextureRefToName.find(texref); + assert( t != m_TextureRefToName.end() ); + return t->second.c_str(); } unsigned int intLOGB2( unsigned int v ) { @@ -140,7 +154,8 @@ unsigned int intLOGB2( unsigned int v ) { void gpgpu_t::gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* texref, const struct cudaArray* array) { - m_TextureRefToCudaArray[texref] = array; + std::string texname = gpgpu_ptx_sim_findNamefromTexture(texref); + m_NameToCudaArray[texname] = array; unsigned int texel_size_bits = array->desc.w + array->desc.x + array->desc.y + array->desc.z; unsigned int texel_size = texel_size_bits/8; unsigned int Tx, Ty; @@ -180,13 +195,14 @@ void gpgpu_t::gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* te texInfo->Ty_numbits = intLOGB2(Ty); texInfo->texel_size = texel_size; texInfo->texel_size_numbits = intLOGB2(texel_size); - m_TextureRefToTexureInfo[texref] = texInfo; + m_NameToTexureInfo[texname] = texInfo; } void gpgpu_t::gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref) { - m_TextureRefToCudaArray.erase(texref); - m_TextureRefToTexureInfo.erase(texref); + std::string texname = gpgpu_ptx_sim_findNamefromTexture(texref); + m_NameToCudaArray.erase(texname); + m_NameToTexureInfo.erase(texname); } unsigned g_assemble_code_next_pc=0; @@ -1246,8 +1262,7 @@ static unsigned get_tex_datasize( const ptx_instruction *pI, ptx_thread_info *th std::string texname = src1.name(); gpgpu_t *gpu = thread->get_gpu(); - const struct textureReference* texref = gpu->get_texref(texname); - const struct textureInfo* texInfo = gpu->get_texinfo(texref); + const struct textureInfo* texInfo = gpu->get_texinfo(texname); unsigned data_size = texInfo->texel_size; return data_size; diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 86951ed..d362231 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -4100,9 +4100,9 @@ void tex_impl( const ptx_instruction *pI, ptx_thread_info *thread ) gpgpu_t *gpu = thread->get_gpu(); const struct textureReference* texref = gpu->get_texref(texname); - const struct cudaArray* cuArray = gpu->get_texarray(texref); - const struct textureInfo* texInfo = gpu->get_texinfo(texref); - const struct textureReferenceAttr* texAttr = gpu->get_texattr(texref); + const struct cudaArray* cuArray = gpu->get_texarray(texname); + const struct textureInfo* texInfo = gpu->get_texinfo(texname); + const struct textureReferenceAttr* texAttr = gpu->get_texattr(texname); //assume always 2D f32 input //access array with src2 coordinates -- cgit v1.3