revert to before texture bug

3 files changed, 4 insertions, 158 deletions

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 struct textureReference*,const struct cudaArray*>::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 struct textureReference*, const struct cudaArray*>::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 struct textureReference*, const struct textureInfo*>::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 struct textureReference*, const struct textureInfo*>::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 struct textureReference*, const struct textureReferenceAttr*>::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<unsigned,function_info*> g_pc_to_finfo;
 std::vector<ptx_instruction*> function_info::s_g_pc_to_insn;