diff --git a/JHA/jha.cu b/JHA/jha.cu
index ec7895c..94172ff 100644
--- a/JHA/jha.cu
+++ b/JHA/jha.cu
@@ -147,6 +147,9 @@ extern "C" int scanhash_jha(int thr_id, struct work *work, uint32_t max_nonce, u
 			CUDA_LOG_ERROR();
 		}
 		cuda_get_arch(thr_id);
+		if (cuda_arch[dev_id] >= 500) {
+			applog(LOG_WARNING, "You are not using the optimal algo, please try -a jackpot");
+		}
 		gpulog(LOG_INFO, thr_id, "Intensity set to %g, %u cuda threads", throughput2intensity(throughput), throughput);
 
 		CUDA_SAFE_CALL(cudaMalloc(&d_hash[thr_id], (size_t) 64 * throughput));
diff --git a/bench.cpp b/bench.cpp
index d147a9e..a91b4c3 100644
--- a/bench.cpp
+++ b/bench.cpp
@@ -59,6 +59,7 @@ void algo_free_all(int thr_id)
 	free_cryptonight(thr_id);
 	free_decred(thr_id);
 	free_deep(thr_id);
+	free_equihash(thr_id);
 	free_keccak256(thr_id);
 	free_fresh(thr_id);
 	free_fugue256(thr_id);
diff --git a/equi/cuda_equi.cu b/equi/cuda_equi.cu
index 6fc864d..b7bcbb5 100644
--- a/equi/cuda_equi.cu
+++ b/equi/cuda_equi.cu
@@ -2081,7 +2081,7 @@ __host__ void eq_cuda_context<RB, SM, SSM, THREADS, PACKER>::solve(const char *t
 // destructor
 template <u32 RB, u32 SM, u32 SSM, u32 THREADS, typename PACKER>
 __host__
-eq_cuda_context<RB, SM, SSM, THREADS, PACKER>::~eq_cuda_context()
+void eq_cuda_context<RB, SM, SSM, THREADS, PACKER>::freemem()
 {
 	if (solutions)
 		free(solutions);
@@ -2104,6 +2104,12 @@ eq_cuda_context<RB, SM, SSM, THREADS, PACKER>::~eq_cuda_context()
 	}
 }
 
+template <u32 RB, u32 SM, u32 SSM, u32 THREADS, typename PACKER>
+__host__
+eq_cuda_context<RB, SM, SSM, THREADS, PACKER>::~eq_cuda_context()
+{
+	freemem();
+}
 
 #ifdef CONFIG_MODE_1
 template class eq_cuda_context<CONFIG_MODE_1>;
diff --git a/equi/eqcuda.hpp b/equi/eqcuda.hpp
index 68bdaf0..7cb10c9 100644
--- a/equi/eqcuda.hpp
+++ b/equi/eqcuda.hpp
@@ -90,7 +90,7 @@ template <u32 RB, u32 SM> struct equi;
 
 struct eq_cuda_context_interface
 {
-	virtual ~eq_cuda_context_interface();
+	//virtual ~eq_cuda_context_interface();
 
 	virtual void solve(const char *tequihash_header,
 		unsigned int tequihash_header_len,
@@ -99,6 +99,7 @@ struct eq_cuda_context_interface
 		fn_cancel cancelf,
 		fn_solution solutionf,
 		fn_hashdone hashdonef);
+
 public:
 	int thread_id;
 	int device_id;
@@ -125,9 +126,9 @@ class eq_cuda_context : public eq_cuda_context_interface
 		fn_cancel cancelf,
 		fn_solution solutionf,
 		fn_hashdone hashdonef);
-
 public:
 	eq_cuda_context(int thr_id, int dev_id);
+	void freemem();
 	~eq_cuda_context();
 };
 
diff --git a/equi/equihash.cpp b/equi/equihash.cpp
index 2a6e514..c9ac1fc 100644
--- a/equi/equihash.cpp
+++ b/equi/equihash.cpp
@@ -281,7 +281,11 @@ void free_equihash(int thr_id)
 	if (!init[thr_id])
 		return;
 
-	delete(solvers[thr_id]);
+	// assume config 1 was used... interface destructor seems bad
+	eq_cuda_context<CONFIG_MODE_1>* ptr = dynamic_cast<eq_cuda_context<CONFIG_MODE_1>*>(solvers[thr_id]);
+	ptr->freemem();
+	ptr = NULL;
+
 	solvers[thr_id] = NULL;
 
 	init[thr_id] = false;
@@ -291,4 +295,3 @@ void free_equihash(int thr_id)
 void eq_cuda_context_interface::solve(const char *tequihash_header, unsigned int tequihash_header_len,
 	const char* nonce, unsigned int nonce_len,
 	fn_cancel cancelf, fn_solution solutionf, fn_hashdone hashdonef) { }
-eq_cuda_context_interface::~eq_cuda_context_interface() { }