关于CUDA限制的一点疑惑，请大牛不吝赐教，感激感激！

[书生小小鱼]

各位大牛，先给大家拜年了！
最近研究碰到个问题，查找一个二维数组的成员在另一个二维数组中的包含情况。
举例来说，
a={{1,2},{2,3},{3,4}};
b={{1,3},{2,3}};
那么查询b中的一元数组包含在a中，期望返回c={0,1}，其中c[1]=1表示b[1]包含在a中
由于实际查询中，a的size大约是33万*12，b行数将近6万，之前用python写过一个查询，大约需要2小时才能计算完，而且汗颜的是，同样的计算需要进行1500次左右，所以python算的话，电脑开机得算半个月。
接触到CUDA之后觉得是个解决问题的好办法，试着写了个CUDA，在测试阶段能跑通，但实际应用起来返回值错的离谱，仔细研究了一下，发现a的最大行数M在12000行左右，b的最大行数O在8192，如果分批计算，涉及a，b的两个循环分割，有点头大，看资料CUDA支持百万线程，因此请教一下各位大牛是我哪里搞错了。

代码附上，静候指教！
由于小弟非计算机专业出身，为用CUDA开始看C语言，学了不到一个月，代码中如有其他问题也请各位大牛不吝赐教！
自己也没什么本事回答其他人的帖子赚积分，剩30分请教这个问题！

#include <stdio.h>
#include <string.h>
#include <stdlib.h>
#include "../book.h"
#define M 12000  
#define N 12    
#define O 8192

/*定义判断二元数组y中哪些项包含在二元数组x中，结果返回至z*/

/*x[M][N];y[O][N];z[O]*/
__global__ void y_in_x(int *x, int *y, int *z){
	int a[M] = { 0 };
	int b[N] = { 0 };
	int c = 0;
	int d = 0;
	int e = threadIdx.x + blockIdx.x * blockDim.x;
	while (e < O){
		for (int i = 0; i < M; i++){
			for (int j = 0; j < N; j++){
				if (*((int*)x + N*i + j) == *((int*)y + N*e + j))
					b[j] = 1;
				else
					b[j] = 0;
			}
			c = 0;
			for (int k = 0; k < N; k++){
				c = c + b[k];
			}
			if (c == N)
				a[i] = 1;
			else
				a[i] = 0;
		}
		d = 0;
		for (int l = 0; l < M; l++){
			d += a[l];
		}
		if (d > 0)
			*(z + e) = 1;
		else
			*(z + e) = 0;
		e += blockDim.x * gridDim.x;
	}
};


int main(void)
{
	int *a = NULL;
	int *b = NULL;
	int *c = NULL;
	int *dev_a = NULL;
	int *dev_b = NULL;
	int *dev_c = NULL;

	a = (int*)malloc(M*N * sizeof(int));
	b = (int*)malloc(O*N * sizeof(int));
	c = (int*)malloc(O * sizeof(int));

	HANDLE_ERROR(cudaMalloc((void**)&dev_a, M*N * sizeof(int)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_b, O*N * sizeof(int)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_c, O * sizeof(int)));

	for (int i = 0; i < M; i++){                                        /*给a在CPU赋值*/
		for (int j = 0; j < N; j++){
			*(a + N*i + j) = i * 10 + j + 1;
		}
	}

	for (int i = 0; i < O; i++){                                       /*给b在CPU赋值*/
		for (int j = 0; j < N; j++){
			*(b + N * i + j) = i * 10 + j + 1;
		}
	}

	HANDLE_ERROR(cudaMemcpy(dev_a, a, M * N * sizeof(int),
		cudaMemcpyHostToDevice));
	HANDLE_ERROR(cudaMemcpy(dev_b, b, O * N * sizeof(int),
		cudaMemcpyHostToDevice));

	y_in_x << <256, 256 >> >(dev_a, dev_b, dev_c);

	HANDLE_ERROR(cudaMemcpy(c, dev_c, O * sizeof(int),
		cudaMemcpyDeviceToHost));

	HANDLE_ERROR(cudaFree(dev_a));
	HANDLE_ERROR(cudaFree(dev_b));
	HANDLE_ERROR(cudaFree(dev_c));

	free(a);
	free(b);
	free(c);
	return 0;
}

[dashuboluo]

我觉得你对 e = threadIdx.x + blockIdx.x * blockDim.x; 中的 e 理解出现了一点问题，它不应该当作一个循环变量来用的。 它对应的是每一个独立的线程，这里最直观的就是对应b中的一个数组。 因此最小白的修改方法（本人小白一名）就是将while循环拆了，并且同时将e += blockDim.x * gridDim.x; 删去，这样就是让每一条线程去执行 b中的一个数组去与a中的全部数组进行比对。（纯想法，没有实践过） 不过这样做对显存的读写次数可能会比较大，可以考虑使用共享内存，每次先读取a中的某几个数组放进共享内存内，比完再取a中另外的数组进行比，直至把a中的所有数组比较过。这样或许可能会快一点。（纯想法，没有实践过） 最后可能要做的是，考虑kernel的<<<>>>要怎么分配了，同时要避免共享内存的不足。