implemented unbind, currently only affects cudaArray map

5 files changed, 17 insertions, 0 deletions

diff --git a/0517_13-24cudaReserved.log b/0517_13-24cudaReserved.log
deleted file mode 100644
index e69de29..0000000
--- a/0517_13-24cudaReserved.log
+++ /dev/null
diff --git a/517-1418cudaReserved.log b/517-1418cudaReserved.log
deleted file mode 100644
index e69de29..0000000
--- a/517-1418cudaReserved.log
+++ /dev/null
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<unsigned,function_info*> g_pc_to_finfo;
 std::vector<ptx_instruction*> function_info::s_g_pc_to_insn;