 	memset(&texDesc2, 0, sizeof(texDesc2));
 	texDesc2.readMode = tex_read_mode;
 	cudaTextureObject_t A_tex2 = 0;
-	cudaCreateTextureObject(&A_tex2, &resDesc2, &texDesc2, NULL);
+	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;
-#if BF_USE_DIAGONAL_KERNEL
 	bf_cherk_N_diagonal_kernel<M_THREAD, N_THREAD_DIAG, 2, 2>
 		<<<grid, block, 0, stream1>>>
 		(N, K, nbatch,

	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)

	resDesc.res.linear.desc.z = A_nbit;
	resDesc.res.linear.desc.w = A_nbit;
	resDesc.res.linear.sizeInBytes = A_nelement_total * 4 * (A_nbit / 8);
	
	cudaTextureDesc texDesc;
	memset(&texDesc, 0, sizeof(texDesc));
	texDesc.readMode = tex_read_mode;
	cudaTextureObject_t A_tex = 0;
	BF_CHECK_CUDA_EXCEPTION(
		cudaCreateTextureObject(&A_tex, &resDesc, &texDesc, NULL),
		BF_STATUS_INTERNAL_ERROR);
	
	cuda::child_stream stream0(stream);
	cuda::child_stream stream1(stream);
