 		 A_tex2, A_nbit, A_stride, A_batchstride,
 		 beta,
 		 (float4*)C_ptr, C_stride, C_batchstride);
+	BF_CHECK_CUDA_EXCEPTION(cudaGetLastError(), BF_STATUS_INTERNAL_ERROR);
+	
+	BF_CHECK_CUDA_EXCEPTION(
+		cudaDestroyTextureObject(A_tex2),
+		BF_STATUS_INTERNAL_ERROR);
 #endif // BF_USE_DIAGONAL_KERNEL
-	
-	cudaDestroyTextureObject(A_tex2);
-	cudaDestroyTextureObject(A_tex);
 }
 
 template<int SIZE> struct shflable_type {};

	BF_ASSERT_EXCEPTION(A_batchstride % 2 == 0, BF_STATUS_UNSUPPORTED_STRIDE);
	BF_ASSERT_EXCEPTION(C_stride      % 2 == 0, BF_STATUS_UNSUPPORTED_STRIDE);
	BF_ASSERT_EXCEPTION(C_batchstride % 2 == 0, BF_STATUS_UNSUPPORTED_STRIDE);
	// TODO: Assert supported limits on N and K based on texture constraints
	
	// Note: The kernel is 2x vectorized (i.e., operates on Jones vectors)

	memset(&texDesc2, 0, sizeof(texDesc2));
	texDesc2.readMode = tex_read_mode;
	cudaTextureObject_t A_tex2 = 0;
	BF_CHECK_CUDA_EXCEPTION(
		cudaCreateTextureObject(&A_tex2, &resDesc2, &texDesc2, NULL),
		BF_STATUS_INTERNAL_ERROR);
	
	// TODO: Clean this up a bit
	grid.x = nblock_m;
	enum { N_THREAD_DIAG = 4 };
	block.y = N_THREAD_DIAG;
	bf_cherk_N_diagonal_kernel<M_THREAD, N_THREAD_DIAG, 2, 2>
		<<<grid, block, 0, stream1>>>
		(N, K, nbatch,
