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 From 7a9c450e6b905af9ca6cdd3c7b79ad5aec535a5a Mon Sep 17 00:00:00 2001 From: Deval Shah Date: Fri, 9 Nov 2018 21:27:34 -0800 Subject: resolving merge conflicts --- libcuda/cuda_runtime_api.cc | 35 +--------------- src/abstract_hardware_model.h | 30 ++------------ src/cuda-sim/cuda-sim.cc | 94 ++++--------------------------------------- src/cuda-sim/instructions.cc | 56 -------------------------- src/gpgpu-sim/gpu-sim.cc | 25 ++---------- 5 files changed, 16 insertions(+), 224 deletions(-) (limited to 'src/abstract_hardware_model.h') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 4b50e34..f00fe52 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1425,40 +1425,7 @@ __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun ) dim3 gridDim = config.grid_dim(); dim3 blockDim = config.block_dim(); - gpgpu_t *gpu = context->get_device()->get_gpgpu(); - checkpoint *g_checkpoint; - g_checkpoint = new checkpoint(); - class memory_space *global_mem; - global_mem = gpu->get_global_memory(); - - if(gpu->resume_option ==1 && (grid->get_uid()==gpu->resume_kernel)) - { - - char f1name[2048]; - snprintf(f1name,2048,"checkpoint_files/global_mem_%d.txt", grid->get_uid()); - - g_checkpoint->load_global_mem(global_mem, f1name); - for (int i=0;iresume_CTA;i++) - grid->increment_cta_id(); - } - if(gpu->resume_option==1 && (grid->get_uid()resume_kernel)) - { - char f1name[2048]; - snprintf(f1name,2048,"checkpoint_files/global_mem_%d.txt", grid->get_uid()); - - g_checkpoint->load_global_mem(global_mem, f1name); - printf("Skipping kernel %d as resuming from kernel %d\n",grid->get_uid(),gpu->resume_kernel ); - g_cuda_launch_stack.pop_back(); - return g_last_cudaError = cudaSuccess; - - } - if(gpu->checkpoint_option==1 && (grid->get_uid()>gpu->checkpoint_kernel)) - { - printf("Skipping kernel %d as checkpoint from kernel %d\n",grid->get_uid(),gpu->checkpoint_kernel ); - g_cuda_launch_stack.pop_back(); - return g_last_cudaError = cudaSuccess; - - } + printf("GPGPU-Sim PTX: pushing kernel \'%s\' to stream %u, gridDim= (%u,%u,%u) blockDim = (%u,%u,%u) \n", kname.c_str(), stream?stream->get_uid():0, gridDim.x,gridDim.y,gridDim.z,blockDim.x,blockDim.y,blockDim.z ); stream_operation op(grid,g_ptx_sim_mode,stream); diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 2350db4..45fba76 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -381,8 +381,7 @@ public: void get_pdom_stack_top_info( unsigned *pc, unsigned *rpc ) const; unsigned get_rp() const; void print(FILE *fp) const; - void resume(char * fname) ; - void print_checkpoint (FILE *fout) const; + protected: unsigned m_warp_id; @@ -502,28 +501,14 @@ public: const char* get_ptx_inst_debug_file() const { return g_ptx_inst_debug_file; } int get_ptx_inst_debug_thread_uid() const { return g_ptx_inst_debug_thread_uid; } unsigned get_texcache_linesize() const { return m_texcache_linesize; } - int get_checkpoint_option() const {return checkpoint_option; } - int get_checkpoint_kernel() const {return checkpoint_kernel; } - int get_checkpoint_CTA() const {return checkpoint_CTA; } - int get_resume_option() const {return resume_option; } - int get_resume_kernel() const {return resume_kernel; } - int get_resume_CTA() const {return resume_CTA; } - int get_checkpoint_CTA_t() const {return checkpoint_CTA_t; } - int get_checkpoint_insn_Y() const {return checkpoint_insn_Y; } + private: // PTX options int m_ptx_convert_to_ptxplus; int m_ptx_use_cuobjdump; int m_experimental_lib_support; unsigned m_ptx_force_max_capability; - int checkpoint_option; - int checkpoint_kernel; - int checkpoint_CTA; - int resume_option; - int resume_kernel; - int resume_CTA; - int checkpoint_CTA_t; - int checkpoint_insn_Y; + int g_ptx_inst_debug_to_file; char* g_ptx_inst_debug_file; int g_ptx_inst_debug_thread_uid; @@ -535,14 +520,7 @@ private: class gpgpu_t { public: gpgpu_t( const gpgpu_functional_sim_config &config ); - int checkpoint_option; - int checkpoint_kernel; - int checkpoint_CTA; - int resume_option; - int resume_kernel; - int resume_CTA; - int checkpoint_CTA_t; - int checkpoint_insn_Y; + void* gpu_malloc( size_t size ); void* gpu_mallocarray( size_t count ); void gpu_memset( size_t dst_start_addr, int c, size_t count ); diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 642e301..6a6b307 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -2125,8 +2125,7 @@ void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL ) //before we execute, we should do PDOM analysis for functional simulation scenario. function_info *kernel_func_info = kernel.entry(); const struct gpgpu_ptx_sim_info *kernel_info = ptx_sim_kernel_info(kernel_func_info); - checkpoint *g_checkpoint; - g_checkpoint = new checkpoint(); + if (kernel_func_info->is_pdom_set()) { printf("GPGPU-Sim PTX: PDOM analysis already done for %s \n", kernel.name().c_str() ); @@ -2143,21 +2142,12 @@ void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL ) - int inst_count=50; - int cp_op= g_the_gpu->checkpoint_option; - int cp_CTA = g_the_gpu->checkpoint_CTA; - int cp_kernel= g_the_gpu->checkpoint_kernel; - cp_count= g_the_gpu->checkpoint_insn_Y; - cp_cta_resume= g_the_gpu->checkpoint_CTA_t; - int cta_launched =0; //we excute the kernel one CTA (Block) at the time, as synchronization functions work block wise while(!kernel.no_more_ctas_to_run()){ unsigned temp=kernel.get_next_cta_id_single(); - if(cp_op==0 || (cp_op==1 && cta_launched= 5000) launch_all_device_kernels(); #endif - } - else - { - kernel.increment_cta_id(); - } - cta_launched++; + } - if(cp_op==1) - { - char f1name[2048]; - snprintf(f1name,2048,"checkpoint_files/global_mem_%d.txt", kernel.get_uid() ); - g_checkpoint->store_global_mem(g_the_gpu->get_global_memory(), f1name , "%08x"); - } + @@ -2227,7 +2207,6 @@ void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL ) void functionalCoreSim::initializeCTA(unsigned ctaid_cp) { int ctaLiveThreads=0; - symbol_table * symtab= m_kernel->entry()->get_symtab(); for(int i=0; i< m_warp_count; i++){ m_warpAtBarrier[i]=false; @@ -2240,10 +2219,7 @@ void functionalCoreSim::initializeCTA(unsigned ctaid_cp) for(unsigned i=0; ithreads_per_cta();i++) { ptx_sim_init_thread(*m_kernel,&m_thread[i],0,i,m_kernel->threads_per_cta()-i,m_kernel->threads_per_cta(),this,0,i/m_warp_size,(gpgpu_t*)m_gpu, true); assert(m_thread[i]!=NULL && !m_thread[i]->is_done()); - char fname[2048]; - snprintf(fname,2048,"checkpoint_files/thread_%d_0_reg.txt",i ); - if(cp_cta_resume==1) - m_thread[i]->resume_reg_thread(fname,symtab); + ctaLiveThreads++; } @@ -2266,40 +2242,25 @@ void functionalCoreSim::createWarp(unsigned warpId) char fname[2048]; snprintf(fname,2048,"checkpoint_files/warp_%d_0_simt.txt",warpId ); - if(cp_cta_resume==1) - { - unsigned pc,rpc; - m_simt_stack[warpId]->resume(fname); - m_simt_stack[warpId]->get_pdom_stack_top_info(&pc,&rpc); - for(int i=warpId*m_warp_size; iset_npc(pc); - m_thread[i]->update_pc(); - } - } m_liveThreadCount[warpId]= liveThreadsCount; } void functionalCoreSim::execute(int inst_count, unsigned ctaid_cp) { - cp_count= m_gpu->checkpoint_insn_Y; - cp_cta_resume= m_gpu->checkpoint_CTA_t; + initializeCTA(ctaid_cp); - int count=0; + while(true){ bool someOneLive= false; bool allAtBarrier = true; for(unsigned i=0;i0 && count>inst_count && (m_kernel->get_uid()==m_gpu->checkpoint_kernel) && (ctaid_cp>=m_gpu->checkpoint_CTA) && (ctaid_cpcheckpoint_CTA_t) && m_gpu->checkpoint_option==1) - { - someOneLive=false; - break; - } + if(!someOneLive) break; if(allAtBarrier){ for(unsigned i=0;ientry()->get_symtab(); - - - unsigned ctaid =m_kernel->get_next_cta_id_single(); - if(m_gpu->checkpoint_option==1 && (m_kernel->get_uid()==m_gpu->checkpoint_kernel) && (ctaid_cp>=m_gpu->checkpoint_CTA) && (ctaid_cpcheckpoint_CTA_t)) - { - char fname[2048]; - snprintf(fname,2048,"checkpoint_files/shared_mem_%d.txt",ctaid-1 ); - g_checkpoint->store_global_mem(m_thread[0]->m_shared_mem, fname , "%08x"); - for(int i=0; i<32*m_warp_count;i++) - { - char fname[2048]; - snprintf(fname,2048,"checkpoint_files/thread_%d_%d_reg.txt",i,ctaid-1 ); - m_thread[i]->print_reg_thread(fname); - char f1name[2048]; - snprintf(f1name,2048,"checkpoint_files/local_mem_thread_%d_%d_reg.txt",i,ctaid-1 ); - g_checkpoint->store_global_mem(m_thread[i]->m_local_mem, f1name , "%08x"); - m_thread[i]->set_done(); - m_thread[i]->exitCore(); - m_thread[i]->registerExit(); - } - - for(int i=0;iprint_checkpoint(fp); - fclose(fp); - } - } } diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 31a33c6..f57a3f7 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -183,63 +183,7 @@ void ptx_thread_info::set_reg( const symbol *reg, const ptx_reg_t &value ) m_last_set_operand_value = value; } -void ptx_thread_info::print_reg_thread(char * fname) -{ - - FILE *fp= fopen(fname,"w"); - assert(fp!=NULL); - - int size = m_regs.size(); - - if(size>0) - { - reg_map_t reg = m_regs.back(); - - typename reg_map_t::const_iterator it; - for (it = reg.begin(); it != reg.end(); ++it) - { - const std::string &name = it->first->name(); - const std::string &dec= it->first->decl_location(); - unsigned size = it->first->get_size_in_bytes(); - fprintf(fp,"%s %llu %s %d\n",name.c_str(),it->second, dec.c_str(),size ); - - } - //m_regs.pop_back(); - } - fclose(fp); - - } - -void ptx_thread_info::resume_reg_thread(char * fname, symbol_table * symtab) -{ - - FILE * fp2 = fopen(fname, "r"); - assert(fp2!=NULL); - //m_regs.push_back( reg_map_t() ); - char line [ 200 ]; - while ( fgets ( line, sizeof line, fp2 ) != NULL ) - { - symbol *reg; - char * pch; - unsigned size; - pch = strtok (line," "); - char * name =pch; - reg= symtab->lookup(name); - ptx_reg_t data; - pch = strtok (NULL," "); - data = atoi(pch); - pch = strtok (NULL," "); - char * decl= pch; - pch = strtok (NULL," "); - size = atoi(pch); - - - m_regs.back()[reg] = data; - } - fclose ( fp2 ); -} - ptx_reg_t ptx_thread_info::get_reg( const symbol *reg ) { diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 79a6fcd..c706f23 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -1414,38 +1414,19 @@ void shader_core_ctx::issue_block2core( kernel_info_t &kernel ) // bind functional simulation state of threads to hardware resources (simulation) warp_set_t warps; unsigned nthreads_in_block= 0; - function_info *kernel_func_info = kernel.entry(); - symbol_table * symtab= kernel_func_info->get_symtab(); - unsigned ctaid= kernel.get_next_cta_id_single(); - checkpoint *g_checkpoint= new checkpoint(); + for (unsigned i = start_thread; iwarp_size; nthreads_in_block += ptx_sim_init_thread(kernel,&m_thread[i],m_sid,i,cta_size-(i-start_thread),m_config->n_thread_per_shader,this,free_cta_hw_id,warp_id,m_cluster->get_gpu()); m_threadState[i].m_active = true; - // load thread local memory and register file - if(m_gpu->resume_option==1 && kernel.get_uid()==m_gpu->resume_kernel && ctaid>=m_gpu->resume_CTA && ctaidcheckpoint_CTA_t ) - { - char fname[2048]; - snprintf(fname,2048,"checkpoint_files/thread_%d_%d_reg.txt",i%cta_size,ctaid ); - m_thread[i]->resume_reg_thread(fname,symtab); - char f1name[2048]; - snprintf(f1name,2048,"checkpoint_files/local_mem_thread_%d_%d_reg.txt",i%cta_size,ctaid); - g_checkpoint->load_global_mem(m_thread[i]->m_local_mem, f1name); - } - // + warps.set( warp_id ); } assert( nthreads_in_block > 0 && nthreads_in_block <= m_config->n_thread_per_shader); // should be at least one, but less than max m_cta_status[free_cta_hw_id]=nthreads_in_block; - if(m_gpu->resume_option==1 && kernel.get_uid()==m_gpu->resume_kernel && ctaid>=m_gpu->resume_CTA && ctaidcheckpoint_CTA_t ) - { - char f1name[2048]; - snprintf(f1name,2048,"checkpoint_files/shared_mem_%d.txt", ctaid); - - g_checkpoint->load_global_mem(m_thread[start_thread]->m_shared_mem, f1name); - } + // now that we know which warps are used in this CTA, we can allocate // resources for use in CTA-wide barrier operations m_barriers.allocate_barrier(free_cta_hw_id,warps); -- cgit v1.3 From 642818ae5ff61c1544bcce9e7ba2dd0aea47ea6a Mon Sep 17 00:00:00 2001 From: Deval Shah Date: Fri, 9 Nov 2018 21:29:18 -0800 Subject: Adding checkpoint support --- libcuda/cuda_runtime_api.cc | 35 +++++++++++++++- src/abstract_hardware_model.h | 30 ++++++++++++-- src/cuda-sim/cuda-sim.cc | 94 +++++++++++++++++++++++++++++++++++++++---- src/cuda-sim/instructions.cc | 56 ++++++++++++++++++++++++++ src/gpgpu-sim/gpu-sim.cc | 25 ++++++++++-- 5 files changed, 224 insertions(+), 16 deletions(-) (limited to 'src/abstract_hardware_model.h') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index f00fe52..4b50e34 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1425,7 +1425,40 @@ __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun ) dim3 gridDim = config.grid_dim(); dim3 blockDim = config.block_dim(); - + gpgpu_t *gpu = context->get_device()->get_gpgpu(); + checkpoint *g_checkpoint; + g_checkpoint = new checkpoint(); + class memory_space *global_mem; + global_mem = gpu->get_global_memory(); + + if(gpu->resume_option ==1 && (grid->get_uid()==gpu->resume_kernel)) + { + + char f1name[2048]; + snprintf(f1name,2048,"checkpoint_files/global_mem_%d.txt", grid->get_uid()); + + g_checkpoint->load_global_mem(global_mem, f1name); + for (int i=0;iresume_CTA;i++) + grid->increment_cta_id(); + } + if(gpu->resume_option==1 && (grid->get_uid()resume_kernel)) + { + char f1name[2048]; + snprintf(f1name,2048,"checkpoint_files/global_mem_%d.txt", grid->get_uid()); + + g_checkpoint->load_global_mem(global_mem, f1name); + printf("Skipping kernel %d as resuming from kernel %d\n",grid->get_uid(),gpu->resume_kernel ); + g_cuda_launch_stack.pop_back(); + return g_last_cudaError = cudaSuccess; + + } + if(gpu->checkpoint_option==1 && (grid->get_uid()>gpu->checkpoint_kernel)) + { + printf("Skipping kernel %d as checkpoint from kernel %d\n",grid->get_uid(),gpu->checkpoint_kernel ); + g_cuda_launch_stack.pop_back(); + return g_last_cudaError = cudaSuccess; + + } printf("GPGPU-Sim PTX: pushing kernel \'%s\' to stream %u, gridDim= (%u,%u,%u) blockDim = (%u,%u,%u) \n", kname.c_str(), stream?stream->get_uid():0, gridDim.x,gridDim.y,gridDim.z,blockDim.x,blockDim.y,blockDim.z ); stream_operation op(grid,g_ptx_sim_mode,stream); diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 45fba76..2350db4 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -381,7 +381,8 @@ public: void get_pdom_stack_top_info( unsigned *pc, unsigned *rpc ) const; unsigned get_rp() const; void print(FILE *fp) const; - + void resume(char * fname) ; + void print_checkpoint (FILE *fout) const; protected: unsigned m_warp_id; @@ -501,14 +502,28 @@ public: const char* get_ptx_inst_debug_file() const { return g_ptx_inst_debug_file; } int get_ptx_inst_debug_thread_uid() const { return g_ptx_inst_debug_thread_uid; } unsigned get_texcache_linesize() const { return m_texcache_linesize; } - + int get_checkpoint_option() const {return checkpoint_option; } + int get_checkpoint_kernel() const {return checkpoint_kernel; } + int get_checkpoint_CTA() const {return checkpoint_CTA; } + int get_resume_option() const {return resume_option; } + int get_resume_kernel() const {return resume_kernel; } + int get_resume_CTA() const {return resume_CTA; } + int get_checkpoint_CTA_t() const {return checkpoint_CTA_t; } + int get_checkpoint_insn_Y() const {return checkpoint_insn_Y; } private: // PTX options int m_ptx_convert_to_ptxplus; int m_ptx_use_cuobjdump; int m_experimental_lib_support; unsigned m_ptx_force_max_capability; - + int checkpoint_option; + int checkpoint_kernel; + int checkpoint_CTA; + int resume_option; + int resume_kernel; + int resume_CTA; + int checkpoint_CTA_t; + int checkpoint_insn_Y; int g_ptx_inst_debug_to_file; char* g_ptx_inst_debug_file; int g_ptx_inst_debug_thread_uid; @@ -520,7 +535,14 @@ private: class gpgpu_t { public: gpgpu_t( const gpgpu_functional_sim_config &config ); - + int checkpoint_option; + int checkpoint_kernel; + int checkpoint_CTA; + int resume_option; + int resume_kernel; + int resume_CTA; + int checkpoint_CTA_t; + int checkpoint_insn_Y; void* gpu_malloc( size_t size ); void* gpu_mallocarray( size_t count ); void gpu_memset( size_t dst_start_addr, int c, size_t count ); diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 6a6b307..642e301 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -2125,7 +2125,8 @@ void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL ) //before we execute, we should do PDOM analysis for functional simulation scenario. function_info *kernel_func_info = kernel.entry(); const struct gpgpu_ptx_sim_info *kernel_info = ptx_sim_kernel_info(kernel_func_info); - + checkpoint *g_checkpoint; + g_checkpoint = new checkpoint(); if (kernel_func_info->is_pdom_set()) { printf("GPGPU-Sim PTX: PDOM analysis already done for %s \n", kernel.name().c_str() ); @@ -2142,12 +2143,21 @@ void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL ) + int inst_count=50; + int cp_op= g_the_gpu->checkpoint_option; + int cp_CTA = g_the_gpu->checkpoint_CTA; + int cp_kernel= g_the_gpu->checkpoint_kernel; + cp_count= g_the_gpu->checkpoint_insn_Y; + cp_cta_resume= g_the_gpu->checkpoint_CTA_t; + int cta_launched =0; //we excute the kernel one CTA (Block) at the time, as synchronization functions work block wise while(!kernel.no_more_ctas_to_run()){ unsigned temp=kernel.get_next_cta_id_single(); + if(cp_op==0 || (cp_op==1 && cta_launched= 5000) launch_all_device_kernels(); #endif - + } + else + { + kernel.increment_cta_id(); + } + cta_launched++; } - + if(cp_op==1) + { + char f1name[2048]; + snprintf(f1name,2048,"checkpoint_files/global_mem_%d.txt", kernel.get_uid() ); + g_checkpoint->store_global_mem(g_the_gpu->get_global_memory(), f1name , "%08x"); + } @@ -2207,6 +2227,7 @@ void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL ) void functionalCoreSim::initializeCTA(unsigned ctaid_cp) { int ctaLiveThreads=0; + symbol_table * symtab= m_kernel->entry()->get_symtab(); for(int i=0; i< m_warp_count; i++){ m_warpAtBarrier[i]=false; @@ -2219,7 +2240,10 @@ void functionalCoreSim::initializeCTA(unsigned ctaid_cp) for(unsigned i=0; ithreads_per_cta();i++) { ptx_sim_init_thread(*m_kernel,&m_thread[i],0,i,m_kernel->threads_per_cta()-i,m_kernel->threads_per_cta(),this,0,i/m_warp_size,(gpgpu_t*)m_gpu, true); assert(m_thread[i]!=NULL && !m_thread[i]->is_done()); - + char fname[2048]; + snprintf(fname,2048,"checkpoint_files/thread_%d_0_reg.txt",i ); + if(cp_cta_resume==1) + m_thread[i]->resume_reg_thread(fname,symtab); ctaLiveThreads++; } @@ -2242,25 +2266,40 @@ void functionalCoreSim::createWarp(unsigned warpId) char fname[2048]; snprintf(fname,2048,"checkpoint_files/warp_%d_0_simt.txt",warpId ); + if(cp_cta_resume==1) + { + unsigned pc,rpc; + m_simt_stack[warpId]->resume(fname); + m_simt_stack[warpId]->get_pdom_stack_top_info(&pc,&rpc); + for(int i=warpId*m_warp_size; iset_npc(pc); + m_thread[i]->update_pc(); + } + } m_liveThreadCount[warpId]= liveThreadsCount; } void functionalCoreSim::execute(int inst_count, unsigned ctaid_cp) { - + cp_count= m_gpu->checkpoint_insn_Y; + cp_cta_resume= m_gpu->checkpoint_CTA_t; initializeCTA(ctaid_cp); - + int count=0; while(true){ bool someOneLive= false; bool allAtBarrier = true; for(unsigned i=0;i0 && count>inst_count && (m_kernel->get_uid()==m_gpu->checkpoint_kernel) && (ctaid_cp>=m_gpu->checkpoint_CTA) && (ctaid_cpcheckpoint_CTA_t) && m_gpu->checkpoint_option==1) + { + someOneLive=false; + break; + } if(!someOneLive) break; if(allAtBarrier){ for(unsigned i=0;ientry()->get_symtab(); + + + unsigned ctaid =m_kernel->get_next_cta_id_single(); + if(m_gpu->checkpoint_option==1 && (m_kernel->get_uid()==m_gpu->checkpoint_kernel) && (ctaid_cp>=m_gpu->checkpoint_CTA) && (ctaid_cpcheckpoint_CTA_t)) + { + char fname[2048]; + snprintf(fname,2048,"checkpoint_files/shared_mem_%d.txt",ctaid-1 ); + g_checkpoint->store_global_mem(m_thread[0]->m_shared_mem, fname , "%08x"); + for(int i=0; i<32*m_warp_count;i++) + { + char fname[2048]; + snprintf(fname,2048,"checkpoint_files/thread_%d_%d_reg.txt",i,ctaid-1 ); + m_thread[i]->print_reg_thread(fname); + char f1name[2048]; + snprintf(f1name,2048,"checkpoint_files/local_mem_thread_%d_%d_reg.txt",i,ctaid-1 ); + g_checkpoint->store_global_mem(m_thread[i]->m_local_mem, f1name , "%08x"); + m_thread[i]->set_done(); + m_thread[i]->exitCore(); + m_thread[i]->registerExit(); + } + + for(int i=0;iprint_checkpoint(fp); + fclose(fp); + } + } } diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index f57a3f7..31a33c6 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -183,7 +183,63 @@ void ptx_thread_info::set_reg( const symbol *reg, const ptx_reg_t &value ) m_last_set_operand_value = value; } +void ptx_thread_info::print_reg_thread(char * fname) +{ + + FILE *fp= fopen(fname,"w"); + assert(fp!=NULL); + + int size = m_regs.size(); + + if(size>0) + { + reg_map_t reg = m_regs.back(); + + typename reg_map_t::const_iterator it; + for (it = reg.begin(); it != reg.end(); ++it) + { + const std::string &name = it->first->name(); + const std::string &dec= it->first->decl_location(); + unsigned size = it->first->get_size_in_bytes(); + fprintf(fp,"%s %llu %s %d\n",name.c_str(),it->second, dec.c_str(),size ); + + } + //m_regs.pop_back(); + } + fclose(fp); + + } + +void ptx_thread_info::resume_reg_thread(char * fname, symbol_table * symtab) +{ + + FILE * fp2 = fopen(fname, "r"); + assert(fp2!=NULL); + //m_regs.push_back( reg_map_t() ); + char line [ 200 ]; + while ( fgets ( line, sizeof line, fp2 ) != NULL ) + { + symbol *reg; + char * pch; + unsigned size; + pch = strtok (line," "); + char * name =pch; + reg= symtab->lookup(name); + ptx_reg_t data; + pch = strtok (NULL," "); + data = atoi(pch); + pch = strtok (NULL," "); + char * decl= pch; + pch = strtok (NULL," "); + size = atoi(pch); + + + m_regs.back()[reg] = data; + } + fclose ( fp2 ); +} + ptx_reg_t ptx_thread_info::get_reg( const symbol *reg ) { diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index c706f23..79a6fcd 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -1414,19 +1414,38 @@ void shader_core_ctx::issue_block2core( kernel_info_t &kernel ) // bind functional simulation state of threads to hardware resources (simulation) warp_set_t warps; unsigned nthreads_in_block= 0; - + function_info *kernel_func_info = kernel.entry(); + symbol_table * symtab= kernel_func_info->get_symtab(); + unsigned ctaid= kernel.get_next_cta_id_single(); + checkpoint *g_checkpoint= new checkpoint(); for (unsigned i = start_thread; iwarp_size; nthreads_in_block += ptx_sim_init_thread(kernel,&m_thread[i],m_sid,i,cta_size-(i-start_thread),m_config->n_thread_per_shader,this,free_cta_hw_id,warp_id,m_cluster->get_gpu()); m_threadState[i].m_active = true; - + // load thread local memory and register file + if(m_gpu->resume_option==1 && kernel.get_uid()==m_gpu->resume_kernel && ctaid>=m_gpu->resume_CTA && ctaidcheckpoint_CTA_t ) + { + char fname[2048]; + snprintf(fname,2048,"checkpoint_files/thread_%d_%d_reg.txt",i%cta_size,ctaid ); + m_thread[i]->resume_reg_thread(fname,symtab); + char f1name[2048]; + snprintf(f1name,2048,"checkpoint_files/local_mem_thread_%d_%d_reg.txt",i%cta_size,ctaid); + g_checkpoint->load_global_mem(m_thread[i]->m_local_mem, f1name); + } + // warps.set( warp_id ); } assert( nthreads_in_block > 0 && nthreads_in_block <= m_config->n_thread_per_shader); // should be at least one, but less than max m_cta_status[free_cta_hw_id]=nthreads_in_block; - + if(m_gpu->resume_option==1 && kernel.get_uid()==m_gpu->resume_kernel && ctaid>=m_gpu->resume_CTA && ctaidcheckpoint_CTA_t ) + { + char f1name[2048]; + snprintf(f1name,2048,"checkpoint_files/shared_mem_%d.txt", ctaid); + + g_checkpoint->load_global_mem(m_thread[start_thread]->m_shared_mem, f1name); + } // now that we know which warps are used in this CTA, we can allocate // resources for use in CTA-wide barrier operations m_barriers.allocate_barrier(free_cta_hw_id,warps); -- cgit v1.3